GPU, CUDA, OpenCL and OpenACC for Parallel Applications

GPU, CUDA e OpenCL for parallel Applications
Ms. Eng. Marcos Amar´ıs Gonz´alez
Dr. Alfredo Goldman vel Lejbman
University of S˜ao Paulo
Institute of Mathematics an Statistics
Department of Science Computing
April, 2014
(gold, amaris)@ime.usp.br (IME - USP) GPU, CUDA e OpenCL April, 2014 0 / 52

Timeline
1 GPUs e GPGPU
2 CUDA
Proﬁling e Optimiza¸c˜oes
3 OpenCL e OpenACC

GPUs
1 GPUs e GPGPU
2 CUDA
3 OpenCL e OpenACC

GPUs
Introdu¸cão
80’ primeiro controlador de v´ıdeo.
Evolu¸cão dos jogos 3D.
Maior poder computacional.
Além de gerar o cenário 3D, é preciso
aplicar texturas, ilumina¸cão, som-
bras, reflexões, etc.
Para tal, as placas gráficas pas-
saram a ser cada vez mais flex´ıveis
e poderosas

GPUs
Introdu¸cão
O termo GPU foi popularizado pela Nvidia em
1999, que inventou a GeForce 256 como a
primeira GPU no mundo.
No 2002 fou lan¸cada a primeira GPU para
propósito geral.
O termo GPGPU foi cunhado por Mark Harris.
Os principais fabricantes de GPUs são a NVIDIA
e a AMD.
2005 NVIDIA lan¸cou CUDA, 2008 grupo
Khronos lan¸cou OpenCL, 2011 foi anunciado
OpenACC.

GPUs
GPU de Propósito Geral
GPGPU (GPU de Propósito Geral).
GPGPU podem atuar em conjunto com CPUs Intel ou AMD.
Paralelismo do tipo SIMD.
Programa principal executa na CPU (host) e é o responsável por iniciar as threads
na GPU (device).
Tem sua própria hierarquia de memória e os dados devem ser transferidos através
de um barramento PCI Express.

GPUs
Lei de Amdahl e Taxonomia de Flynn
Lei de Amdahl - 1967
A Lei de Amdahl é a lei que governa o speedup na utiliza¸cão de proces-
sadores paralelos em rela¸cão ao uso de apenas um processador.
Speedup:
S = Speed-up
P = Number of Processors
T = Time
Sp =
T1
Tp
(1)
Taxonomia de Flynn - 1966
Single Instruction Multiple Instruction
Single Data SISD - Sequential MISD
Multiple Data SIMD [SIMT] - GPU MIMD - Multicore

GPUs
Medidas de desempenho: FLOPS
FLOPS - FLoating-point Operations Per Second
Opera¸cões de ponto flutuante por segundo
Forma simples: medir o número de instru¸cões por unidade de tempo.
No caso, instru¸cões de ponto flutuante.
Linpack (sistemas densos de equa¸cões lineares)
Melhor rela¸cão desempenho/custo (Gflops/ $$);
Melhor rela¸cão desempenho/consumo (Gflops/Watts);
Melhor rela¸cão desempenho/volume (Gflops/Volume).

GPUs
GPU Versus CPU
Hoje em dia elas são capaz de realizar a computa¸cão paralela mais eficiente
do que CPUs multicore.

GPUs
Top 500 Supercomputers
Intel Core i7 990X: 6 núcleos, US$ 1000 Desempenho teórico máximo de 0.4 TFLOP
GTX680: 1500 núcleos e 2GB, pre¸co US$500 Desempenho teórico máximo de 3.0 TFLOP
Aceleradores e co-processadores no ranking dos 500 Supercomputadores mais rápidos do mundo.

GPUs
Top 500 Green Supercomputers $$$$$$
Ranking dos supercomputadores mais eﬁcientes no mundo em termos de
energia.

GPUs
RoadMap de Arquiteturas de GPU NVIDIA
Em GPUs modernas o consumo de energia é uma restri¸cão importante.
Projetos de GPU são geralmente altamente escalável.

GPUs
RoadMap de Arquiteturas de GPU NVIDIA
Compute Capability é uma diferencia¸cão entre arquiteturas e modelos de
GPUs da NVIDIA, para certas especifica¸cões de hardware e software em
elas.

GPUs
Arquitetura de uma GPU NVIDIA

GPUs
Arquitetura Tesla
Arquitetura Tesla, ela é umas das primeiras com suporte a CUDA. Seu mul-
tiprocessador tem 8 processadores e até 16 Kb de memória compartilhada.

GPUs
Arquitetura Fermi

GPUs
Escalonador de Warps
Ocultar latˆencia
1 escalonador de Warps em arquiteturas Tesla, 2 na arquitetura Fermi e
Kepler tem 4 escalonadores de Warp.

GPUs
Arquitetura Kepler
Arquitetura Kepler tem um multiprocessador (SMX) de 192 processadores,
32 SFU, 32 unidades de escrita e leitura, arquivo de 64 kb de registradores,
etc...

CUDA
1 GPUs e GPGPU
2 CUDA
3 OpenCL e OpenACC

CUDA
Compute Unified Device Architecture
CUDA - Compute Unified Device Architecture
CUDA é uma linguagem proprietária para programa¸cão em GPUs desenvolvida pela
NVIDIA.
O CUDA está na versão 7.0 atualmente, avan¸ca também segundo o Compute
Capability das GPUs.
Ela é uma extensão da linguagem C, e permite controlar a execu¸cão de threads na
GPU e gerenciar sua memória.
Ambiente CUDA
CUDA Driver
CUDA Toolkit
CUDA SKD

CUDA
Compute Capability
As GPUs com compute capability 3.5 podem fazer uso do paralelismo
dinˆamico e Hyper-Q.

CUDA
Paralelismo dinˆamico
Permite que segmentos de processamento da GPU gerem dinamicamente
novos segmentos, possibilitando que a GPU se adapte de modo dinˆamico
aos dados

CUDA
Hyper-Q
Permite que até 32 processos MPI sejam lan¸cados simultaneamente em uma
GPU. O Hyper-Q é ideal para aplica¸cões de cluster que usam MPI.

CUDA
Fun¸cões Kernel
Threads executam o código definido em uma fun¸cão kernel. A chamada da fun¸cão
kernel, dispara a execu¸cão de N instâncias paralelas para N threads.
Executado sobre Só é chamada desde
device float deviceFunction() device device.
global float KernelFunction() device device/host.
host float hostFunction() host host.

CUDA
Fun¸cões de Transferência e Manipula¸cão de dados
Se declaram e alocam as variáveis no host e no device.
* cudaMalloc(void **pointer, size_t nbytes)
* cudaMemcpy(void *dst, const void *src, size_t count, enum
cudaMemcpyKind kind)
* cudaFree(void *pointer)
cudaMemcpyKind
cudaMemcpyHostToHost Host ⇒ Host
cudaMemcpyHostToDevice Host ⇒ Device
cudaMemcpyDeviceToHost Device ⇒ Host
cudaMemcpyDeviceToDevice Device ⇒ Device

CUDA
Exemplo: Soma de Vetores
Processo de Transferência de dados para a GPU, execu¸cão do kernel da
Soma de Vetores e transferência da solu¸cão de cada bloco de threads.
// allocate the memory on the GPU
cudaMalloc( (void **)& dev_a , N*sizeof(float ));
cudaMalloc( (void **)& dev_b , N*sizeof(float ));
cudaMalloc( (void **)& dev_partial_c , GridSize*sizeof(float ));
// copy the arrays ’a’ and ’b’ to the GPU
cudaMemcpy( dev_a , host_a , N*sizeof(float), cudaMemcpyHostToDevice );
cudaMemcpy( dev_b , host_b , N*sizeof(float), cudaMemcpyHostToDevice );
VecAdd <<<GridSize ,BlockSize >>>( dev_a , dev_b , dev_partial_c , N );
// copy the array ’c’ back from the GPU to the CPU
cudaMemcpy( host_partial_c , dev_partial_c , GridSize*sizeof(float),
cudaMemcpyDeviceToHost );

CUDA
Organiza¸cão da Execu¸cão
Uma chamada a uma fun¸cão kernel cria um grid de blocos de threads, as
quais executam o código.
Um SM, SMX ou SMM executa um ou mais blocos de threads e os cores e
outras unidades de execu¸cão no multiprocessador executam as instru¸cões de
cada thread.
Variáveis e palavras reservadas são identificadores de cada thread dentro de
um bloco.

CUDA
Modelo de Programa¸cão
Organizado em grids, blocos e threads. Threads são agrupadas em blocos e estes
são agrupados em um grid.
Tradu¸cão para endere¸camento linear para saber o id de uma thread em um grid.
Para saber a posi¸cão de uma thread usamos as palavras reservadas:

CUDA
Espa¸cos de Mem´oria sobre um dispositivo CUDA

CUDA
Espa¸cos de Memória sobre um dispositivo CUDA
A latência de acesso a memória global é 100x que da memória com-
partilhada.
Tem palavras reservadas para a declara¸cão das variaveis a ser alocadas
em cada n´ıvel de memória.

CUDA
./devicequery
Device 0: "GeForce GTX 295"
CUDA Driver Version / Runtime Version 6.0 / 5.5
CUDA Capability Major/Minor version number: 1.3
Total amount of global memory: 896 MBytes (939327488 bytes)
(30) Multiprocessors, ( 8) CUDA Cores/MP: 240 CUDA Cores
GPU Clock rate: 1242 MHz (1.24 GHz)
Memory Clock rate: 1000 Mhz
Memory Bus Width: 448-bit
Maximum Texture Dimension Size (x,y,z) 1D=(8192), 2D=(65536, 32768), 3D=(2048, 2048, 2048)
Maximum Layered 1D Texture Size, (num) layers 1D=(8192), 512 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(8192, 8192), 512 layers
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 16384
Warp size: 32
Maximum number of threads per multiprocessor: 1024
Maximum number of threads per block: 512
Max dimension size of a thread block (x,y,z): (512, 512, 64)
Max dimension size of a grid size (x,y,z): (65535, 65535, 1)
Maximum memory pitch: 2147483647 bytes
Texture alignment: 256 bytes
Concurrent copy and kernel execution: Yes with 1 copy engine(s)
Run time limit on kernels: No
Integrated GPU sharing Host Memory: No
Support host page-locked memory mapping: Yes
Alignment requirement for Surfaces: Yes
Device has ECC support: Disabled
Device supports Unified Addressing (UVA): No
Device PCI Bus ID / PCI location ID: 4 / 0
Compute Mode:
< Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

CUDA
./devicequery
Device 1: "GeForce GT 630"
( 2) Multiprocessors, ( 48) CUDA Cores/MP: 96 CUDA Cores
L2 Cache Size: 131072 bytes
Warp size: 32
Device has ECC support: Disabled
Device supports Unified Addressing (UVA): Yes
Compute Mode:
> Peer access from GeForce GTX 660 (GPU0) -> GeForce GT 630 (GPU1) : No
> Peer access from GeForce GT 630 (GPU1) -> GeForce GTX 660 (GPU0) : No

CUDA
./deviceQuery Starting...
Device 0: "Tesla K40c"
(15) Multiprocessors, (192) CUDA Cores/MP: 2880 CUDA Cores
L2 Cache Size: 1572864 bytes
Warp size: 32
Device has ECC support: Enabled
Device supports Unified Addressing (UVA): Yes
Compute Mode:
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 6.5, CUDA Runtime Version = 5.5, NumDevs = 1,
Device0 = Tesla K40c Result = PASS

CUDA
Multiplica¸cão de Matrizes em CUDA - I
Esquema de paraleliza¸cão e kernel da multiplica¸cão de matrizes por padrão
com CUDA.
__global__ void matMul(float* Pd , float* Md ,
float* Nd , int N) {
float Pvalue = 0.0;
int j = blockIdx.x * tWidth + threadIdx.x;
int i = blockIdx.y * tWidth + threadIdx.y;
for (int k = 0; k < N; ++k)
Pvalue += Md[j * N + k] * Nd[k * N + i];
Pd[j * N + i] = Pvalue;
}

CUDA
Multiplica¸cão de Matrizes em CUDA - II
Esquema da Multiplica¸cão de Matrizes usando memória compartilhada em
CUDA:

CUDA
Kernel da Multiplica¸c˜ao de Matrizes usando mem´oria compartilhada em
CUDA:
__global__ void matMul(float* Pd , float* Md ,
float* Nd , int N){
__shared__ float Mds[tWidth ][ tWidth ];
__shared__ float Nds[tWidth ][ tWidth ];
int tx = threadIdx.x;
int ty = threadIdx.y;
int Col = blockIdx.x * tWidth + tx;
int Row = blockIdx.y * tWidth + ty;
float Pvalue = 0;
for (int m = 0; m < N/tWidth; ++m) {
Mds[ty][tx] = Md[Row*N + (m*tWidth + tx)];
Nds[ty][tx] = Nd[Col + (m*tWidth + ty)*N];
__syncthreads ();
for (int k = 0; k < Tile_Width; ++k)
Pvalue += Mds[ty][k] * Nds[k][tx];
__syncthreads ();
}
Pd[Row * N + Col] = Pvalue;
}

CUDA Profiling e Optimiza¸cões
Ferramentas de Profiling
Ferramentas de Profiling oferecidas pela NVIDIA
Figura : Profiling Tools provided by NVIDIA

NVIDIA Visual Proﬁling
Figura : Proﬁle Discrete Cosine Transform

Acessos Agrupados a Memória
A partir de GPUs de CC superior a 1.2.
Acessos agrupados realmente melhora o desempenho da aplica¸cão.
Se o endere¸co de base de um bloco é n, então qualquer thread i dentro desse bloco deve acessar
o endere¸co: (n + i) ∗ typeOfRead.
Acessos Agrupados Acessos No Agrupados

Conflito de Bancos na Memória Compartilhada
A memória compartilhada é dividida em módulos (também chamados de
bancos). Se duas posi¸cões de memória ocorrem no mesmo banco, então
temos um conflito de banco.

Figura : Achieved Occupancy metric in GTX-Titan

Figura : Global Load Transactions metric in GTX-Titan

Figura : Global Store Transactions metric in GTX-Titan

OpenCL
1 GPUs e GPGPU
2 CUDA
3 OpenCL e OpenACC

OpenCL
Open Computing Language
Modelo de plataforma
A linguagem serve como uma camada de abstra¸cão ao hardware heterogêneo.
É composto por um host e um ou mais dispositivos OpenCL (OpenCL devices).
Cada dispositivo possui uma ou mais unidades de computa¸cão (compute units).
Estes são compostos por um conjunto de elementos de processamento (processing ele-
ments).

OpenCL
Open Computing Language
Aplica¸cão e Fun¸cões Kernel em OpenCL
Uma aplica¸cão OpenCL deve seguir os seguintes passos:
1 Descobrir os componentes heterogêneos;
2 Detectar suas caracter´ısticas;
3 Criar os blocos de instru¸cões (kernels) que irão executar na plataforma
heterogênea;
4 Iniciar e manipular objetos de memória;
5 Executar os kernels na ordem correta e nos dispositivos adequados pre-
sentes no sistema;
6 Coletar os resultados finais.
( 

OpenCL
Contexto
Define o ambiente de execu¸cão no qual os kernels são definidos e executam.
Assim, um contexto é todo o conjunto de recursos que um kernel vai
utilizar durante sua execu¸cão.
// Get platform and device information
cl_platform_id platform_id = NULL;
cl_device_id device_id = NULL;
cl_uint ret_num_devices ;
cl_uint ret_num_platforms ;
cl_int ret = clGetPlatformIDs (1, &platform_id , & ret_num_platforms );
ret = clGetDeviceIDs ( platform_id , CL_DEVICE_TYPE_DEFAULT , 1,
&device_id , & ret_num_devices );
// Create an OpenCL context
cl_context context = clCreateContext ( NULL , 1, &device_id , NULL , NULL ,

OpenCL
Fila de Comandos
Os comandos são colocados nesta fila e aguardam seu momento de executar.
Esta fila aceita três tipos de comandos:
1 Execu¸cão de kernel,
2 Transferência de dados (objetos de memória)
3 Sincroniza¸cão, sé é necessária.
// Create a command queue
cl_command_queue command_queue = clCreateCommandQueue (context , device_id , 0, &ret );
// Create memory buffers on the device for each vector
cl_mem a_mem_obj = clCreateBuffer (context , CL_MEM_READ_ONLY ,
LIST_SIZE * sizeof(int), NULL , &ret );
cl_mem b_mem_obj = clCreateBuffer (context , CL_MEM_READ_ONLY ,
cl_mem c_mem_obj = clCreateBuffer (context , CL_MEM_WRITE_ONLY ,
// Copy the lists A and B to their respective memory buffers
ret = clEnqueueWriteBuffer (command_queue , a_mem_obj , CL_TRUE , 0,
LIST_SIZE * sizeof(int), A, 0, NULL , NULL );
ret = clEnqueueWriteBuffer (command_queue , b_mem_obj , CL_TRUE , 0,
LIST_SIZE * sizeof(int), B, 0, NULL , NULL );

OpenCL
Execute the OpenCL kernel
// Create a program from the kernel source
cl_program program = clCreateProgramWithSource (context , 1,
(const char **)& source_str , (const size_t *)& source_size , &ret );
// Build the program
ret = clBuildProgram (program , 1, &device_id , NULL , NULL , NULL );
// Create the OpenCL kernel
cl_kernel kernel = clCreateKernel (program , "vector_add", &ret );
// Set the arguments of the kernel
ret = clSetKernelArg (kernel , 0, sizeof(cl_mem), (void *)& a_mem_obj );
ret = clSetKernelArg (kernel , 1, sizeof(cl_mem), (void *)& b_mem_obj );
ret = clSetKernelArg (kernel , 2, sizeof(cl_mem), (void *)& c_mem_obj );
// Execute the OpenCL kernel on the list
size_t global_item_size = LIST_SIZE; // Process the entire lists
size_t local_item_size = 64; // Divide work items into groups of 64
ret = clEnqueueNDRangeKernel (command_queue , kernel , 1, NULL ,
&global_item_size , &local_item_size , 0, NULL , NULL );
Kernel de Soma de Vetores.

OpenCL
OpenCL - Tipos de Execu¸cão de Kernels
2 tipos de execu¸cão: Dara Parallel e Task
Parallel. A hierarquia de execu¸cão de
OpenCL é também parecida que em CUDA.
N-Dimensional Range
CUDA OpenCL
grid NDRange
block threads work group
thread work item

OpenCL
OpenCL - Modelo de Memória
Parecido que em CUDA, em OpenCL existem 4 locais diferentes para a
memória que é enviada para o device:

OpenCL
CUDA - OpenCL
Semelhan¸cas
O host inicia o ambiente de execu¸cão na GPU.
As threads são identificadas por ´ındices.
As threads são agrupadas.
O host aloca e preenche dados na memória do device
A execu¸cão dos kernels pode ser s´ıncrona ou ass´ıncrona.
Existem 4 diferentes tipos de memória no device: Global, constante, local
(shared), private.
Diferen¸cas
No OpenCL existem 2 tipos de execu¸cão diferentes:
1 Data Parallel
2 Task Parallel
O CUDA implementa só o modelo SIMT(gold, amaris)@ime.usp.br (IME - USP) GPU, CUDA e OpenCL April, 2014 47 / 52

OpenCL
OpenACC
Anunciado em novembro de 2011 na conferência SuperComputing.
É um padrão para programa¸cão paralela.
O padrão tem como base o compilador PGI (Portland Group)
Cole¸cão de diretivas para especificar la¸cos e regiões de código paralelizáveis
em aceleradores.

OpenCL
Modelo de execu¸cão de OpenACC
O modelo de execu¸cão do OpenACC tem três n´ıveis: gang, worker e vector.
Em GPU pode ser mapeado como:
gang → bloco de threads
worker → warp
vector → threads em um warp
As Diretivas em C/C++ são especificadas usando #pragma.
Se o compilador não utilizar pré-processamento, as anota¸cões são ignoradas
na compila¸cão.

OpenCL
Exemplo

OpenCL
Compila¸c˜ao com PGI usando acc

OpenCL
S´o isso... Obrigado.
O EP 2 sobre GPUs, deve estar pronto para a sexta 17 de abril
com data de entrega 1 de maio!

GPU, CUDA, OpenCL and OpenACC for Parallel Applications

Mais conteúdo relacionado

Destaque

Semelhante a GPU, CUDA, OpenCL and OpenACC for Parallel Applications

Mais de Marcos Gonzalez

GPU, CUDA, OpenCL and OpenACC for Parallel Applications