CUDA/Open CL

  • Published on
    14-Jun-2015

  • View
    173

  • Download
    45

Embed Size (px)

DESCRIPTION

Conceitos bsicos de CUDA e OpenCL: das arquiteturas de computadores s diretivas de programao.

Transcript

<ul><li> 1. CUDA/OpenCLArquiteturas Avancadas de ComputadoresKrissia de ZawadzkiInstituto de Fsica de Sao Carlos - Universidade de Sao Paulo06 de Maio 2014Krissia de Zawadzki CUDA/OpenCL 1 / 61</li></ul><p> 2. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoOutline1 CUDA - Introduc~ao2 GPU e CUDA3 Programando em CUDA4 OpenCL5 Caos6 Conclus~aoKrissia de Zawadzki CUDA/OpenCL 2 / 61 3. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoCUDA - IntroducaoCUDACompute Uni 4. ed Device Architecture3 Plataforma de computac~ao paralela emodelo de programac~ao3 Desenvolvido pela NVIDIA eimplementada para GPU's NVIDIA3 Conjunto de instruc~oes e memoriadiretamente acessveis ao programadorKrissia de Zawadzki CUDA/OpenCL 3 / 61 5. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoBackground historicoBackground historico2002: Stanford UniversitySteam processingPrototipo e arquitetura muito parecidacom GPUBaseline programmablestream processorKrissia de Zawadzki CUDA/OpenCL 4 / 61 6. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoBackground historicoBackground historico2002: GeForce 3 e ATI Radeon 97003 Shaders programaveisKrissia de Zawadzki CUDA/OpenCL 5 / 61 7. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoBackground historicoBackground historico2006: GeForce 8Arquitetura uni 8. cada CUDA!Krissia de Zawadzki CUDA/OpenCL 6 / 61 9. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoBackground historicoBackground historico2006: GeForce 83 Programabilidade realmenteexvel3 Revolucionou os conceitos depipeline de pixel e vertices3 Cadeia de processadoresKrissia de Zawadzki CUDA/OpenCL 7 / 61 10. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoGPUs com suporte para CUDAGPUs com suporte para CUDAGeForce83 8, 9, 100, 200, 400, 500 e600-series (mn 256MB memloc)3 1.0, 1.1, 1.2, 1.3, 2.0, 2.1, 3.0,3.5 e 5.0GeForce GTX-750 (5.0)NVS3 Quadro 295, 420, 4503 NVIDIA NVS 300, 315, 5103 1.1, 1.2, 2.1 e 3.0NVIDIA NVS 510 (3.0)Krissia de Zawadzki CUDA/OpenCL 8 / 61 11. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoGPUs com suporte para CUDAGPUs com suporte para CUDAQUADRO3 Quadro NVS, Quadro FX, Quadro K3 1.0, 1.1, 1.2, 1.3, 2.0, 2.1, 3.0 e3.5QUADRO K600 (3.5)Tesla3 D780, C870, C1060, C2050/2070,C20753 K20, K403 1.0, 1.3, 2.0, 3.5Tesla K20 (3.5)Krissia de Zawadzki CUDA/OpenCL 9 / 61 12. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoAplicacoesAlem de processamento graficoKrissia de Zawadzki CUDA/OpenCL 10 / 61 13. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoAplicacoesKrissia de Zawadzki CUDA/OpenCL 11 / 61 14. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoAplicacoesGanho de desempenho em aplicacoes cientficasKrissia de Zawadzki CUDA/OpenCL 12 / 61 15. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoGPUGPU vs. CPUCPU3 Codigos sequenciais3 Baixa lat^encia3 Controle complexoGPU3 Paralelismo de dados3 Alto throughtput3 Aritmetica com pouco controleKrissia de Zawadzki CUDA/OpenCL 13 / 61 16. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoGPUGPU vs. CPUCPU3 Fluxo iterativo7 Tempo de computacao7GPU3 Operacoes simult^aneas7 Desvio de fluxo7Krissia de Zawadzki CUDA/OpenCL 13 / 61 17. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoGPUCPU vs. GPUKrissia de Zawadzki CUDA/OpenCL 14 / 61 18. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoGPUCPU vs. GPUKrissia de Zawadzki CUDA/OpenCL 14 / 61 19. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoLinguagens e modelos de programacao paralelaLinguagens e modelos de programacao paralelaOpenMPshared memorylimite de centenas de nosCUDAalta escalabilidadeportabilidade e mais simples!MPIcapacidade de nos100.000esforco para portar o codigoOpenCLmodelo de programacao padronizadosuporte para AMD/ATI, NVIDIA,Apple e IntelKrissia de Zawadzki CUDA/OpenCL 15 / 61 20. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoArquitetura da GPUGPU - Unified processor array (GeForce 8800 GTX)Krissia de Zawadzki CUDA/OpenCL 16 / 61 21. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoEstrutura de um programa CUDAEstrutura de um programa CUDATrechos seriais ou com fraco paralelismo no codigo C do hostPorc~ao altamente paralela no codigo C do kernel associado ao deviceKrissia de Zawadzki CUDA/OpenCL 17 / 61 22. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoEstrutura de um programa CUDACUDA Threads paralelas// de dados: Todas as threads rodam o mesmo codigothreadIdx: identi 23. cador da thread de memoria e controleKrissia de Zawadzki CUDA/OpenCL 18 / 61 24. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoEstrutura de um programa CUDACUDA Thread BlocksThread Blocks: Array (dim 2x2) the threads que cooperam entre sibloco: memoria compartilhada, operac~oes at^omicas ebarreiras de sincronizac~aoblockIdx: identi 25. cador do bloco em um gridKrissia de Zawadzki CUDA/OpenCL 19 / 61 26. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoEstrutura de um programa CUDACUDA IdsIds sao uteis para identificar osdados sob os quais cada thread iratrabalharConveniente para simplificar o de memoria em dadosmultidimensionaisblockIdx:1D (blockIdx.x)2D (blockIdx.x, blockIdx.y)threadIdx:1D (threadIdx.x),2D (threadIdx.x, threadIdx.y),3D (threadIdx.x, threadIdx.y, threadIdx.z)Krissia de Zawadzki CUDA/OpenCL 20 / 61 27. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoEstrutura de um programa CUDAPar^ametros de configuracaoblock:dimBlock ( Widthx , Widthy , , Widthz )grid:dimGrid (Wgridx , Wgridy , Wgridz )Kernel launching// Setup the execution configurationdim3 dimBlock (Width , Width , 1);dim3 dimGrid (1, 1, 1);// Launch the device computation threads !MyKernelFunctiondimGrid , dimBlock ( args );Krissia de Zawadzki CUDA/OpenCL 21 / 61 28. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoModelo de Memoria CUDAModelo de Memoria CUDAMemoria Global:comunicacao host-deviceR/Wconteudo visvel por todas asthreadstipicamente implementadacomo DRAMacesso de longa lat^encia(400-800 ciclos)7 congestionamentothroughput limitado (a 177GB/s na GTX8800)Krissia de Zawadzki CUDA/OpenCL 22 / 61 29. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoModelo de Memoria CUDAModelo de Memoria CUDAMemoria Constante:read onlybaixa lat^encia e alta largurade banda quando todas asthreadas acessam o mesmolocalKrissia de Zawadzki CUDA/OpenCL 22 / 61 30. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoModelo de Memoria CUDAModelo de Memoria CUDAMemoria Compartilhada:3 rapidaaltamente paralelaapenas um bloco tem acessoKrissia de Zawadzki CUDA/OpenCL 22 / 61 31. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoModelo de Memoria CUDAModelo de Memoria CUDARegistradores3 o componente da memoriada GPU mais rapidoacessvel por uma threadaKrissia de Zawadzki CUDA/OpenCL 22 / 61 32. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoVariaveis CUDAVariaveis CUDAkernel: a variavel deve ser declarada no escopo da funcao kernel ficadisponvel somente no kernelapplication: a variavel deve ser declarada fora de qualquer funcaoconstant: a variavel deve ser declarada fora de qualquer funcao limitado (a 64KB na GTX8800)Krissia de Zawadzki CUDA/OpenCL 23 / 61 33. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoCUDA Kernel FunctionsCUDA Kernel FunctionsKernel functions:implementam o trechoparalelo de codigo a serexecutado no deviceSua chamada pode ser feitacom as con 34. gurac~oes deblocos e de threadsKrissia de Zawadzki CUDA/OpenCL 24 / 61 35. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoCUDA Kernel FunctionsAtribuicao de threadsbase block-by-block7 Runtime systemcoordena os blocos e asthreads a seremexecutadas: mantem alista de blocos e associanovos blocos a SM'slivresrecursos do SMunidades aritmeticasnumero de threads quepodem ser rastreadas eescalonadassimultaneamenteKrissia de Zawadzki CUDA/OpenCL 25 / 61 36. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoCUDA Kernel FunctionsCUDA WarpsWarps : conjunto de threads comndices consecutivosA capacidade do warp (num.de threads) e dependente daimplementac~aoWarp e a unidade paraescalonar threads no SMSIMDordem qualquer entre warps7 diverg^encia causada por branchsKrissia de Zawadzki CUDA/OpenCL 26 / 61 37. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoCUDA Kernel FunctionsCompilador NVCCBaseado no Open64(opensource originario doMIPSPro - SGI).Implementado em C e C++.NVidia atualmente investindono LLVM.Existe um utilitario queconverte LLVM IR (geradopor qualquer frontend decompilador LLVM) em PTX,que pode ser programado nasGPUs NVidia.Krissia de Zawadzki CUDA/OpenCL 27 / 61 38. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoCUDA Kernel FunctionsDefici^enciasEquipe preferiu implementar codigo para o desa 39. o ECC2K-130diretamente em codigo de maquina.BERNSTEIN, D. J. et al. Usable Assembly Language for GPUs: ASuccess Story. In: Workshop Records of Special-Purpose Hardwarefor Attacking Cryptographic Systems SHARCS 2012. [s.n.], 2012.p. 169178.Compilador NVCC muito lento para lidar com kernels contendomuitas instruc~oes.Registradores alocados de forma pouco e 40. ciente { muitas variaveisacabaram tendo de ser alocadas pelo NVCC na memoriacompartilhada.Varios truques necessarios para obter uma implementac~ao em Caceitavel. Implementac~ao em Assembly 148% mais rapida quemelhor implementac~ao em C.Krissia de Zawadzki CUDA/OpenCL 28 / 61 41. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploMultiplicacao Matricial: o Hello World do CUDA = * =,,3 Paralelismo de dados!Cada elemento de pode sercalculado simultaneamente aosdemais!Krissia de Zawadzki CUDA/OpenCL 29 / 61 42. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploRepresentacao matricial em CAlocac~ao de memoria no C para arrays bidimensionais:Krissia de Zawadzki CUDA/OpenCL 30 / 61 43. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploCodigo main C sequencial (host)int main ( void ){// 1. Alocamos e inicializamos as matrizes M, N e P// Funcoes I/O leem as matrizes M e N...// 2. Multiplicacao M * NMatMul (M,N,P, Width );...// 3. Funcao I/O para escrever a saida P// Liberamos a memoria de M, N e Preturn 0;}Krissia de Zawadzki CUDA/OpenCL 31 / 61 44. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploFuncao C sequencial (host)void MatMul ( float *M; float *N, float *P, int Width ){for ( int i = 0; iWidth ; ++i)for (int j = 0; jWidth ; ++j){float sum = 0;for ( int k = 0; kWidth ; ++k){float m = M[i* Width + k];float n = N[k* width + j];sum += m * n;}P[i * Width + j] = sum ;}}Krissia de Zawadzki CUDA/OpenCL 32 / 61 45. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploPortando o codigo para CUDA - alocacao de memoriacudaMalloc()aloca um objeto naMemoria Globalpar^ametros: endereco deum ponteiro para o objetoalocado, tamanho do objetocudaFree()libera um objeto naMemoria Globalpar^ametro: ponteiro para oobjetoKrissia de Zawadzki CUDA/OpenCL 33 / 61 46. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploPortando o codigo para CUDA - alocacaoExemplo:int Width =64;float * Md , Nd;int size = Width * Width * sizeof ( float );cudaMalloc (( void **) Md , size );cudaMalloc (( void **) Nd , size );...cudaFree (Md );cudaFree (Nd );Krissia de Zawadzki CUDA/OpenCL 34 / 61 47. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploPortando o codigo para CUDA - transfer^encia de dadoscudaMemcpy()transfere dados entre o hoste o deviceAssncronapar^ametros:ponteiro para o destinoponteiro para a fonte numerode bytes a serem copiadostipo de transfer^enciatipos:Host to HostHost to DeviceDevice to HostDevice to DeviceKrissia de Zawadzki CUDA/OpenCL 35 / 61 48. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploPortando o codigo para CUDA - transfer^encia de dadosExemplo:...int size = Width * Width * sizeof ( float );...cudaMemcpy (Md , M, size , cudaMemcyHostToDevice );cudaMemcpy (Nd , N, size , cudaMemcyHostToDevice );...cudaMemcpy (Pd , P, size , cudaMemcyDeviceToHost );Krissia de Zawadzki CUDA/OpenCL 36 / 61 49. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploPortando o codigo para CUDA - funcao MatMul no devicevoid MatMul ( float *M; float *N, float *P, int Width ){int size = Width * Width * sizeof ( float );float * Md , Nd , Pd;// 1. Alocamos memoria no device para M, N e PcudaMalloc (( void **)) Md , size );cudaMemcpy (Md , M, size , cudaMemcyHostToDevice );cudaMalloc (( void **)) Nd , size );cudaMemcpy (Nd , N, size , cudaMemcyHostToDevice );cudaMalloc (( void **)) Pd , size );// 2. Evocamos a funcao kernel para a multiplicacao// 3. Copiamos o resultado P para a memoria do hostcudaMemcpy (Pd , P, size , cudaMemcyDeviceToHost );// Liberamos as memorias de M, N e P no devicecudaFree (Md ); cudaFree (Nd ); cudaFree (Pd );}Krissia de Zawadzki CUDA/OpenCL 37 / 61 50. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploFuncao kernel MatMulKrissia de Zawadzki CUDA/OpenCL 38 / 61 51. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploKernel function - um pouco mais sobre especificacoesglobal de 52. ne uma func~ao kerneldevice e host podem ser usadas simultaneamente7 recurs~oes7 variaveis estaticas7 chamadas indiretas de func~oes por ponteirosKrissia de Zawadzki CUDA/OpenCL 39 / 61 53. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploConfiguracao de execucaoExemplo: De 54. nir a multiplicac~ao matricial quando = 32 emblocos Grids 2D com (2x2) blocos Blocos 2D com (16x16) threadsKrissia de Zawadzki CUDA/OpenCL 40 / 61 55. CUDA - Introducao GPU e CUDA Programando em CUDA OpenCL Caos ConclusaoProgramando em CUDA: exemploConfiguracao de execucaoA configuracao define a dimensao do problema!!!! No exemplo anterior, usando blocos1D podemos apenas trabalhar com = 16 !Solucao: Manipular dimGrid e dimBlock e dividir o calculo de pedacos da matrizresultado entre threads e blocos!Krissia de Zawadzki CUDA/OpenCL 41 / 61 56. CUDA - Introducao GPU e CUDA Programando em...</p>