1 / 66

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

Programação Paralela em Ambientes Computacionais Heterogêneos com OpenCL. César L. B. Silveira Prof. Dr. Luiz G. da Silveira Jr. Prof. Dr. Gerson Geraldo H. Cavalheiro 28 de outubro de 2010. contato@v3d.com.br cesar@v3d.com.br. Introdução.

wattan
Download Presentation

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

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript


  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);

More Related