Programação Paralela em OpenCL

1. ProgramaçãoParalelaemOpenCL César L. B. Silveira cesar@v3d.com.br 19 de novembro de 2010

2. Introdução • Open Computing Language (OpenCL) • Padrãoaberto, livre de royalties, paraprogramaçãoparalelaemambientesheterogêneos • Ambientesheterogêneos • CPUs multicore, GPUs many-core • Framework • Arquitetura • Linguagem • API

3. Introdução • Neste tutorial... • Conceitos da arquitetura OpenCL • Breve visão da linguagem OpenCL C • APIs fundamentais para o desenvolvimento • Exemplo completo de aplicação • Integração com OpenGL

4. Hello World • Códigosequencial void ArrayDiff(constint* a, const int* b, int* c, intn) { for (inti = 0; i < n; ++i) { c[i] = a[i] - b[i]; } }

5. Hello World • CódigoOpenCL (kernel) __kernel void ArrayDiff(__global const int* a, __global const int* b, __global int* c) { int id = get_global_id(0); c[id] = a[id] - b[id]; }

6. Aplicações OpenCL Camadas de uma aplicação OpenCL

7. ArquiteturaOpenCL • Arquiteturaabstrata de baixonível • Implementaçõesmapeiamparaentidadesfísicas • Quatromodelos • Plataforma • Execução • Programação • Memória

8. ArquiteturaOpenCL • Modelo de plataforma • Descreveentidades do ambienteOpenCL

9. Arquitetura OpenCL Hospedeiro Dispositivo Dispositivo Unidade de Computação Unidade de Computação ... Elemento de Processamento Elemento de Processamento Elemento de Processamento Elemento de Processamento ... ... ... Modelo de Plataforma

10. Arquitetura OpenCL • Exemplo: GPU NVIDIA • GPU: dispositivo • Streaming Multiprocessor: unidade de computação • Processor/Core: elemento de processamento

11. ArquiteturaOpenCL • Modelo de execução • Descreveinstanciaçãoeidentificação de kernels

12. Arquitetura OpenCL • Modelo de execução • Espaços de índices 1D 2D 3D

13. ArquiteturaOpenCL Item de trabalho (work-item) Grupo de trabalho (work-group) Modelo de execução: itens e grupos de trabalho

14. Arquitetura OpenCL Instância Item de trabalho __kernelvoid f(...) { ... Kernel

15. ArquiteturaOpenCL 0 1 0 Grupo de trabalho: (1, 0) 1 Modelo de execução: identificadores

16. ArquiteturaOpenCL 2 0 1 3 0 Item de trabalho: (1, 2) 1 2 Identificador global 3 Modelo de execução: identificadores

17. ArquiteturaOpenCL Item de trabalho: (1, 0) 0 1 0 Identificador local 1 Modelo de execução: identificadores

18. Arquitetura OpenCL Contexto mem Filas de comandos Dispositivos Kernels Objetos de programa Objetos de memória Modelo de execução: objetos

19. ArquiteturaOpenCL • Modelo de execução: objetos • Toda comunicação com um dispositivoéfeitaatravésdasuafila de comandos • Objetos de memória • Buffers: acessodireto via ponteiros • Images: acesso especial via samplers

20. ArquiteturaOpenCL • Modelo de programação • Paralelismo de dados (data parallel): múltiplositens de trabalhopara um kernel • Paralelismo de tarefas (task parallel): um item de trabalhopara um kernel

21. Arquitetura OpenCL • Modelo de memória • Descreve uma hierarquia de memória e o compartilhamento de cada nível entre itens e grupos de trabalho

22. Arquitetura OpenCL Memória privada - Exclusiva de cada item de trabalho - Leitura e Escrita Modelo de memória

23. Arquitetura OpenCL Memória local - Compartilhada pelo grupo de trabalho - Leitura e Escrita Modelo de memória

24. Arquitetura OpenCL Memória global - Compartilhada por todos os itens de trabalho - Leitura e Escrita Modelo de memória

25. Arquitetura OpenCL Memória constante - Compartilhada por todos os itens de trabalho - Somente leitura Modelo de memória

26. ArquiteturaOpenCL • Modelo de memória: consistência • Um item de trabalholêcorretamenteoqueoutrosescrevem?

27. Arquitetura OpenCL - Memória é consistente para um item de trabalho x x Item de trabalho Memória Modelo de memória: consistência

28. Arquitetura OpenCL - Memória é consistente para um grupo de trabalho após uma barreira (barrier) x ? ? y Itens de trabalho Memória Modelo de memória: consistência

29. Arquitetura OpenCL - Memória é consistente para um grupo de trabalho após uma barreira (barrier) x y Itens de trabalho Memória Modelo de memória: consistência

30. Arquitetura OpenCL - Memória é consistente para um grupo de trabalho após uma barreira (barrier) Itens de trabalho Memória barrier() Modelo de memória: consistência

31. Arquitetura OpenCL - Memória é consistente para um grupo de trabalho após uma barreira (barrier) y x Itens de trabalho Memória Modelo de memória: consistência

32. Arquitetura OpenCL • Sincronização • Barreira: execução só prossegue após todos os itens de trabalho de um mesmo grupo de trabalho a terem atingido • Não há sincronização entre grupos de trabalho

33. LinguagemOpenCL C • Utilizadapara a escrita de kernels • Baseada no padrão C99 • Acrescentaextensõeserestrições

34. LinguagemOpenCL C • Extensão: tiposvetoriais • Notação: tipo[# de componentes] • 1, 2, 4, 8 ou 16 componentes • Exemplos: float4 int8 short2 uchar16

35. LinguagemOpenCL C • Operações com tiposvetoriais • Entre vetores de mesmonúmero de componentes • Entre vetoreseescalaresfloat4 v = (float4)(1.0, 2.0, 3.0, 4.0); float4 u = (float4)(1.0, 1.0, 1.0, 1.0);float4 v2 = v * 2;float4 t = v + u;

36. LinguagemOpenCL C • Definição de kernels • Qualificador__kernel • Todokerneldeve ser void__kernel void f(…){ ...}

37. LinguagemOpenCL C • Qualificadores de espaço de endereçamento • Definemníveldamemóriaapontadapor um ponteiro • __global, __local, __privatee__const • Default: __private

38. LinguagemOpenCL C • Restrições • Ponteirosparafunçãonãosãosuportados • Funçõese macros com númerovariável de argumentosnãosãosuportados • Qualificadoresextern, staticeautonãosãosuportados • Nãohásuporte a recursão

39. API de suporte: kernels • Funções de identificação • Informaçõessobreespaço de índices, item egrupo de trabalhoget_global_id(uintdimindx)get_local_id(uintdimindx)get_group_id(uintdimindx)get_global_size(uintdimindx)get_local_size(uintdimindx)get_num_groups()get_work_dim()

40. API de suporte: kernels • Funções de sincronizaçãobarrier(cl_mem_fence_flags flags)Flags:CLK_LOCAL_MEM_FENCECLK_GLOBAL_MEM_FENCE

41. API de suporte: hospedeiro • Ilustração através de um exemplo prático • Aplicação completa • Kernelpara subtração dos elementos de dois arrays

42. Exemploprático • Exemplo de códigoexecutado no hospedeiro • Ilustra as principaisetapasparaodesenvolvimento de soluçõesOpenCL • Multi-plataforma

43. Exemploprático #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/opencl.h> #endif

44. Exemploprático cl_platform_idplatformId; cl_device_iddeviceId; clGetPlatformIDs(1, &platformId, NULL); clGetDeviceIDs( platformId, CL_DEVICE_TYPE_GPU, 1, &deviceId, NULL);

45. Exemploprático cl_context context = clCreateContext( NULL, 1, &deviceId, NULL, NULL, NULL); cl_command_queue queue = clCreateCommandQueue( context, deviceId, NULL, NULL);

46. Exemploprático const char* source = "__kernel void ArrayDiff( \ __global const int* a, \ __global const int* b, \ __global int* c) \ { \ int id = get_global_id(0);\ c[id] = a[id] - b[id]; \ }";

47. Exemploprático cl_program program = clCreateProgramWithSource( context, 1, &source, NULL, NULL);

48. Exemploprático clBuildProgram( program, 0, NULL, NULL, NULL, NULL); cl_kernel kernel = clCreateKernel( program,"ArrayDiff", NULL);

49. Exemploprático #define N ... cl_membufA = clCreateBuffer( context, CL_MEM_READ_ONLY, N * sizeof(int), NULL, NULL); cl_membufB = ...; cl_membufC = ...; CL_MEM_READ_WRITE

50. Exemploprático int* hostA; clEnqueueWriteBuffer( queue, bufA, CL_TRUE, 0, N * sizeof(int), hostA, 0, NULL, NULL);