Capítulo 3: Modelos de Programação Paralela em Design de GPU
As Unidades de Processamento Gráfico (GPUs) evoluíram de aceleradores gráficos de função fixa para motores de computação altamente paralelos e programáveis, capazes de acelerar uma ampla gama de aplicações. Para permitir que os programadores aproveitem efetivamente o paralelismo maciço nas GPUs, vários modelos e APIs de programação paralela foram desenvolvidos, como NVIDIA CUDA, OpenCL e DirectCompute. Esses modelos de programação fornecem abstrações que permitem que os programadores expressem o paralelismo em suas aplicações, ocultando os detalhes de baixo nível do hardware da GPU.
Neste capítulo, exploraremos os conceitos e princípios-chave por trás dos modelos de programação paralela para GPUs, com foco no modelo de execução SIMT (Single Instruction, Multiple Thread), no modelo de programação e APIs CUDA, e na estrutura OpenCL. Também discutiremos técnicas para mapear algoritmos para arquiteturas de GPU a fim de alcançar alto desempenho e eficiência.
Modelo de Execução SIMT (Single Instruction, Multiple Thread)
O modelo de execução SIMT é o paradigma fundamental usado pelas GPUs modernas para alcançar um paralelismo maciço. No modelo SIMT, um grande número de threads executa o mesmo programa (chamado de kernel) em paralelo, mas cada thread tem seu próprio contador de programa e pode seguir caminhos de execução diferentes, com base em seu ID de thread e nos dados que opera.
Kernels e Hierarquia de Threads
Um kernel de GPU é uma função executada em paralelo por um grande número de threads. Ao lançar um kernel, o programador especifica o número de threads a serem criadas e como elas são organizadas em uma hierarquia de grids, blocos (ou matrizes de threads cooperativas - CTAs) e threads individuais.
- Um grid representa todo o espaço do problema e consiste em um ou mais blocos.
- Um bloco é um grupo de threads que podem cooperar e sincronizar umas com as outras via memória compartilhada e barreiras. As threads dentro de um bloco são executadas no mesmo núcleo de GPU (chamado de streaming multiprocessor).Arquivo em Português:
Unidade de Computação Paralela (ou Unidade de Processamento)
- Cada thread tem um ID único dentro do seu bloco e grade, que pode ser usado para calcular endereços de memória e tomar decisões de fluxo de controle.
Essa organização hierárquica permite que os programadores expressem tanto o paralelismo de dados (onde a mesma operação é aplicada a múltiplos elementos de dados) quanto o paralelismo de tarefas (onde diferentes tarefas são executadas em paralelo).
A Figura 3.1 ilustra a hierarquia de threads no modelo de execução SIMT.
Grade
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Bloco |
| | | |
Thread Thread ...Figura 3.1: Hierarquia de threads no modelo de execução SIMT.
Execução SIMT
No modelo de execução SIMT, cada thread executa a mesma instrução, mas opera sobre diferentes dados. No entanto, ao contrário do SIMD (Single Instruction, Multiple Data), onde todos os elementos de processamento executam em lockstep, o SIMT permite que as threads tenham caminhos de execução independentes e divirjam em instruções de ramificação.
Quando uma warp (um grupo de 32 threads em GPUs da NVIDIA ou 64 threads em GPUs da AMD) encontra uma instrução de ramificação, o hardware da GPU avalia a condição de ramificação para cada thread na warp. Se todas as threads seguirem o mesmo caminho (convergentes), a warp continua a execução normalmente. No entanto, se algumas threads seguirem caminhos diferentes (divergentes), a warp é dividida em duas ou mais subwarps, cada uma seguindo um caminho diferente. O hardware da GPU serializa a execução dos caminhos divergentes, mascarando as threads inativas em cada subwarp. Quando todos os caminhos são concluídos, as subwarps reconvergem e continuam a execução em lockstep.
A Figura 3.2 ilustra a execução SIMT com fluxo de controle divergente.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Ramificação |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
```Reconvergência
Figura 3.2: Execução SIMT com fluxo de controle divergente.
Este mecanismo de tratamento de divergência permite que o SIMT suporte um fluxo de controle mais flexível do que o SIMD, mas isso vem com o custo de uma eficiência SIMD reduzida quando ocorre a divergência. Os programadores devem se esforçar para minimizar a divergência dentro de um warp para alcançar um desempenho ideal.
Hierarquia de memória
As GPUs têm uma hierarquia de memória complexa para suportar os requisitos de alta largura de banda e baixa latência das cargas de trabalho paralelas. A hierarquia de memória geralmente consiste em:
- Memória global: O maior, mas o espaço de memória mais lento, acessível por todos os threads em um kernel. A memória global geralmente é implementada usando memória GDDR ou HBM de alta largura de banda.
- Memória compartilhada: Um espaço de memória rápido e on-chip compartilhado por todos os threads em um bloco. A memória compartilhada é usada para comunicação entre threads e compartilhamento de dados dentro de um bloco.
- Memória constante: Um espaço de memória somente leitura usado para difundir dados somente leitura para todos os threads.
- Memória de textura: Um espaço de memória somente leitura otimizado para localidade espacial e acessado via caches de textura. A memória de textura é mais comumente usada em cargas de trabalho gráficas.
- Memória local: Um espaço de memória privado para cada thread, usado para transbordamento de registros e estruturas de dados grandes. A memória local geralmente é mapeada para a memória global.
A utilização eficaz da hierarquia de memória é crucial para obter um alto desempenho nas GPUs. Os programadores devem ter como objetivo maximizar o uso da memória compartilhada e minimizar os acessos à memória global para reduzir a latência de memória e os gargalos de largura de banda.
A Figura 3.3 ilustra a hierarquia de memória da GPU.
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________Figura 3.3: Hierarquia de memória da GPU.
Modelo de Programação CUDA e APIs
CUDA (Compute Unified Device Architecture) é uma plataforma de computação paralela e modelo de programação desenvolvido pela NVIDIA para computação de propósito geral em GPUs. CUDA fornece um conjunto de extensões para linguagens de programação padrão, como C, C++ e Fortran, que permitem que os programadores expressem o paralelismo e aproveitem o poder computacional das GPUs NVIDIA.
Modelo de Programação CUDA
O modelo de programação CUDA é baseado no conceito de kernels, que são funções executadas em paralelo por um grande número de threads na GPU. O programador especifica o número de threads a serem lançadas e sua organização em um grid de blocos de threads.
CUDA introduz várias abstrações-chave para facilitar a programação paralela:
- Thread: A unidade básica de execução no CUDA. Cada thread tem seu próprio contador de programa, registradores e memória local.
- Bloco: Um grupo de threads que podem cooperar e sincronizar umas com as outras. As threads dentro de um bloco são executadas no mesmo streaming multiprocessador e podem se comunicar via memória compartilhada.
- Grid: Uma coleção de blocos de threads que executam o mesmo kernel. O grid representa todo o espaço do problema e pode ser unidimensional, bidimensional ou tridimensional.
CUDA também fornece variáveis embutidas (por exemplo, threadIdx, blockIdx, blockDim, gridDim) que permitem que as threads se identifiquem e calculem endereços de memória com base em sua posição na hierarquia de threads.
A Figura 3.4 ilustra o modelo de programação CUDA.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Bloco |
| | | |
Threads Threads ...Figura 3.4: Modelo de programação CUDA.
Hierarquia de Memória CUDAAqui está a tradução em português para o arquivo markdown "archy":
CUDA expõe a hierarquia de memória da GPU ao programador, permitindo o controle explícito sobre o posicionamento e movimento de dados. Os principais espaços de memória no CUDA são:
- Memória global: Acessível por todas as threads em um kernel e persiste através de lançamentos de kernel. A memória global tem a maior latência e é tipicamente usada para grandes estruturas de dados.
- Memória compartilhada: Uma memória rápida, on-chip, compartilhada por todas as threads em um bloco. A memória compartilhada é usada para comunicação entre threads e compartilhamento de dados dentro de um bloco.
- Memória constante: Um espaço de memória somente leitura usado para transmitir dados somente leitura para todas as threads. A memória constante é armazenada em cache e fornece acesso de baixa latência.
- Memória de textura: Um espaço de memória somente leitura otimizado para localidade espacial e acessado através de caches de textura. A memória de textura é mais comumente usada em cargas de trabalho gráficas.
- Memória local: Um espaço de memória privado para cada thread, usado para derramamento de registradores e grandes estruturas de dados. A memória local geralmente é mapeada para a memória global.
Os programadores podem alocar e transferir dados entre a memória do host (CPU) e do dispositivo (GPU) usando as APIs de runtime do CUDA, como cudaMalloc, cudaMemcpy e cudaFree.
A Figura 3.5 ilustra a hierarquia de memória do CUDA.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________Figura 3.5: Hierarquia de memória do CUDA.
Sincronização e Coordenação do CUDA
O CUDA fornece primitivas de sincronização e coordenação para permitir a cooperação e comunicação entre threads:
-
Sincronização de barreira: O __syncthreadA função s() atua como uma barreira que garante que todos os threads em um bloco tenham alcançado o mesmo ponto antes de prosseguir.
-
Operações atômicas: o CUDA suporta operações atômicas (por exemplo, atomicAdd, atomicExch) que permitem que os threads realizem operações de leitura-modificação-escrita na memória compartilhada ou global sem interferência de outros threads.
-
Primitivas de nível de warp: o CUDA fornece intrínsecos de nível de warp (por exemplo, __shfl, __ballot) que permitem uma comunicação e sincronização eficientes dentro de uma warp.
O uso adequado de primitivas de sincronização e coordenação é essencial para a escrita de programas paralelos corretos e eficientes em CUDA.
O Exemplo 3.1 mostra um kernel CUDA simples que realiza a adição de vetores.
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
int main() {
int *a, *b, *c;
int n = 1024;
// Alocar memória no host
a = (int*)malloc(n * sizeof(int));
b = (int*)malloc(n * sizeof(int));
c = (int*)malloc(n * sizeof(int));
// Inicializar os vetores de entrada
for (int i = 0; i < n; i++) {
a[i] = i;
b[i] = i * 2;
}
// Alocar memória no device
int *d_a, *d_b, *d_c;
cudaMalloc(&d_a, n * sizeof(int));
cudaMalloc(&d_b, n * sizeof(int));
cudaMalloc(&d_c, n * sizeof(int));
// Copiar os vetores de entrada do host para o device
cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
// Lançar o kernel
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
// Copiar o vetor de resultado do device para o host
cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
// Liberar a memória do device
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
// Liberar a memória do host
free(a);
free(b);
free(c);
returAqui está a tradução em português do arquivo Markdown, com os comentários do código traduzidos para o português, enquanto o código em si não foi traduzido:
n 0;
}Este código CUDA lança o kernel vectorAdd com numBlocks blocos e blockSize threads por bloco. O kernel realiza a adição elemento a elemento dos vetores de entrada a e b e armazena o resultado no vetor c. A sintaxe <<<...>>> é usada para especificar as dimensões da grade e do bloco ao executar um kernel.
Streams e Eventos CUDA
Streams e eventos CUDA fornecem um mecanismo para a execução concorrente e sincronização de kernels e operações de memória:
- Streams: Uma sequência de operações (lançamentos de kernel, cópias de memória) que são executadas em ordem. Diferentes streams podem ser executados concorrentemente, permitindo a sobreposição de computação e transferências de memória.
- Eventos: Marcadores que podem ser inseridos em um stream para registrar a conclusão de operações específicas. Eventos podem ser usados para fins de sincronização e temporização.
Streams e eventos permitem que os programadores otimizem o desempenho de seus aplicativos CUDA, sobrepondo a computação e as transferências de memória e explorando toda a capacidade do hardware da GPU.
O Exemplo 3.2 demonstra o uso de streams CUDA para sobrepor a execução de kernel e as transferências de memória.
// Criar dois streams
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// Copiar os dados de entrada para o dispositivo de forma assíncrona
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
// Executar kernels em diferentes streams
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
// Copiar os resultados de volta para o host de forma assíncrona
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
// Sincronizar streams
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);Neste exemplo, dois streams CUDA são criados. Os dados de entrada são copiados para o dispositivo de forma assíncrona usando cada stream. Em seguida, os kernels são executados nos diferentes streams, permitindo a sobreposição da computação e das transferências de memória.Aqui está a tradução em português deste arquivo Markdown. Para o código, não traduzi o código, apenas os comentários.
Estrutura de Trabalho OpenCL
OpenCL (Open Computing Language) é um padrão aberto e gratuito para programação paralela em plataformas heterogêneas, incluindo CPUs, GPUs, FPGAs e outros aceleradores. O OpenCL fornece um modelo de programação unificado e um conjunto de APIs que permitem que os desenvolvedores escrevam código paralelo portátil e eficiente.
Modelo de Programação OpenCL
O modelo de programação OpenCL é semelhante ao CUDA, com algumas diferenças-chave na terminologia e abstrações:
- Kernel: Uma função executada em paralelo por um grande número de work-items (threads) em um dispositivo OpenCL.
- Work-item: A unidade básica de execução no OpenCL, análoga a um thread no CUDA.
- Work-group: Uma coleção de work-items que podem se sincronizar e compartilhar dados através da memória local. Os work-groups são análogos aos thread blocks no CUDA.
- NDRange: Define o espaço de índice e a organização de work-items para a execução de um kernel. Pode ser unidimensional, bidimensional ou tridimensional.
O OpenCL também define um modelo de memória hierárquico semelhante ao CUDA:
- Memória global: Acessível por todos os work-items em todos os work-groups, análoga à memória global no CUDA.
- Memória local: Compartilhada por todos os work-items em um work-group, análoga à memória compartilhada no CUDA.
- Memória privada: Privada a um único work-item, análoga aos registradores no CUDA.
- Memória constante: Memória somente leitura acessível por todos os work-items.
Os kernels OpenCL são compilados em tempo de execução pelo runtime do OpenCL. O programa host pode consultar os dispositivos OpenCL disponíveis, selecionar um dispositivo apropriado, criar um contexto e compilar o kernel para esse dispositivo específico. Isso permite que aplicativos OpenCL sejam altamente portáveis em diferentes plataformas de hardware.
O Exemplo 3.3 mostra um kernel OpenCL que realiza a adição de vetores, semelhante ao exemplo CUDA no Exemplo 3.1.
__kernel void vectorAdd(__global const int *a, __global int *b, __global int *c) {
// Obtem o índice do work-item atual
int i = get_global_id(0);
// Realiza a adição de vetores
c[i] = a[i] + b[i];
}
```Aqui está a tradução em português do arquivo Markdown, com os comentários traduzidos, mas o código mantido intacto:
```c
__kernel void vector_add(
__global const int *a,
__global const int *b,
__global int *c,
int n) {
int i = get_global_id(0);
if (i < n) {
c[i] = a[i] + b[i];
}
}A palavra-chave __kernel define uma função de kernel OpenCL. A palavra-chave __global especifica que um ponteiro aponta para a memória global. A função get_global_id retorna o índice global do trabalho atual, que é usado para calcular os endereços de memória para os vetores de entrada e saída.
Mapeando Algoritmos para Arquiteturas de GPU
Mapear eficientemente algoritmos para a arquitetura da GPU é crucial para alcançar alto desempenho. Considerações-chave incluem:
-
Expor paralelismo suficiente: O algoritmo deve ser decomposto em muitos threads de granularidade fina que podem ser executados concorrentemente para utilizar totalmente as capacidades de processamento paralelo da GPU.
-
Minimizar divergência de ramificação: O fluxo de controle divergente dentro de uma carga de trabalho/wavefront pode levar à serialização e reduzir a eficiência do SIMD. Os algoritmos devem ser estruturados para minimizar a divergência de ramificação sempre que possível.
-
Explorar a hierarquia de memória: Acessar a memória global é caro. Os algoritmos devem maximizar o uso da memória compartilhada e dos registradores para reduzir os acessos à memória global. Os dados também devem ser organizados na memória para permitir acessos de memória coalescidos.
-
Equilibrar computação e acessos à memória: Os algoritmos devem ter uma alta proporção de operações aritméticas em relação a operações de memória para ocultar efetivamente a latência da memória e alcançar alto rendimento computacional.
-
Minimizar as transferências de dados entre host e dispositivo: Transferir dados entre a memória do host e do dispositivo é lento. Os algoritmos devem minimizar essas transferências, executando o máximo de computação possível na GPU.
Vários padrões de projeto de algoritmos paralelos são comumente usados no desenvolvimento de kernels de GPU:
-
Mapa: Cada thread executa a mesma operação em um elemento de dados diferente, permitindo o processamento paralelo simples de grandes conjuntos de dados.
-
Redução: A redução paralela é usada para calcular eficientemente um único valor (por exemplo, soma, máximo) a partir de um grande conjunto de dados de entrada.Threads perform local reductions, which are then combined to produce the final result.
-
Scan: Also known as prefix sum, scan is used to compute the running sum of elements in an array. Efficient parallel scan algorithms are key building blocks for many GPU-accelerated applications.
-
Stencil: Each thread computes a value based on neighboring data elements. Stencil computations are common in scientific simulations and image processing applications.
-
Gather/Scatter: Threads read from (gather) or write to (scatter) arbitrary locations in global memory. Careful data layout and access patterns are required for efficiency.
Conclusão
Modelos de programação GPU como CUDA e OpenCL expõem as capacidades de processamento paralelo de GPUs modernas aos desenvolvedores, permitindo-lhes acelerar uma ampla gama de aplicações. Esses modelos de programação fornecem abstrações que permitem que cargas de trabalho paralelas de grão fino sejam mapeadas de forma eficiente para o hardware da GPU.
Compreender o modelo de execução, a hierarquia de memória e os primitivos de sincronização fornecidos por esses modelos de programação é essencial para escrever código de GPU de alto desempenho. Os desenvolvedores devem considerar cuidadosamente fatores como organização de threads, divergência de ramificação, padrões de acesso à memória e design de algoritmos para aproveitar ao máximo o poder computacional das GPUs.
À medida que as arquiteturas de GPU continuam a evoluir, os modelos e ferramentas de programação também devem avançar para permitir que os desenvolvedores utilizem efetivamente novos recursos e capacidades de hardware. A pesquisa contínua em áreas como design de linguagem de programação, otimização do compilador e autotuning será crucial para melhorar a produtividade do programador e a portabilidade de desempenho na era da computação heterogênea.