Capítulo 2: Modelos de Programação 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 de programação paralela e APIs 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, ao mesmo tempo que escondem 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, arquiteturas de conjunto de instruções (ISAs) da GPU, ISAs da GPU da NVIDIA e a ISA Graphics Core Next (GCN) da AMD. Também forneceremos exemplos para ilustrar como esses conceitos são aplicados na prática.
Modelo de Execução
O modelo de execução dos modelos de programação de GPU modernos é baseado no conceito de kernels, que são funções executadas em paralelo por um grande número de threads na GPU. Ao iniciar 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 por meio de memória compartilhada e barreiras. As threads dentro de um bloco são executadas no mesmo núcleo da GPU (chamado de multiprocessador de streaming ou unidade de computação).
- Cada thread tem um ID único dentro de seu bloco e grid, 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 vários elementos de dados) quanto o paralelismo de tarefas (onde diferentes tarefas são executadas em paralelo).
FiguraAqui está a tradução em português deste arquivo markdown. Para o código, não traduzi o código, apenas comentários.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Bloco |
| | | |
Thread Thread ...Figura 2.1: Hierarquia de threads no modelo de execução GPU.
Execução SIMT
Modelos de programação de GPU como CUDA e OpenCL seguem um modelo de execução de Instrução Única, Múltiplas Threads (SIMT). No modelo SIMT, as threads são executadas em grupos chamados de warps (terminologia da NVIDIA) ou wavefronts (terminologia da AMD). Todas as threads dentro de um warp executam a mesma instrução ao mesmo tempo, mas cada thread opera em dados diferentes.
No entanto, diferentemente do modelo tradicional de Instrução Única, Múltiplos Dados (SIMD), onde todos os elementos de processamento são executados em sincronia, o SIMT permite que as threads tenham caminhos de execução independentes e divergir em instruções de ramificação. Quando um warp encontra uma instrução de ramificação, o hardware da GPU avalia a condição de ramificação para cada thread no warp. Se todas as threads seguirem o mesmo caminho (convergido), o warp continua a execução normalmente. Se algumas threads seguirem caminhos diferentes (divergido), o warp é dividido em dois ou mais subwarps, cada um seguindo um caminho diferente. O hardware da GPU serializa a execução dos caminhos divergentes, desativando as threads inativas em cada subwarp. Quando todos os caminhos forem concluídos, os subwarps reconvergerão e continuarão a execução em sincronia.
A Figura 2.2 ilustra a execução SIMT com fluxo de controle divergente.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Ramificação |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
ReconvergênciaFigura 2.2: Execução SIMT com fluxo de controle divergente.
Esse mecanismo de tratamento de divergência permite que o SIMT suporte um fluxo de controle mais flexível do que o SIMD tradicional.Traduza o arquivo Markdown, mantendo o código inalterado e traduzindo apenas os comentários.
Hierarquia de Memória
GPUs têm uma hierarquia de memória complexa para suportar os requisitos de alta largura de banda e baixa latência de cargas de trabalho paralelas. A hierarquia de memória geralmente é composta por:
- Memória global: O maior, mas mais lento espaço de memória, 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 transmitir 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 por meio 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 transbordamento de registradores e estruturas de dados grandes. A memória local geralmente é mapeada para a memória global.
A utilização efetiva da hierarquia de memória é fundamental para atingir alto desempenho em 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 da memória e os gargalos de largura de banda.
A Figura 2.3 ilustra a hierarquia de memória da GPU.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________FigAqui está a tradução em português deste arquivo markdown. Para o código, não foi traduzido o código em si, apenas os comentários.
Arquiteturas de Conjunto de Instruções de GPU
As arquiteturas de conjunto de instruções de GPU (ISAs) definem a interface de baixo nível entre o software e o hardware. Elas especificam as instruções, registradores e modos de endereçamento de memória suportados pela GPU. Compreender as ISAs de GPU é essencial para desenvolver código de GPU eficiente e otimizar o desempenho.
Nesta seção, vamos explorar as ISAs de dois principais fornecedores de GPU: NVIDIA e AMD. Nos concentraremos nas ISAs Parallel Thread Execution (PTX) e SASS da NVIDIA, e na ISA Graphics Core Next (GCN) da AMD.
ISAs de GPU da NVIDIA
As GPUs da NVIDIA suportam dois níveis de ISAs: PTX (Parallel Thread Execution) e SASS (Streaming ASSembler). O PTX é uma ISA virtual que fornece um destino estável para os compiladores CUDA, enquanto o SASS é a ISA nativa das GPUs da NVIDIA.
PTX (Parallel Thread Execution)
O PTX é uma ISA virtual de baixo nível projetada para as GPUs da NVIDIA. É semelhante ao LLVM IR ou ao bytecode Java, pois fornece um destino estável e independente da arquitetura para os compiladores. Os programas CUDA são normalmente compilados para o código PTX, que é então traduzido para as instruções SASS nativas pelo driver da GPU da NVIDIA.
O PTX suporta uma ampla gama de instruções aritméticas, de memória e de fluxo de controle. Ele tem um número ilimitado de registradores virtuais e suporta predição, o que permite uma implementação eficiente do fluxo de controle. O PTX também fornece instruções especiais para sincronização de threads, operações atômicas e amostragem de texturas.
Aqui está um exemplo de código PTX para um kernel simples de adição de vetores:
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
# Carrega os parâmetros da função
ld.param.u64 %rd1, [vecAdd_param_0];
ld.param.u64 %rd2, [vecAdd_param_1];
ld.param.u64 %rd3, [vecAdd_param_2];
ld.param.u32 %r1, [vecAdd_param_3];
cvta.to.global.u64 %rd4, %rd1;
cvta
```Aqui está a tradução em português do arquivo Markdown, com os comentários traduzidos, mas o código não traduzido:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;
ret; }
Este código PTX define uma função de kernel `vecAdd` que recebe quatro parâmetros: ponteiros para os vetores de entrada e saída, e o tamanho dos vetores. O kernel calcula o ID do thread global, carrega os elementos correspondentes dos vetores de entrada, realiza a adição e armazena o resultado no vetor de saída.
#### SASS (Streaming ASSembler)
SASS é o ISA nativo dos GPUs da NVIDIA. É um ISA de baixo nível e específico da máquina, que mapeia diretamente para o hardware do GPU. As instruções SASS são geradas pelo driver do GPU da NVIDIA a partir do código PTX e geralmente não são visíveis para os programadores.
As instruções SASS são codificadas em um formato compacto para reduzir o consumo de largura de banda de memória e o uso do cache de instruções. Elas suportam uma ampla gama de tipos de operandos, incluindo registradores, valores imediatos e vários modos de endereçamento para acesso à memória.
Aqui está um exemplo de código SASS para o kernel de adição de vetores:
```sass
code_version_number 90
@P0 LDG.E R2, [R8];
@P0 LDG.E R0, [R10];
@P0 FADD R0, R2, R0;
@P0 STG.E [R12], R0;
EXIT;Este código SASS corresponde ao código PTX mostrado anteriormente. Ele carrega os elementos do vetor de entrada da memória global (LDG.E), realiza a adição (FADD), armazena o resultado de volta na memória global (STG.E) e sai do kernel.
ISA do AMD Graphics Core Next
Os GPUs da AMD usam a arquitetura e o ISA Graphics Core Next (GCN). O GCN é um ISA baseado em RISC que suporta cargas de trabalho gráficas e de computação. Ele é projetado para alto desempenho, escalabilidade e eficiência energética.
O GCN introduz várias características-chave, como:
- Um Aqui está a tradução em português do arquivo markdown, com os comentários traduzidos, mas mantendo o código inalterado:
Arquitetura ALU eficiente para execução de operações escalares e controle de fluxo.
- Uma ALU vetorial para execução paralela de operações data-parallel.
- Um sistema de memória de alta largura de banda com suporte para operações atômicas e acesso de baixa latência à memória compartilhada.
- Um modo de endereçamento flexível para operações de memória, suportando endereçamento base+offset e escalar+vetor.
Aqui está um exemplo de código da ISA GCN para um kernel de adição de vetores:
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
# Carrega os argumentos do kernel da memória
s_load_dwordx4 s[0:3], s[4:5], 0x0
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
# Calcula o endereço da memória para o primeiro elemento do vetor de entrada
v_lshlrev_b32 v0, 2, v0
v_add_u32 v1, vcc, s1, v0
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc
flat_load_dword v1, v[1:2]
# Calcula o endereço da memória para o segundo elemento do vetor de entrada
v_add_u32 v3, vcc, s0, v0
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc
flat_load_dword v0, v[3:4]
# Executa a adição dos elementos do vetor
v_add_f32 v0, v0, v1
flat_store_dword v[3:4], v0
s_endpgmEste código GCN carrega os elementos dos vetores de entrada usando flat_load_dword, executa a adição usando v_add_f32 e armazena o resultado de volta na memória usando flat_store_dword. As instruções s_load_dwordx4 e s_load_dword são usadas para carregar os argumentos do kernel da memória.
Mapeando Algoritmos para Arquiteturas GPU
Mapear eficientemente algoritmos para a arquitetura GPU é crucial para alcançar alto desempenho. Considerações-chave incluem:
Expondo Paralelismo Suficiente
O algoritmo deve ser decomposto em muitos threads de grão fino que possam ser executados concorrentemente para utilizar totalmente as capacidades de processamento paralelo da GPU. Isso muitas vezes envolve identificar partes data-parallel do algoritmo que podem ser executadas independentemente em diferentes elementos de dados.
Minimizando Divergência de Ramificação
O fluxo de controle divergente dentro de uma warp/wavefront pode levar a serialização e eficiência SIMD reduzida. Os algoritmos devem ser estruturados para minimizar a divergência de ramificação sempre que possível. Isso pode ser alcançado reduzindo o uso de fluxo de controle dependente de dados.### Explorando a Hierarquia de Memória
Acessar a memória global é caro. Os algoritmos devem maximizar o uso de memória compartilhada e registradores para reduzir os acessos à memória global. Os dados também devem ser organizados na memória para permitir acessos coalesced à memória, onde os threads em uma warp acessam locais de memória contíguos. O uso eficaz da hierarquia de memória pode reduzir significativamente a latência da memória e os gargalos de largura de banda.
Balanceando 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 um alto desempenho computacional. Isso pode ser alcançado maximizando a reutilização de dados, pré-buscando dados e sobrepondo computação com acessos à memória.
Minimizando Transferências de Dados entre Host e Dispositivo
Transferir dados entre a memória do host (CPU) e do dispositivo (GPU) é lento. Os algoritmos devem minimizar essas transferências, realizando o máximo de computação possível na GPU. Os dados devem ser transferidos para a GPU em lotes grandes e mantidos no dispositivo pelo maior tempo necessário, para amortizar o overhead da transferência.
Vários padrões de design de algoritmos paralelos são comumente usados no desenvolvimento de kernels GPU:
-
Map: Cada thread executa a mesma operação em um elemento de dados diferente, permitindo o processamento paralelo simples de grandes conjuntos de dados.
-
Reduce: 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. Os threads realizam reduções locais, que são então combinadas para produzir o resultado final.
-
Scan: Também conhecido como prefixo de soma, o scan é usado para calcular a soma acumulada dos elementos em um array. Algoritmos eficientes de scan paralelo são blocos de construção essenciais para muitas aplicações aceleradas por GPU.
-
Stencil: Cada thread calcula um valor com base em elementos de dados vizinhos. Os cálculos de stencil são comuns em simulações científicas e aplicações de processamento de imagem.Aqui está a tradução em português do arquivo Markdown fornecido, com os comentários traduzidos, mas sem a tradução do código:
-
Gather/Scatter: As threads leem de (gather) ou gravam em (scatter) locais arbitrários na memória global. Um layout de dados e padrões de acesso cuidadosos são necessários para eficiência.
A Figura 3.20 ilustra um exemplo do padrão de mapeamento, onde cada thread aplica uma função (por exemplo, raiz quadrada) a um elemento diferente da matriz de entrada.
Matriz de Entrada:
| | | | | | | |
v v v v v v v v
______________________________
Threads: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Matriz de Saída: Figura 3.20: Exemplo do padrão de mapeamento na programação de GPU.
Conclusão
Modelos de programação de GPU, como CUDA e OpenCL, expõem os recursos de processamento paralelo das 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.
Entender 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 de programação e as ferramentas também devem avançar para permitir que os desenvolvedores utilizem de forma eficaz os novos recursos e capacidades do hardware. A pesquisa contínua em áreas como design de linguagem de programação, otimização de compiladores e autoajuste serão cruciais para melhorar a produtividade do programador e a portabilidade de desempenho na era da computação heterogênea.