CUDA/Open CL

CUDA/OpenCLArquiteturas Avancadas de Computadores

Krissia de Zawadzki

Instituto de Fısica de Sao Carlos - Universidade de Sao Paulo

06 de Maio 2014

Krissia de Zawadzki CUDA/OpenCL 1 / 61

CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos Conclusao

Outline

1 CUDA - Introducao

2 GPU e CUDA

3 Programando em CUDA

4 OpenCL

5 Caos

6 Conclusao



CUDA - Introducao

CUDA

Compute Unified Device Architecture

3 Plataforma de computacao paralela emodelo de programacao

3 Desenvolvido pela NVIDIA eimplementada para GPU’s NVIDIA

3 Conjunto de instrucoes e memoriadiretamente acessıveis ao programador



Background historico


2002: Stanford University

Steam processing

Prototipo e arquitetura muito parecidacom GPU

Baseline programmablestream processor





2002: GeForce 3 e ATI Radeon 9700

3 Shaders programaveis





2006: GeForce 8

Arquitetura unificada → CUDA!





2006: GeForce 8

3 Programabilidade realmente flexıvel

3 Revolucionou os conceitos depipeline de pixel e vertices

3 Cadeia de processadores



GPU’s com suporte para CUDA


GeForce > 8

3 8, 9, 100, 200, 400, 500 e

600-series (mın 256MB memloc)

3 1.0, 1.1, 1.2, 1.3, 2.0, 2.1, 3.0,3.5 e 5.0

GeForce GTX-750 (5.0)

NVS

3 Quadro 295, 420, 450

3 NVIDIA NVS 300, 315, 510

3 1.1, 1.2, 2.1 e 3.0

NVIDIA NVS 510 (3.0)





QUADRO

3 Quadro NVS, Quadro FX, Quadro K

3 1.0, 1.1, 1.2, 1.3, 2.0, 2.1, 3.0 e3.5

QUADRO K600 (3.5)

Tesla

3 D780, C870, C1060, C2050/2070,

C2075

3 K20, K40

3 1.0, 1.3, 2.0, 3.5

Tesla K20 (3.5)



Aplicacoes

Alem de processamento grafico



Aplicacoes



Aplicacoes

Ganho de desempenho em aplicacoes cientıficas



GPU

GPU vs. CPU

CPU

3 Codigos sequenciais

3 Baixa latencia

3 Controle complexo

GPU

3 Paralelismo de dados

3 Alto throughtput

3 Aritmetica com pouco controle



GPU

GPU vs. CPU

CPU

3 Fluxo iterativo

7 Tempo de computacao

7

GPU

3 Operacoes simultaneas

7 Desvio de fluxo

7



GPU

CPU vs. GPU



GPU

CPU vs. GPU



Linguagens e modelos de programacao paralela

Linguagens e modelos de programacao paralela

OpenMP

shared memory

limite de centenas de nos

CUDA

alta escalabilidade

portabilidade e mais simples!

MPI

capacidade de nos > 100.000

esforco para portar o codigo

OpenCL

modelo de programacao padronizado

suporte para AMD/ATI, NVIDIA,

Apple e Intel



Arquitetura da GPU

GPU - Unified processor array (GeForce 8800 GTX)



Estrutura de um programa CUDA


Trechos seriais ou com fraco paralelismo no codigo C do host

Porcao altamente paralela no codigo C do kernel associado ao device




CUDA Threads paralelas

// de dados: Todas as threads rodam o mesmo codigo

threadIdx: identificador da thread → & de memoria e controle




CUDA Thread Blocks

Thread Blocks: Array (dim 2x2) the threads que cooperam entre si

bloco: memoria compartilhada, operacoes atomicas ebarreiras de sincronizacao

blockIdx: identificador do bloco em um grid




CUDA Id’s

Ids sao uteis para identificar os

dados sob os quais cada thread ira

trabalhar

Conveniente para simplificar o &

de memoria em dados

multidimensionaisblockIdx:1D (blockIdx.x)

2D (blockIdx.x, blockIdx.y)

threadIdx:1D (threadIdx.x),2D (threadIdx.x, threadIdx.y),

3D (threadIdx.x, threadIdx.y, threadIdx.z)




Parametros de configuracao

block:dimBlock(Widthx , Widthy , , Widthz)

grid:

dimGrid(Wgridx , Wgridy , Wgridz)

Kernel launching

// Setup the execution configuration

dim3 dimBlock(Width , Width , 1);

dim3 dimGrid(1, 1, 1);

// Launch the device computation threads!

MyKernelFunction <<<dimGrid , dimBlock >>>(args);



Modelo de Memoria CUDA


Memoria Global:

comunicacao host-device

R/W

conteudo visıvel por todas as

threads

tipicamente implementada

como DRAM

acesso de longa latencia

(400-800 ciclos)

7 congestionamento

throughput limitado (a 177

GB/s na GTX8800)





Memoria Constante:

read only

baixa latencia e alta largura

de banda quando todas as

threadas acessam o mesmo

local





Memoria Compartilhada:

3 rapida

altamente paralela

apenas um bloco tem acesso





Registradores

3 o componente da memoriada GPU mais rapido

acessıvel por uma threada



Variaveis CUDA

Variaveis CUDA

kernel: a variavel deve ser declarada no escopo da funcao kernel → fica

disponıvel somente no kernel

application: a variavel deve ser declarada fora de qualquer funcao

constant: a variavel deve ser declarada fora de qualquer funcao →limitado (a 64KB na GTX8800)



CUDA Kernel Functions


Kernel functions:

implementam o trechoparalelo de codigo a serexecutado no device

Sua chamada pode ser feitacom as configuracoes deblocos e de threads




Atribuicao de threads

base block-by-block

7 Runtime systemcoordena os blocos e asthreads a seremexecutadas: mantem alista de blocos e associanovos blocos a SM’slivres

recursos do SM

unidades aritmeticas

numero de threads quepodem ser rastreadas eescalonadassimultaneamente




CUDA Warps

Warps : conjunto de threads comındices consecutivos

A capacidade do warp (num.de threads) e dependente daimplementacao

Warp e a unidade paraescalonar threads no SM

SIMD

ordem qualquer entre warps

7 divergencia causada por branchs




Compilador NVCC

Baseado no Open64(opensource originario doMIPSPro - SGI).

Implementado em C e C++.

NVidia atualmente investindono LLVM.

Existe um utilitario queconverte LLVM IR (geradopor qualquer frontend decompilador LLVM) em PTX,que pode ser programado nasGPUs NVidia.




Deficiencias

Equipe preferiu implementar codigo para o desafio ECC2K-130diretamente em codigo de maquina.

BERNSTEIN, D. J. et al. Usable Assembly Language for GPUs: ASuccess Story. In: Workshop Records of Special-Purpose Hardwarefor Attacking Cryptographic Systems – SHARCS 2012. [s.n.], 2012.p. 169–178.

Compilador NVCC muito lento para lidar com kernels contendomuitas instrucoes.

Registradores alocados de forma pouco eficiente – muitas variaveisacabaram tendo de ser alocadas pelo NVCC na memoriacompartilhada.

Varios truques necessarios para obter uma implementacao em Caceitavel. Implementacao em Assembly 148% mais rapida quemelhor implementacao em C.



Programando em CUDA: exemplo

Multiplicacao Matricial: o Hello World do CUDA

𝑃 = 𝑀 *𝑁

𝑃𝑖𝑗 =∑𝑘

𝑀𝑖,𝑘𝑁𝑘,𝑗

3 Paralelismo de dados!

Cada elemento 𝑃𝑖𝑗 de 𝑃 pode sercalculado simultaneamente aosdemais!




Representacao matricial em C

Alocacao de memoria no C para arrays bidimensionais:




Codigo main C sequencial (host)

int main(void){//1. Alocamos e inicializamos as matrizes M, N e P// Funcoes I/O leem as matrizes M e N...//2. Multiplicacao M * NMatMul(M,N,P,Width );...//3. Funcao I/O para escrever a saida P// Liberamos a memoria de M, N e Preturn 0;

}




Funcao C sequencial (host)

void MatMul(float *M; float *N, float *P, int Width){

for(int i = 0; i < Width; ++i)for (int j = 0; j < Width; ++j){

float sum = 0;for(int k = 0; k < Width; ++k){

float m = M[i*Width + k];float n = N[k*width + j];sum += m * n;

}P[i * Width + j] = sum;

}}




Portando o codigo para CUDA - alocacao de memoria

cudaMalloc()

aloca um objeto naMemoria Global

parametros: endereco deum ponteiro para o objetoalocado, tamanho do objeto

cudaFree()

libera um objeto naMemoria Global

parametro: ponteiro para oobjeto




Portando o codigo para CUDA - alocacao

Exemplo:

int Width =64;float* Md , Nd;int size = Width*Width*sizeof(float );

cudaMalloc ((void **)&Md, size);cudaMalloc ((void **)&Nd, size);...cudaFree(Md);cudaFree(Nd);




Portando o codigo para CUDA - transferencia de dados

cudaMemcpy()

transfere dados entre o hoste o device

Assıncrona

parametros:ponteiro para o destino

ponteiro para a fonte numero

de bytes a serem copiados

tipo de transferencia

tipos:

Host to Host

Host to Device

Device to Host

Device to Device




Portando o codigo para CUDA - transferencia de dados

Exemplo:

...int size = Width*Width*sizeof(float );...cudaMemcpy(Md, M, size , cudaMemcyHostToDevice );cudaMemcpy(Nd, N, size , cudaMemcyHostToDevice );...cudaMemcpy(Pd, P, size , cudaMemcyDeviceToHost );




Portando o codigo para CUDA - funcao MatMul no device

void MatMul(float *M; float *N, float *P, int Width){int size = Width * Width * sizeof(float );float* Md , Nd, Pd;//1. Alocamos memoria no device para M, N e PcudaMalloc ((void **)) &Md, size);cudaMemcpy(Md, M, size , cudaMemcyHostToDevice );cudaMalloc ((void **)) &Nd, size);cudaMemcpy(Nd, N, size , cudaMemcyHostToDevice );cudaMalloc ((void **)) &Pd, size);

//2. Evocamos a funcao kernel para a multiplicacao

//3. Copiamos o resultado P para a memoria do hostcudaMemcpy(Pd, P, size , cudaMemcyDeviceToHost );// Liberamos as memorias de M, N e P no devicecudaFree(Md); cudaFree(Nd); cudaFree(Pd);

}




Funcao kernel MatMul




Kernel function - um pouco mais sobre especificacoes

global define uma funcao kernel

device e host podem ser usadas simultaneamente

7 recursoes

7 variaveis estaticas

7 chamadas indiretas de funcoes por ponteiros




Configuracao de execucao

Exemplo: Definir a multiplicacao matricial quando 𝑊𝑖𝑑𝑡ℎ = 32 emblocos Grids 2D com (2x2) blocos Blocos 2D com (16x16) threads




Configuracao de execucao

A configuracao define a dimensao do problema!!!! No exemplo anterior, usando blocos

1D podemos apenas trabalhar com 𝑊𝑖𝑑𝑡ℎ = 16 !

Solucao: Manipular dimGrid e dimBlock e dividir o calculo de pedacos da matriz

resultado entre threads e blocos!




Usando blockIdx e threadIdx

Solucao: tiles!




Usando blockIdx e threadIdx: nova funcao kernel

__global__ void MatMulK(float*Md, float*Nd , float*Pd, int Width){

// linha e colunas do elemento de Pdint Row = blockIdx.y*TILE_WIDTH + threadIdx.y;int Col = blockIdx.x*TILE_WIDTH + threadIdx.x;

float Pvalue = 0;// cada thread calcula um elemento da sub -matriz no blocofor(int k = 0; k < Width; ++k)

Pvalue += Md[Row*Width+k] * Nd[k*Width+Col];

Pd[Row * Width + Col] = Pvalue;}




Usando blockIdx e threadIdx: parametros de configuracao

// configuracao para varios blocosdim3 dimBlock(Width/TILE_WIDTH , Width/TILE_WIDTH );dim3 dimGrid(TILE_WIDTH , TILE_WIDTH );

// Lancamento do KernelMatMulK <<<dimGrid , dimBlock >>>(Md, Nd, Pd, Width);




Sincronizacao

syncthreads()

primitiva chamada por umafuncao kernel

o kernel que chamou fica emespera ate que todas as threadsterminem sua execucao

7 conditionals if-then-else

7 threads em blocos diferentesnao podem sincronizar




Escalabilidade Transparente

3 Potencial para executar o mesmo codigo no hardware com umnumero diferente de recursos de execucao e escalabilidadetransparente.




Usando a memoria para obter performance

3 reduzir o trafico de dados damemoria global evitacongestionamento

3 podemos aproveitar a localidadede dados para otimizar o acessoa dados na memoria da GPU

3 Threads que usam dadoscomuns podem colaborar!

3 Solucao:tiling + shared memory









3 Threads 𝑃1,0 e 𝑃1,1 compartilham o elemento 𝑁1,0

3 Threads 𝑃0,0 e 𝑃1,0 compartilham o elemento 𝑁1,0

3 threads com elementos em comum devem estar associadas aomesmo bloco e, assim, os dados comuns podem ser guardados namemoria compartilhada!





3 reducao de trafico dedados na mem. global∝ TILE WIDTH

3 P/ NxN blocos, areducao ∝ N

3 num. de fases e Width/TILE WIDTH




Kernel MatMul com memoria compartilhada

__global__ voidMatMulK(float *Md, float *Nd, float *Pd, int Width){

__shared__ float Mds[TILE_WIDTH ][ TILE_WIDTH ];__shared__ float Nds[TILE_WIDTH ][ TILE_WIDTH ];

int bx = blockIdx.x; int by = blockIdx.y;int tx = threadIdx.x; int ty = threadIdx.y;

// Identificamos a linha e a coluna do elemento de Pdint Row = by * TILE_WIDTH + ty;int Col = bx * TILE_WIDTH + tx;

float Pvalue = 0;// Loop sobre os tiles Nd e Mdfor(int m = 0; m < Width/TILE_WIDTH; ++){

Mds[ty][tx] = Md[Row * Width + (m*TILE_WIDTH + tx)];Nds[ty][tx] = Nd[(m+TILE_WIDTH + ty)*Width + Col];_syncthreads ();

for(int k = 0; k < TILE_WIDTH; ++k)Pvalue += Mds[ty][k] * Nds[k][tx]_syncthreads ();

}Pd[Row*Width + Col] = Pvalue;

}



OpenCL

OpenCL

extensao de linguagem e API’sp/ GPU’s

applicacoes OpenCL saoportaveis para todos osprocessadores com suporte

3 sintaxe e primitivas semelhantesao CUDA

7 performance ≈ CUDA



OpenCL - arquitetura do device

OpenCL - arquitetura do device



OpenCL e CUDA

OpenCL e CUDA

OpenCL API call Explicacao equivalente em CUDA

get global id(0); ındice global do work item blockIdx.x ×blocDim.x + threadIdx.x

get local id(0); ındice local do work group threadIdx.x

get global size(0); tamanho do range ND gridDim.x ×blocDim.x

get local size(0); tamanho de cada work group blockDim.x

OpenCL conceito de paralelismo equivalente em CUDA

Kernel Kernel

programa Host programa Host

ND range (espaco de ındice) Grid

work item Thread

work group Block

__kernel void vadd(__global const float *a,__global const float *b,__global float *result){

int id = get_global_id (0);result[id] = a[id]+b[id];

}



Exemplo de codigo OpenCL: multiplicacao matricial

Exemplo de codigo OpenCL: multiplicacao matricial

#define BLOCK_SIZE 16

__kernel void

matrixMul(__global float* P, __global float* M, __global float* N, int Width)

{

int bx = get_group_id (0), by = get_group_id (1);

int tx = get_local_id (0), ty = get_local_id (1);

int mBegin = Width * BLOCK_SIZE * by;

int mEnd = aBegin + Width - 1;

int mStep = BLOCK_SIZE;

int nBegin = BLOCK_SIZE * bx;

int nStep = BLOCK_SIZE * Width;

for (int m = mBegin , n = nBegin; m <= mEnd; m += mStep , n += nStep)

{

__local float Ms[BLOCK_SIZE ][ BLOCK_SIZE ];

__local float Ns[BLOCK_SIZE ][ BLOCK_SIZE ];

Ms[ty][tx] = M[m + Width * ty + tx];

Ns[ty][tx] = N[n + Width * ty + tx];

barrier(CLK_LOCAL_MEM_FENCE );

for (int k = 0; k < BLOCK_SIZE; ++k)

Psub += Ms[ty][k] * Ns[k][tx];

barrier(CLK_LOCAL_MEM_FENCE );

}

int p = Width * BLOCK_SIZE * by + BLOCK_SIZE * bx;

P[p + Width * ty + tx] = Psub;

}



Equacao diferencial com caos


��+ 𝑥3 = sin(Ω𝑡) (Ω1,Ω2,Ω3, · · · ,Ω𝑁−1)

(𝑡1, 𝑡2, · · · , 𝑡𝑓 )𝑑𝑥

𝑑𝑡= 𝑢

𝑥0 𝑢0

integracao

Runge-Kutta 4

· · ·

𝑡0 𝑡𝑓𝑑𝑡1 𝑑𝑡

Expoente de Lyapunov

|𝛿Z(𝑡)| ≈ 𝑒𝜆𝑡|𝛿Z0|




GPU

CUDA version: v5050

CUDA Devices: 1

0: GeForce GTX 650: 3.0

Global memory: 2047mb

Shared memory: 48kb

Constant memory: 64kb

Block registers: 65536

Multiprocessors: 2

Max threads per multiprocessor:

2048

Warp size: 32

Threads per block: 1024

Max block dimensions: [ 1024,

1024, 64 ]

Max grid dimensions: [

2147483647, 65535, 65535 ]




Caos - resultados

𝑥 e 𝑢 como funcao de 𝑡




Caos - resultados

Espaco de fase 𝑢 por 𝑥




Caos - resultados

Expoente de Lyapunov




Caos - analise de desempenho











Conclusoes

3 Uso de GPU’s e altamente recomendado em aplicacoes comparalelismo de dados

3 Cada vez mais aplicacoes exigirao alta performance e fomentarao odesenvolvimento de GPU’s e de modelos de programacao paralelos

3 CUDA e um modelo de programacao inteligıvel e permite explorareficientemente o paralelismo de aplicacoes e os recursos da GPU

3 Parallel Thinking: antes de portar um codigo para rodar na GPU eimportante reconhecer quais os trechos sequenciais e os paralelos eexplorar ao maximo este usando os recursos GPU



Referencias

3 Kirk, D. ; Hwu, W.W. Programming massively parallel processors

3 www.nvidia.com/object/gpu-applications.html

3 https://developer.nvidia.com/cuda-gpus

3 http://cs.nyu.edu/courses/spring12/CSCI-GA.3033-012/



Tesla vs. Fermi - Arquiteturas CUDA 6 Curiosidades

GPU Tesla: Alto paralelismo da arquitetura

Cada GPU G200b (com arquitetura Tesla) contem 30 SMs(Streaming Multiprocessors).

Uma placa de vıdeo pode ter mais de uma GPU. Por exemplo, aGTX295 possui duas dessas GPUs.

Cada um desses SMs contem 8 ALUs (Arithmetic-Logic Units).

NVidia chama ALUs de “CUDA cores”, mas esse nome e ilusorio,pois nao se tratam de cores completos (com unidade de execucao,etc.)

Alto throughput – e possıvel requisitar para cada ALU umaoperacao logica ou aritmetica por ciclo.



GPU Tesla: Limitacoes impostas pela arquitetura

Entretanto, cada ALU e construıda na forma de um pipeline devarios estagios, apresentando muitos ciclos de latencia(da ordem de 24 ciclos).

As 8 ALUs de um SM precisam executar a mesma instrucao ao longode 4 ciclos, que e o tempo de resposta do dispatcher de instrucoes.

Por isso, cada instrucao precisa ser executada pelo menos um totalde 32 vezes. Isso e chamado de um warp.

Para superar ambas as limitacoes, o programa deve ser organizadode forma a executar em cada SM um numero de threads bem maiorque o numero de ALUs disponıveis.

O dispatcher da Tesla precisa de pelo menos 2 warps para funcionarde forma contınua, entao o mınimo e de 64 threads para naodesperdicar oportunidades de inserir operacoes nas ALUs.Recomenda-se de 128 a 192 threads. Dessas threads, apenas 8executam simultaneamente por vez. Entao, cada thread fica ociosaciclos o suficiente para compensar a latencia de cada operacao.



GPU Tesla: Instrucoes condicionais

Instrucoes subsequentes a umacomparacao podem verificar se osoperandos eram iguais (eq),diferentes (ne), um menor que ooutro (lt), etc.

A instrucao so executa se a condicaofor verdadeira, caso contrario a ALUfica ociosa por um ciclo.

Inspirado na arquitetura ARM.Porem capaz de memorizar oresultado de ate 4 comparacoes($p0–$p3) em vez de somente aultima realizada.

No ARM, evita esvaziar o pipelineem pequenas condicionais difıceis deprever.

Na GPU, evita a divergencia dewarps.



GPU Tesla: Recursos compartilhados de uma SM

Cada SM tem 16384 registradores de 32 bits que sao divididosigualmente entre as threads.

O conjunto de instrucoes e capaz de enderecar no maximo 128registradores por thread. Porem o ultimo registrador ($r127) esomente-escrita, apelidado de bitbucket.

Se alocarmos menos de 128 threads por SM, estaremosdesperdicando registradores!

Cada SM tem 16384 bytes de memoria compartilhada. Essamemoria e dividida em 16 bancos de memoria. Apenas enderecossituados em bancos diferentes podem ser acessados simultamente nomesmo ciclo.

Essa memoria e entrelacada. Cada 4 bytes (32 bits) adjacentes saocolocados em um banco diferente.Se duas threads tentarem acessar simultameamente o mesmo banco,o acesso e serializado pelo hardware e perde-se paralelismo.



GPU Tesla: Memoria global

A memoria global e compartilhada entre todas as SMs e e acessıveltambem pela CPU hospedeira.

A latencia de acesso e na faixa de 400 a 600 ciclos.

O throughput e de no maximo um acesso de 32 bits para cada SMpor ciclo.

Na arquitetura Tesla, nao existe cache. O programador deveexplicitamente copiar os dados para memorias mais locais, conformenecessario para obter desempenho.



GPU Fermi: Principais diferencas

Cache para memoria global; mais memoriacompartilhada; dobro de registradores.

Quadruplo de ALUs em cada SM.

NVidia aumentou o poder de uma SM

em vez de aumentar muito o numero de

SMs, reduzindo numero de transistores

necessarios para o mesmo pico de

GFLOPS.

Warps agora duram 2 ciclos e, devido as 2

unidades de dispatch, agora e possıvel

executar 2 warps simultaneamente em uma

SM, um em cada grupo de 16 ALUs.

Portanto, o tamanho do warp continua

sendo de 32 threads.

Melhor suporte a precisao dupla.

Porem apenas um warp por vez (metade

do throughput da precisao simples).



Unified memory - CUDA6

Unified memory - CUDA6

7 memorias da CPU e da GPU eram fisicamente distintas e separadas pelo PCI-Express bus

3 Unified memory permite a CPU o acesso direto a dados da GPU e vice-versa

3 Ha uma managed memory que torna dads acessıveis para CPU e para GPU por um unico

ponteiro



Top 500 - maquinas com CUDA

Top 500 - maquinas com CUDA



Performance em aplicacoes cientıticas - Tesla K40

Performance em aplicacoes cientıticas - Tesla K40



Papers Web of Science

Papers Web of Science



Tendencias de salario IT

Tendencias de salario IT



Tendencias de emprego IT

Tendencias de emprego IT
