GPGPU

1. GPGPU

2. GPU

3. GPU – Quebra de paradigma radical • Você mudaria todo o teu paradigma de desenvolvimento e hardware apenas para ganhar 5% de desempenho?

4. GPGPU – a evolução ... • Há 3 fases históricas das GPUs: • Fixed Function GPU • Programmable GPU • Arquitetura Unificada

5. Fixed Function GPUs • Arquitetura incapaz de aceitar programação • Impossível realizar cálculos sem ser de computação gráfica • Incapacidade de acesso ao processador • Conceito de APIs

6. Back Buffer Front Buffer Fixed Function GPUs Processador(es) Engine de Geometria Engines de Rasterização d e v í d e o M e m ó r i a Front Buffer Back Buffers Interface CPU - GPU Z Buffer Stencil Buffer CPU Texture Buffer Interface GPU - Video

7. Programmable GPU • Vertex and Pixel Shaders • Arquitetura orientada a estrutura de dados de computação gráfica

8. Programmable GPU

9. Programmable GPU

10. FERMI – GF100

11. FERMI – GF100

12. Arquitetura Unificada - CUDA

13. Programmable GPU • Limitações: • Shaders • Modo de endereçamento limitado • Pequeno grupo de instruções • Dificuldade de comunicação entre processadores e processos

14. GPGPU • Problema de ter que mapear tudo para APIs • Usar funções OpenGL ou DirectX para efetuar todo tipo de operações

15. Arquitetura Unificada - CUDA

16. Arquitetura Unificada - CUDA • Paralelismo sem esforço, baixo custo...

17. Por que mais rápido? Tarefa 100 vezes mais rápido • Tarefa de 1 ano de duração cai para 3 dias • Tarefa de 1 dia cai para 15 minutos • Tarefa de 3 segundos cai para 30 vezes por segundo

18. Exemplo: Simulação de Multidão

19. GPU • Operações aritiméticas ++ • Operações de memória --

20. Threads • Porque programar em Threads? • Para fazer distribuição de carga em arquitetura single core • Para fazer distribuição de carga entre múltiplos núcleos • Desafio de ter que manter o máximo de uso de cada núcleo

21. Threads • Quantas threads voce já criou para seu programa?

22. Threads • CUDA permite até 12 mil threads • CUDA é basicamente um cluster de threads

23. Threads – Custo de gerenciamento • Em CPU, como fazemos pouca troca de threads, podemos achar natural gastar 1000 instruções para fazer a troca de uma thread para outra. Em CUDA há outro paradigma.... • Não é necessário gerenciar as threads, a priori. • Sincronismo deve ser explicito

24. Modelo de Programação • CUDA estende a linguagem C através de kernels • Kernels equivalem a funções, que serão executadas N vezes em paralelo • N é o número de threads

25. Arquitetura do Software

26. Funções em CUDA • Não pode haver recursão no __device__ • Sem variáveis estáticas • Sem numero variável de parâmetros • A chamada do kernel é assíncrona • Sincronismo deve ser feito explicitamente • __device__ __host__ : podem ser usados juntos

27. Limite de banda de memória Importância do reuso de dados

28. Threads, Blocos e Grids Um kernel corresponde a um grid de Blocos Cada Bloco é composto por threads Todas as threads compartilham a mesma área de memória As threads de um mesmo bloco podem compartilhar umas com as outras Threads de blocos diferentes não podem compartilhar memória entre si

29. Threads, Blocos e Grids - memórias

30. Palavras reservadas

31. Hierarquia de Threads • Todos os threads de um bloco usam da mesma memória compartilhada. • O número de threads num bloco é limitado pela memória: GPUs atuais possuem até 512 threads.

32. Funções em CUDA • __ global__ void KernelFunction (...) • dim3 DimGrid (100, 10); // Grid com 1000 blocos • dim3 DimBlock (4, 8, 8); // Cada bloco tem 256 threads • Size_t SharedMemBytes = 32 • KernelFun << DimGrid, DimBlock, SharedMemBytes>> (...);

33. Kernel – será compilado para a GPU // Kernel definition __global__ void vecAdd(float* A, float* B, float* C) { } int main() { // Kernel invocation vecAdd<<<1, N>>>(A, B, C); } __global  define que é um kernel

34. kernel __global__ void vecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { vecAdd<<<1, N>>>(A, B, C); } threadIdx  define um ID de um dos threads << n, m>>  numero de blocos (n) e threads (m) solicitados para o kernel

35. Hierarquia de Threads • threadIdx é um vetor de 3 componentes • Threads podem ser identificados com índices de 1, 2 ou 3 dimensões (formando thread blocks de uma, duas ou três dimensões) • Índice de uma thread: • Se for um bloco 1D: é a mesma coisa • Se for um bloco 2D (Dx, Dy): threadId de um thread de índice (x, y) é x + yDx • Se for um bloco 3D (Dx, Dy, Dz): threadId de uma thread de índice (x, y, z) é x + yDx + zDxDy

36. Primeiro exemplo de programa • cudaMalloc  aloca espaço de memória global da GPU

37. Primeiro exemplo de programa

38. Exercicio – Julia set (Fractais em GPU) Idéia do Julia Set: iterar uma equação para pontos no espaço dos numeros complexos. Um ponto não pertence ao conjunto, se depois de infinitas iterações o valor cresce para o infinito. Caso contrário, ele está preso a uma região. A Equação de iteração será dada por Z(n+1) = Z(n)*Z(n) + C

39. Versao em CPU int main (void) { CPUBitmap bitmap (DIM, DIM); unsigned char *ptr = bitmap.get_ptr(); kernel (ptr); bitmap.display_and_exit (); }

40. Versao em CPU (2) void kernel (unsigned char *ptr) { for (int y = 0; y<DIM; y++) { for (int x=0; x<DIM; x++){ int offset = x + y * DIM; int juliaValue = julia (x, y); ptr [offset*3 + 0] = 255 * juliaValue; ptr [offset*3 + 1] = 255 * juliaValue; ptr [offset*3 + 1] = 255 * juliaValue; } } }

41. Versao em CPU (3) int julia (int x, int y) { cont float scale = 1.5; float jx = scale * (float) (DIM/2 – x)/(DIM/2); float jy = scale * (float) (DIM/2 – y)/(DIM/2); // assumimos que o complexo C = -0.6, 0.121i // assumimos que o complexo a = jx, jy i ; int i = 0; for (i=0; i<200; i++) { a = a * a + c; // multiplicacao de complexos: a.real*C.real – a.i*C.i, a.i*C.real + a.real*C.i // soma de complexos: a.real + C.real, a.i +C.i if (a.magnitude() > 1000) // magnitude = a.r*a.r + a.i*a.i return 0; } return 1; }

42. Versao em GPU (1) Int main (void) { CPUBitmap bitmap (DIM, DIM); unsigned char *dev_bitmap; cudaMalloc ( (void**)&dev_bitmap, bitmap.image_size() ) ); dim3 grid (DIM, DIM); kernel <<<grid, 1>>> (dev_bitmap); cudaMemcpy (bitmap.get_ptr(), dev_bitmapo, bitmap.image_size(), cudaMemcpyDeviceToHost ) ); bitmap.display_and_exit(); cudaFree (dev_bitmap); }

43. Versao em GPU (2) __global__ void kernel (unsigned char *ptr) { int x = blockIdx.x int y = blockIdx.y; int offset = x+y*gridDim.x int juliaValue = julia (x, y); ptr[offset*3 +0] = 255 *juliaValue; ptr[offset*3 +1] = 255 *juliaValue; ptr[offset*3 +2] = 255 *juloaValue; cudaMalloc ( (void**)&dev_bitmap, bitmap.image_size() ) ); }

44. Versao em GPU (3) __device__ int julia (int x, int y) { cont float scale = 1.5; float jx = scale * (float) (DIM/2 – x)/(DIM/2); float jy = scale * (float) (DIM/2 – y)/(DIM/2); // assumimos que o complexo C = -0.6, 0.121i // assumimos que o complexo a = jx, jy i ; int i = 0; for (i=0; i<200; i++) { a = a * a + c; // multiplicacao de complexos: a.real*C.real – a.i*C.i, a.i*C.real + a.real*C.i // soma de complexos: a.real + C.real, a.i +C.i if (a.magnitude() > 1000) // magnitude = a.r*a.r + a.i*a.i return 0; } return 1; }