Programação Paralela em Ambientes Computacionais Heterogêneos com OpenCL

1. ProgramaçãoParalelaemAmbientesComputacionaisHeterogêneos com OpenCL César L. B. Silveira Prof. Dr. Luiz G. daSilveira Jr. Prof. Dr. Gerson Geraldo H. Cavalheiro 28 de outubro de 2010

2. contato@v3d.com.br cesar@v3d.com.br

3. Introdução • Ambientes computacionais atuais: CPUs e GPUs • Ambientes heterogêneos • CPUsmulticore e GPUsmany-core • Processamento paralelo • Alto desempenho • Como utilizar todo este poder computacional?

4. Introdução • Programação paralela • CPUs • Multithreading • Múltiplos processos • GPUs • APIs gráficas (OpenGL, DirectX) • Shaders • Tecnologias proprietárias • NVIDIA CUDA • AMD Streaming SDK

5. Introdução • Problema: múltiplas ferramentas e paradigmas • Alto custo de implementação • Aplicações dependentes de plataforma • Necessidade de um padrão para programação paralela

6. Introdução • Open Computing Language (OpenCL) • Padrãoaberto, livre de royalties, paraprogramaçãoparalelaemambientesheterogêneos • MantidopeloKhronos Group desde 2008 • Define framework • Arquitetura • Linguagem • API

7. Introdução Camadas de uma aplicação OpenCL

8. 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]; } }

9. 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]; }

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

11. ArquiteturaOpenCL • Modelo de plataforma • Entidades do ambienteOpenCL • Hospedeiro (host) • Agregadispositivos (devices) quesuportamopadrão • Coordenaexecução • Dispositivo (device) • Composto de unidades de computação (compute units)

12. ArquiteturaOpenCL • Modelo de plataforma • Unidade de computação • Composta de elementos de processamento (processing elements) • Elemento de processamento • Realizacomputação

13. ArquiteturaOpenCL

14. ArquiteturaOpenCL • Modelo de execução • Descreveinstanciaçãoeidentificação de kernels • Espaço de índices (NDRange) • Até 3 dimensões • Itens de trabalho (work-items) organizadosemgrupos de trabalho (work-groups)

15. ArquiteturaOpenCL • Modelo de execução • Item de trabalho • Instância de um kernel • Grupo de trabalho • Permitesincronização entre itens de trabalho

16. ArquiteturaOpenCL • Modelo de execução: identificação • Item de trabalho • Identificadoresglobais no espaço de índices • Identificadoreslocais no grupo de trabalho • Grupo de trabalho • Identificadores no espaço de índices • Um identificadorpordimensão do espaço de índices

17. ArquiteturaOpenCL

18. ArquiteturaOpenCL • Modelo de execução: objetos • Contexto (context): agregadispositivos, kernelseoutrasestruturas • Fila de comandos (command queue): comandospara um dispositivo • Comandos de I/O paramemória, execução de kernels, sincronizaçãoeoutros

19. ArquiteturaOpenCL • Modelo de execução: objetos • Objeto de programa (program object): encapsulacódigo de um oumaiskernels • Base paracompilação de kernels • Objeto de memória(memory object): comunicação com a memória dos dispositivos • Buffers: acessodireto (ponteiros) • Images: acesso especial via samplers, texturas

20. ArquiteturaOpenCL • Modelo de execução • Relacionado com modelo de plataforma • Item de trabalho: elemento de processamento • Grupo de trabalho: unidade de computação

21. 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 • Execuçãoparalela de tarefasdistintas • Útilquandohámuitooverhead de I/O entre hospedeiroedispositivo

22. ArquiteturaOpenCL • Modelo de memória • Memória global (global memory) • Compartilhada entre todositens de trabalho do espaço de índices • Escritaeleitura • Memória local (local memory) • Compartilhadaapenas entre itens de trabalho do mesmogrupo de trabalho • Escritaeleitura

23. ArquiteturaOpenCL • Modelo de memória • Memóriaprivada (private memory) • Exclusiva de cada item de trabalho • Escritaeleitura • Memóriaconstante (constant memory) • Compartilhada entre todositens de trabalho do espaço de índices • Somenteleitura

24. ArquiteturaOpenCL

25. ArquiteturaOpenCL • Modelo de memória: consistência • Um item de trabalholêcorretamenteoqueoutrosescrevem? • Memóriaéconsistentepara: • Um único item de trabalho • Grupo de trabalhoemumabarreira (sincronização)

26. Arquitetura OpenCL • Item de trabalho: memória consistente

27. Arquitetura OpenCL • Grupo de trabalho: memória consistente após barreira (barrier)

28. ArquiteturaOpenCL • Modelo de memória: consistência • Nãohásincronização entre grupos de trabalho

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

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

31. 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;

32. LinguagemOpenCL C • Definição de kernels • Qualificador__kernel__kernel void f(…){ …}

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

34. LinguagemOpenCL C • Restrições • Ponteirosparafunçãonãosãosuportados • Argumentosparakernels nãopodem ser ponteirosparaponteiros • Funçõese macros com númerovariável de argumentosnãosãosuportados • Qualificadoresextern, staticeautonãosãosuportados

35. LinguagemOpenCL C • Restrições • Nãohásuporte a recursão • Elementos de structseunions devempertenceraomesmoespaço de endereçamento • Cabeçalhospadrão (stdio.h, stdlib.h, etc.) nãoestãopresentes

36. 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()

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

38. API de suporte: kernels • Sincronização__kernel void aKernel( __global float* in, __global float* out){intgid = get_global_id(0); int lid = get_local_id(0); intlsize = get_local_size(0); __local float a[lsize];

39. API de suporte: kernels a[lid] = in[gid];barrier(CLK_LOCAL_MEM_FENCE);out[gid] =a[(lid + 4) % lsize] * 2;barrier(CLK_GLOBAL_MEM_FENCE);}

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

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

42. Exemploprático • Etapas • Aquisição do dispositivo (GPU) • Criação do contexto • Compilação do kernel • Criação dos buffersparaentradaesaída • Escrita dos dados de entrada • Execução do kernel • Leitura dos dados de saída • Liberação de recursos

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 = …;

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