Sumário

High Performance
Programming: Multicore, • GPGPU
Clusters and GPU ― Introdução
― CPU X GPU
― Aplicações
GPGPU - CUDA • Como usar
• CUDA
Professor Marcelo Trindade Rebonatto ― Principais funções da API
Curso de Ciência da Computação • Exemplos

Introdução Jogos ...

• Placas gráficas
― Jogos 3D evoluíram e passaram a exigir muito
poder computacional
• Jogos
― Além de gerar cenário 3D
― Aplicação de texturas, iluminação, sombras,

reflexões, ...
# Folhas individuais são desenhadas

# Sombras calculadas dinamicamente

3 4

GPGPU CPU x GPU

• O que é
― General Purpose Computing on GPUs
Modelo Core 2 Quad 3 Ghz GeForce 8800
― Uso de hardware gráfico para computação não-gráfica
Ano 2007 2006
• GPU (Graphics Processing Unit) Núcleos 4 128
― Tradicionalmente usada para renderização em tempo Desempenho 96 GFlops 330 GFlops
real Aritmético
― Excelente atrativo para a exploração da programação

paralela de aplicações

5 6
 CPU x GPU CPU x GPU?

Intel Core I7 980X NVIDIA GTX 480

Núcleos 6 480

Desempenho aritmético 0.1 Tflop (double) 0.5 Tflop (double)

(Máximo teórico) 1.3 Tflop (single)

Acesso Memória 25.6 GB/s 177 GB/s

7 8

CPU X GPU CPU x GPU

• CPU
― Quantidade de Caches (L1, L2, L3)
― Previsão de salto
― Alto desempenho (pelas previsões)
• GPU
― Muitas ALUs
― Memória OnBoard Rápida
― Grande quantidade de tarefas paralelas Operações aritméticas ++
Operações de memória --
9 10

GPUs Exemplos de aplicações

• Maior número de transístores para colocar em ULAs

simplificadas
• Permite a realização de um maior número de cálculos
ao mesmo tempo
• Controle de fluxo mais simples
• Aplicações paralelas onde as mesmas operações são
aplicadas sobre um grande conjunto de dados
• Por que não usar esse poder de
processamento em aplicações de propósito
geral?????
11 12
 TESLA S2070 Como Usar?

• Modelo Heterogeneos Computing

• CUDA
4 x 448 processadores
Form Factor 1U
# of Tesla GPUs 4 • OpenCL
GPU Memory Speed 1.55 GHz
GPU Memory Interface 384-bit
GPU Memory Bandwidth 148 GB/s
Double Precision floating point performance (peak) 2 Tflops
• Direct Computer
Single Precision floating point performance (peak) 4.13 Tflops
Total Dedicated Memory 12 GB GDDR5 13 14

Heterogeneos Computing
Heterogeneos Computing
Terminologia
• Host • Device
― CPU ― CPU
― Memória ― Memória (própria)

15 16

Funcionamento Funcionamento

17 18
 Funcionamento CUDA
• Compute Unified Device Architecture
― Arquitetura de computação paralela
― Criada pela NVidia
― Tirar proveito das GPUs (ou outros dispositivos)

19 20

O que é CUDA Modelo de Execução

• Arquitetura paralela de propósito geral • Execução do programa controlada pela CPU

• Objetivo: utilizar o poder computacional de GPUs ― Lança kernels executados na GPU
nVIDIA # Trechos de código executados em paralelo por múltiplas
threads
• Extensão da linguagem C que permite o uso de • Execução composta por 2 ciclos
GPUs
― CPU → GPU → … → CPU → ...→ GPU
― Suporte a uma hierarquia de threads
― Definição de kernels que são executados em GPUs
― API com funções para gerenciamento da memória e
outros tipos de controle

21 22

Grids e Thread Blocks Exemplo Grid, Bloco, Threads

• As threads são organizadas em • Kernel executando em 72 threads

blocos, dentro de grids ― Grid 2D
• Um bloco (de threads) é um # Dimensão 3 x 2 x 1
array de threads que cooperam # 6 Blocos
por:
― Memória Compartilhada
― Sincronização ― Blocos 2D
• Blocos de threads de um grid # Dimensão: 4 x 3 x1
executam de forma # 12 threads cada
independente
23 24
 Modelo de memória Modelo de Memória

• Um device possui uma área de memória

relativa ao device
• Cada Grid possui sua memória própria
― Todas as threads do grid podem acessá-la
• Cada bloco possui um espaço de memória
compartilhada
• Cada thread possui uma memória própria
― Além dos registradores de controle
© NVIDIA Corporation

25 26

Modelo de memória Memória e Threads

Device
• Execução Kernel
Block (0,0) Block (1,0)
― Grid → Blocos →
Shared Memory Shared Memory Threads
Registers Registers Registers Registers

Thread (0,0) Thread (1,0) Thread (0,0) Thread (1,0)

• Memória
Local Local Local Local
Memory Memory Memory Memory ― Registradores por thread
Global ― Memória compartilhada
Memory
por bloco
Constant
Host Memory ― Memória Global
Texture
Memory
27 28

Modelo de programação
Arquitetura Software
Threads
• Programas em Cuda • Porque programar em Threads?
utilizam CUDA Runtime
― Primitivas de alto
― Para fazer distribuição de carga entre
nível múltiplos núcleos
• É possível utilizar a API • Desafio de ter que manter o máximo de
do Driver
uso de cada núcleo
― Melhor controle da
aplicação • CUDA permite até 12 mil* threads
― CUDA é basicamente um cluster de threads

29 30
 Modelo de programação
Threads
Modelo de Programação

• Em CPU • Kernel
― Função geralmente escrita em C para CUDA
― Poucas threads (troca contexto)
― Executa no dispositivo N vezes em N threads em paralelo
― Natural gastar 1000 instruções para fazer a
• Blocos
troca de uma thread para outra
― São arranjos 1D, 2D ou 3D de threads
• Em CUDA há outro paradigma.... ― Cada thread de um Bloco possui um índice 1D, 2D ou 3D
• Não é necessário gerenciar as threads ― Organizados em grids
― Realizado em hardware • Grid é um arranjo 1D, 2D ou 3D de blocos
― Sincronismo deve ser explicito ― Cada bloco de um Grid possui um índice 1D 2D ou 3D

31 32

Identificando Threads e Blocos Limites

• Threads e Blocos possuem um ID • Cada device possui N multiprocessing units

― Uso de built-in variables • Cada multiprocessador pode executar um limite de
• Blocos e threads threads
― blockIdx.x, blockIdx.y, blockIdx.z • Exemplo: 4 MP com 768 threads cada
― threaIdx.x, threadIdx.y, threadIdx.z • Limite = 4 * 768 = 3072 threads simultaneas
• Dimensões • Threads num bloco (x, y e z)
― gridDim.x, gridDim.y, gridDim.z ― X * y * z <= 768
― blockDim.x, blockDim.y, blockDim.z • Blocos ficam alocados no mesmo multiprocessador

33 34

Modelo de Programação Invocando um kernel

• Um kernel é uma função kernel <<< #blocos, #threads >>> (argumentos)

― Começa com o especificador __global__
― Tem tipo de retorno void • Exemplos
• Kernels podem invocar funções ― nome_kernel <<< 1, 1 >>> ( void )
― Especificadas como __device__ # Cria um bloco com apenas 1 thread = å threads = 1

• Funções que executam no dispositivo ― kernel <<<2, 32>>>(void)

― Não admitem número variável de argumentos Cria 2 blocos com 32 threads cada
#

å threads = 64
#
― Não admitem recursão
• dim3 usado para especificar dimensões (x, y, z)
― Não admitem variáveis do tipo endereço de função
― Dimensões não definidas = 1
35 36
 Entendo Blocos e Threads Entendo Blocos e Threads

kernel<<< 1, 32 >>>( ) dim3 bloco (4, 8);

kernel<<< 1, bloco >>>( )

(1, 1, 1) (1, 1, 1)
(32, 1, 1) (4, 8, 1)

37 38

Interface de Programação Programa básico para CUDA

• API Cuda Runtime e API de Driver • Seleção do dispositivo a ser usado

― cudaSetDevice()
― Gerência de memória do dispositivo • Alocação de memória no dispositivo
― Transferência de dados entre host e device ― cudaMalloc()
― Gerência de sistemas com vários dispositivos • Transferência de dados entre host e dispositivo
― . . . (Manual de referência CUDA) ― cudaMemcpy()
• Liberação de memória no dispositivo
• Runtime
― cudaFree()
― Gerenciamento de Threads • Finalização
― Detecção de erros ― cudaThreadExit()

39 40

Modelo de programa
CUDA API
Dispositivo
__global__ void matrix_mul(…){ cudaError_t cudaGetDevice(int *device)
…
Definição de um Kernel /* Returns in *device the current device for the
[GPU (Parallel) code]
calling host thread. */
…
}

cudaError_t cudaSetDevice(int device)

void main(){
… /* Sets device as the current device for the
[CPU (serial) code] calling host thread. */
… Código na CPU
matrix_mul<<<dimGrid, dimBlock>>>(…)
…
cudaError_t cudaGetDeviceCount(int *count)
[CPU (serial) code]
…
/* Returns in *count the number of devices that
}
are available for execution. */
41 42
 CUDA API CUDA API
Memória Memória
cudaError_t cudaGetDeviceProperties( cudaError_t cudaMalloc(void ** devPtr,
struct cudaDeviceProp *prop,
int device) size_t size )
/* Returns in *prop the properties of device dev. */ /* Allocates size bytes of linear memory on
the device and returns in *devPtr a pointer
struct cudaDeviceProp { int maxGridSize[3];
to the allocated memory. The memory is not
char name[256]; int multiProcessorCount;
cleared. */
size_t totalGlobalMem; int computeMode;
size_t int concurrentKernels;
sharedMemPerBlock;
int regsPerBlock;
int ECCEnabled; cudaError_t cudaFree(void *devPtr)
int pciBusID;
int warpSize;
int
/* Frees the memory space pointed to by
int maxThreadsPerBlock; maxThreadsPerMultiProcessor; devPtr, which must have been returned by a
int maxThreadsDim[3]; } previous call to cudaMalloc(). */
43 44

CUDA API Programa básico para CUDA

Memória
cudaError_t cudaMemcpy( void * dst, • Select Device to use
const void *src, • Allocate host memory(malloc) for Array(s)
size_t count,
enum cudaMemcpyKind • Initialize host Array(s)
kind) • Allocate device memory(cudaMalloc) for Array(s)
/* Copies count bytes from the memory area • Transfer data from host to device memory(cudaMemCpy)
pointed to by src to the memory area • Specify kernel execution Configuaration
pointed to by dst.*/ ― This is very important. Depending upon it blocks and
kind: cudaMemcpyHostToHost threads automatically get assigned numbers
cudaMemcpyHostToDevice • Call Kernel
cudaMemcpyDeviceToHost • Transfer result from device to host memory (cudaMemCpy)
cudaMemcpyDeviceToDevice • Deallocate host(free) and device(cudaFree) memories
45 46

Referências

Site NVIDEA, Cuda Zone

[Link]/~castelo/CUDA/[Link]
[Link]
[Link]
[Link]

Video com Uso de CUDA

[Link]
ipelines.mp4

47