Tópicos sobre GPGPU em CUDA

Tópicos sobre GPGPU em CUDA Paulo A. Pagliosa pagliosa@facom.ufms.br Faculdade de Computação Universidade Federal de Mato Grosso do Sul Natal, maio de 2013

Apresentação • Conteúdo • Introdução a GPGPU (www.gpgpu.org) • O que é CUDA • Fundamentos de programação CUDA • Práticas de programação • APIs e ferramentas • Estudo de caso

Introdução a GPGPU • GPUs: unidades de processamento gráfico • Originalmente projetadas para processamento gráfico 3D • Exemplos • NVIDIA GeForce 8800 GTX • NVIDIA Tesla C1060

Introdução a GPGPU • Pipeline gráfico fixo (OpenGL) Software Hardware Cena Dados de geometria Processamento de geometria Aplicação Rasterização Processamento de fragmentos Dados de pixels Fontes de luz Framebuffer Materiais Exibição

Introdução a GPGPU • Pipeline gráfico fixo (OpenGL) Hardware Processamento de geometria Transformação Iluminação Rasterização Projeção Processamento de fragmentos Montagem de primitivos Framebuffer Recorte Exibição Primitivos

Introdução a GPGPU • Pipeline gráfico fixo (OpenGL) Hardware Processamento de geometria Rasterização Processamento de fragmentos Framebuffer Exibição

Introdução a GPGPU • Pipeline gráfico fixo (OpenGL) Hardware Processamento de geometria Mapeamento de textura Rasterização Tonalização Aplicação de neblina Processamento de fragmentos Composição Framebuffer Exibição

Introdução a GPGPU • GPUs de primeira geração (1998) • Pipeline fixo • Renderização de triângulos pré-transformados • Exemplos: TNT2, Voodoo3 • GPUs de segunda geração (1999-2000) • Transformações geométricas • Iluminação • Velocidade de renderização maior • Exemplos: GeForce 2, ATI 5700

Introdução a GPGPU • Pipeline gráfico programável (OpenGL 3.1) Hardware Processamento de geometria Transformação Iluminação Rasterização Projeção Processamento de fragmentos Montagem de primitivos Framebuffer Recorte Exibição Primitivos

Introdução a GPGPU • Pipeline gráfico programável (OpenGL 3.1) Hardware Processamento de geometria Mapeamento de textura Rasterização Tonalização Aplicação de neblina Processamento de fragmentos Composição Framebuffer Exibição

Introdução a GPGPU • Shaders • Módulos que executam em GPU • Um shader pode ser de: • Vértice • Geometria (SM 4) • Fragmento • Substituem a funcionalidade do pipeline fixo • Programa em GPU: um ou mais shaders Programa GPU Shader de vértice Shader de geometria Shader de fragmento

Introdução a GPGPU • Shader de fragmento • Opera isoladamente sobre um fragmento • Entrada • Variáveis pré-calculadas pela OpenGL • Variáveis definidas pela aplicação • Saída: cor do fragmento • Operações • Mapeamento de textura • Tonalização • Aplicação de neblina • Fragmento pode ser descartado • Coordenadas do fragmento não podem ser mudadas

Introdução a GPGPU • GPUs de terceira geração (2001-2002) • Pipeline programável com shaders de vértices • Número limitado de instruções sem ponto flutuante • Programação em linguagem de montagem • Exemplos: GeForce3 e 4, ATI 8500 • GPUs de quarta geração (2003-2006) • Shaders de vértices e fragmentos • Expansão do número de instruções • Ponto flutuante de 32 bits e dados em texturas • Surgimento de linguagens de shaders • Exemplos: GeForce FX, 6 e 7; ATI 9700 e 9800

Introdução a GPGPU • NVIDIA GeForce 6800 • Até 6 processadores de vértices • 16 processadores de fragmentos

Introdução a GPGPU • Linguagens de shaders • Cg (C for graphics), NVIDIA • HLSL (High Level Shader Language), Microsoft • GLSL (OpenGL Shader Language)

Introdução a GPGPU • GPGPU: programação genérica em GPU • GPGPU com shaders • Modelo de programação: streams e kernels • Streams: dados de entrada e saída • Vetores em CPU — texturas em GPU • Kernels: shaders (de fragmento) • Saída: renderização em textura • Execução: rasterização • Mapping Computational Concepts to GPUs (GPU Gems 2) • Dificuldades

Introdução a GPGPU • Exemplos • Simulação dinâmica de corpos rígidos • Multiplicação de matrizes potência de dois

Introdução a GPGPU • Desvantagens de GPGPU com shaders • GPU é programada através de uma API gráfica • Curva de aprendizado da API • Overhead para aplicações não gráficas • Flexibilidade • Memória da GPU pode ser lida (gather) mas não pode ser escrita (scatter) de maneira geral • Shader de fragmento produz apenas saídas RGBA • CUDA: Compute Unified Device Architecture • GPUs de quinta geração (2007-) • Computação de propósito geral • Exemplos: GeForce 8, 9, 100 e 200

Introdução a GPGPU • NVIDIA GeForce 8800 GTX • 16 multiprocessadores (SMs) • 8 processadores (SPs) por multiprocessador

Introdução a GPGPU • Aplicações CUDA • Visualização e simulação • http://www.nvidia.com/object/cuda_home_new.html

Introdução a GPGPU • CUDA: arquitetura de computação paralela para GPGPU

Introdução a GPGPU • Por que CUDA? • Desempenho versus custo • NVIDIA é líder de mercado: mais de 108 GPUs • Programação paralela para muitos • Tesla, Fermi e Kepler • Instalação de CUDA • Driver • CUDA toolkit (compilador e bibliotecas) • CUDA SDK (utilitários e exemplos) • https://developer.nvidia.com/cuda-downloads

Introdução a GPGPU • Bibliografia (http://docs.nvidia.com/cuda/index.html) • CUDA Programming Guide • CUDA Reference Manual • CUDA Best Practices Guide

Fundamentos de programação CUDA • Modelo de programação • Host: executa a aplicação (CPU) • Dispositivo (GPU) • Coprocessador da CPU • Executa kernels • Host e dispositivo tem DRAMs próprias • Host: • Aloca memória no dispositivo • Transfere dados de entrada para o dispositivo • Dispara a execução de kernels • Transfere dados resultantes do dispositivo • Libera memória no dispositivo

Fundamentos de programação CUDA • Modelo de programação • Kernel • Função geralmente escrita em C para CUDA • Executa no dispositivo N vezes em N threads em paralelo • Threads são organizadas em blocos • Um bloco é um arranjo 1D, 2D ou 3D de threads • Cada thread de um bloco tem um índice 1D, 2D ou 3D • Blocos são organizados em grids • Um grid é um arranjo 1D ou 2D de blocos • Cada bloco de um grid tem um índice 1D ou 2D • Os blocos de um grid têm o mesmo número de threads

Fundamentos de programação CUDA • Exemplo: kernel executando em 72 threads • Grid 2D com: • Dimensão 3×2×1 • 6 blocos • Blocos 2D com: • Dimensão 4×3×1 • 12 threads cada

Fundamentos de programação CUDA • Modelo de programação • Um kernel é uma função que: • Começa com o especificador __global__ • Tem tipo de retorno void • Kernels podem invocar outras funções que: • São especificadas como __device__ • Podem invocar outras especificadas como __device__(GPUs Kepler podem invocar kernels) • Funções que executam no dispositivo: • Não admitem número variável de argumentos • Não admitem variáveis estáticas • GPUs não Fermi não admitem recursão nem variáveis do tipo endereço de função

Fundamentos de programação CUDA • Modelo de programação • Kernels são invocados do host (ou de Kepler) • Um dispositivo executa um kernel de cada vez (Fermi pode mais)

Fundamentos de programação CUDA • Modelo de programação • Configuração de execução de um kernel • Dimensões do grid e dos blocos • Tamanho da memória compartilhada (opcional) • Especificada na invocação (ou lançamento) do kernel • Dimensão 3D é representada por objeto do tipo dim3 • Um grid: • 1D de dimensão dim3 dG tem dG.x × 1 × 1 blocos • 2D de dimensão dim3 dG tem dG.x × dG.y × 1 blocos • Um bloco: • 1D de dimensão dim3 dB tem dB.x × 1 × 1 threads • 2D de dimensão dim3dB tem dB.x × dB.y × 1 threads • 3D de dimensão dim3dB tem dB.x × dB.y × dB.z threads

Fundamentos de programação CUDA • Modelo de programação • Identificador global de uma thread • Pode ser usado para indexar vetores e matrizes em funções __global__ou __device__ • Determinado a partir das variáveis pré-definidas: • dim3 gridDim • Dimensão do grid • dim3 blockDim • Dimensão do bloco • dim3 blockIdx • Índice do bloco no grid • dim3 threadIdx • Índice da thread no bloco

Fundamentos de programação CUDA • Modelo de programação • Hierarquia de memória acessada por uma thread • Memória compartilhada do bloco da thread • Visível para todas as threads do bloco • Tempo de vida do bloco • Memória local da thread • Memória global • Memória constante • Memória de textura Tempo de vida da aplicação Somente leitura

Fundamentos de programação CUDA • Modelo de programação • Capacidade de computação • 1.x (Tesla) • 2.x (Fermi) • 3.x (Kepler) • Especificações: depende do dispositivo • Número de multiprocessadores e processadores • Dimensões de grids e blocos • DRAM • Memória compartilhada, etc.

Fundamentos de programação CUDA • Interface de programação • C para CUDA e API de runtime • API de driver • Ambas APIs têm funções para: • Gerência de memória no dispositivo • Transferência de dados entre host e dispositivo • Gerência de sistemas com vários dispositivos, etc. • API de runtime tem também funções para: • Gerenciamento de threads • Detecção de erros • Manual de referência de CUDA APIs exclusivas

Fundamentos de programação CUDA • Programa básico em C para CUDA • Seleção do dispositivo a ser usado • Alocação de memória no host • Dados de entrada na memória do host • Alocação de memória no dispositivo • Transferência de dados do host para dispositivo • Invocação do(s) kernel(s) • Transferência de dados do dispositivo para host • Liberação de memória no host • Liberação de memória no dispositivo • Finalização

Fundamentos de programação CUDA • Programa básico em C para CUDA • Seleção do dispositivo a ser usadocudaSetDevice() • Alocação de memória no dispositivocudaMalloc() • Transferência de dados entre host e dispositivocudaMemcpy() • Liberação de memória no dispositivocudaFree() • FinalizaçãocudaDeviceReset()

Fundamentos de programação CUDA • Exemplo: multiplicação de matrizes • Sem memória compartilhada

Fundamentos de programação CUDA • Memória compartilhada • Disponível para threads de um mesmo bloco • Acesso 400 a 600 mais rápido que memória global • Permite colaboração entre threads (do bloco) • Cada thread do bloco transfere dados da memória global para memória compartilhada • Após a transferência, threads devem ser sincronizadas com __syncthreads() • Kernel efetua operações usando os dados da memória compartilhada: cada thread pode usar dados carregados por outras threads do bloco • Se necessário, threads podem ser sincronizadas com __syncthreads(); cada thread do bloco transfere dados da memória compartilhada para memória global

Fundamentos de programação CUDA • Exemplo: multiplicação de matrizes • Com memória compartilhada

Fundamentos de programação CUDA • Exemplo: transposição de matrizes Mem. compartilhada __syncthreads()

Fundamentos de programação CUDA • Exemplo: transposição de matrizes Mem. compartilhada

Fundamentos de programação CUDA • Implementação em hardware • Arquitetura: arranjo de multiprocessadores (SMs) • Cada SM consiste de: • 8 processadores (SPs) • 1 unidade de instrução • Memória compartilhada • Cada SM: • Executa threads de um bloco em grupos de 32: warp • Implementa barreira de sincronização:__syncthreads() • Emprega arquitetura SIMT:Single Instruction Multiple Thread

Dicas de performance • Estratégias básicas • Maximização da execução em paralelo • Estruturação do algoritmo • Escolha da configuração de execução do kernel • Número de threads por bloco múltiplo do tamanho do warp • Mínimo de 64 threads por bloco • Configuração inicial: entre 128 e 256 threads por bloco • Evitar divergência dentro do mesmo warp • Otimização do uso de memória • Minimização de transferência de dados host/dispositivo • Acesso coalescido à memória global • Uso de mem. compartilhada • Acesso sem conflitos de bancos à mem. compartilhada • Otimização do uso de instruções

Dicas de performance • Memória global • O dispositivo é capaz de ler palavras de 4, 8 ou 16 bytes da memória global para registradores com UMA única instrução, DESDE que o endereço de leitura seja alinhado a (múltiplo de) 4, 8, ou 16. • A largura de banda da memória global é mais eficientemente usada quando acessos simultâneos à memória por threads de um meio-warp (durante a execução de uma instrução de leitura ou escrita) podem ser coalescidos em uma única transação de 32, 64 ou 128 bytes de memória.

Dicas de performance • Memória global • Coalescência em dispositivos 1.0 e 1.1 • Ocorre em transação de 64 ou 128 ou duas transações de 128 bytes • Threads de um meio-warp devem acessar: • Palavras de 4 bytes, resultando numa transação de 64 bytes, • Ou palavras de 8 bytes, resultando numa transação de 128 bytes, • Ou palavras de 16 bytes, resultando em duas transações de 128 bytes • Todas as 16 palavras (cada uma acessada por uma das 16 threads do meio-warp) devem estar no mesmo segmento de tamanho igual ao tamanho das transações (ou seja, 64, 128 ou 256 bytes). Como consequência, o segmento deve ser alinhado a este tamanho • Threads devem acessar palavras na sequência: a thread k deve acessar a palavra k

Dicas de performance • Memória global • Com coalescência

Dicas de performance • Memória global • Sem coalescência

Dicas de performance • Memória global • Coalescência com vetor de estruturas (AOS) • Tamanho da estrutura até 16 bytes • Alinhamento deve ser 4, 8 ou 16 bytes, dependendo do tamanho • Elementos devem estar no mesmo segmento da transação • Estruturas maiores que 16 bytes: reorganizar em estrutura de vetores (SOA) com elementos de tamanho até 16 bytes

Dicas de performance • Memória compartilhada • Dividida em módulos chamados bancos • Em dispositivos 1.x o número de bancos é 16 • Bancos tem largura de 32 bits • Palavras sucessivas de 32 bits estão em bancos sucessivos • Acessos a n endereços que estão em n bancos distintos são efetuados simultaneamente • Acessos a dois endereços que estão no mesmo banco geram um conflito de banco: acessos são serializados • Threads de um meio-warp devem acessar endereços em bancos distintos para evitar conflitos de bancos

Dicas de performance • Mem. compartilhada • Sem conflitos

Tópicos sobre GPGPU em CUDA