Você clica “run” e o kernel CUDA executa em 10 microssegundos. Mas o que acontece entre o seu código e os 128 SMs da GPU?

A maioria dos devs que trabalha com CUDA trata o <<>> como mágica. Você declara as dimensões do grid, chama o kernel, e de alguma forma milhares de threads processam dados em paralelo. Funciona. Mas por baixo dessa sintaxe amigável existe uma orquestra de compiladores, drivers, barramentos PCIe e schedulers de hardware que transformam seu código C++ em instruções executadas por 16.384 cores simultaneamente.

Eu passei anos escrevendo kernels CUDA sem realmente entender essa jornada. Quando finalmente mergulhei nos detalhes — graças a uma análise brilhante de Fergus Finn usando profilers e reverse engineering do driver — percebi que estava otimizando código às cegas.

Vamos desmontar essa caixa preta.

O pipeline de compilação que ninguém vê

Quando você compila um arquivo .cu com nvcc, o que parece um único compilador na verdade é um orquestrador que chama pelo menos três compiladores diferentes:

O PTX (Parallel Thread Execution) é uma espécie de “bytecode” da NVIDIA — uma ISA virtual com registradores tipados e ilimitados. Ele existe para garantir compatibilidade entre gerações de GPU. Seu código PTX compilado para uma RTX 3090 pode ser recompilado just-in-time para uma RTX 5090, porque a GPU recebe o PTX comprimido no fatbinary e gera SASS nativo na hora.

O SASS (Streaming Assembler) é onde a mágica acontece de verdade. Cada instrução SASS tem 128 bits, e os 21 bits de controle que acompanham cada instrução contêm informações críticas:

• Stall counts: quantos ciclos o scheduler deve esperar antes de tentar executar a próxima instrução
• Yield hints: se o warp deve ceder prioridade para outros warps
• Scoreboard barriers: 6 barreiras físicas (0-5) por warp para operações de latência variável

Esses bits de controle são decididos pelo ptxas em tempo de compilação. É ele, não o hardware, quem define grande parte do scheduling de instruções. Quando você vê alguém falar que “o GPU scheduler é muito sofisticado”, a verdade é que metade da inteligência está embutida nos bits de controle gerados pelo compilador.

Do host à GPU: uma dança de ioctls

Quando seu código host chama vadd<<>>(d_a, d_b, d_c, n), o compilador já transformou isso em algo bem diferente do que parece. Nos bastidores, o que acontece é:

// O que o compilador gera (simplificado):
__cudaPushCallConfiguration(dim3(4096), dim3(256), 0, 0);
// Pack argumentos em offsets específicos: 0x160, 0x168, 0x170, 0x178
__cudaLaunch(&vadd_stub);

Os argumentos do kernel são empacotados em um buffer host com offsets que correspondem às posições dos constant banks da GPU. Não é um detalhe aleatório — esses offsets determinam onde cada parâmetro vai parar na memória constante do SM.

A partir daqui, o controle passa para a libcuda.so.1, o driver closed-source da NVIDIA. E esse driver é… prolixo. Durante a inicialização, ele faz 948 chamadas ioctl nos device files /dev/nvidiactl e /dev/nvidia-uvm. Novecentas e quarenta e oito. Para configurar o contexto GPU, alocar memória, estabelecer canais de comunicação.

Cada vez que eu reclamo que CUDA “demora para inicializar”, agora eu entendo por quê.

Pushbuffer, GPFIFO e o doorbell: o mecanismo de submissão

O driver não fala diretamente com os SMs da GPU. Existe uma camada intermediária de hardware chamada Host Engine que consome comandos de uma estrutura chamada pushbuffer.

O pushbuffer é uma região de memória no host onde o driver escreve GPU methods — pares de (endereço de registrador, valor). Pense neles como comandos remotos: “escreva este valor neste registrador da GPU”.

Para coordenar a produção (driver) e o consumo (GPU) desses comandos, existe um ring buffer chamado GPFIFO com dois cursores:

• GP_PUT: onde o driver está escrevendo (avança no host)
• GP_GET: onde a GPU está lendo (avança no device)

Ambos vivem no USERD, uma pequena estrutura por canal na memória do device.

O momento mais interessante é o doorbell ring. Depois de escrever os comandos no pushbuffer e atualizar o GP_PUT, o driver escreve um token em um registrador mapeado em memória (memory-mapped I/O). Esse write é o sinal para o Host Engine da GPU: “tem trabalho novo, pode consumir”.

É um mecanismo elegante. Zero polling do lado da GPU. Um único write no barramento PCIe acorda o engine inteiro.

QMD: o “cartão de embarque” do kernel

Antes de disparar o kernel, o driver monta uma estrutura de 256 bytes chamada QMD (Queue Meta Data). Ela é o cartão de embarque do kernel — contém tudo que a GPU precisa para executar:

QMD (256 bytes):
├── Grid dimensions: 4096 blocos
├── Block size: 256 threads
├── Registradores por thread: 16
├── Entry point do programa
├── Endereços dos constant buffers
├── Shared memory size
└── Campos de completion semaphore

O QMD é transmitido inline no pushbuffer através de 64 words de LOAD_INLINE_QMD_DATA, precedidos por SET_INLINE_QMD_ADDRESS_A/B. Nada de DMA separado — os 256 bytes vão direto no fluxo de comandos.

Quando você otimiza registradores ou shared memory no seu kernel, está diretamente afetando os valores nesse QMD. E esses valores determinam quantos blocos podem rodar simultaneamente em cada SM.

Compute Work Distributor: o dispatcher dos SMs

Com o QMD em mãos, o Compute Work Distributor (CWD) assume. Ele é o componente que decide como distribuir 4096 blocos entre os 128 SMs de uma RTX 4090.

A distribuição não é trivial. O CWD precisa respeitar resource constraints para determinar a ocupância máxima:

Limitação por registradores:

Cada SM tem 65.536 registradores de 32 bits. Cada bloco precisa de 256 threads × 16 registradores = 4.096 registradores. Resultado: 65.536 ÷ 4.096 = 16 blocos caberiam pela perspectiva de registradores.

Limitação por threads:

Cada SM suporta no máximo 1.536 threads ativos. Com blocos de 256 threads: 1.536 ÷ 256 = 6 blocos residentes.

Resultado final: O gargalo é o limite de threads — 6 blocos por SM, totalizando 6 × 128 = 768 blocos residentes simultaneamente. Os outros 3.328 blocos ficam em uma fila esperando slots livres.

Com 4096 blocos divididos entre 128 SMs, cada SM processa 32 blocos no total, mas apenas 6 de cada vez. Os blocos restantes são enfileirados e despachados conforme outros terminam.

Warps: a unidade real de execução

Aqui mora uma confusão comum. Você programa em threads. Você configura blocks. Mas a GPU executa warps — grupos de 32 threads que executam a mesma instrução simultaneamente no modelo SIMT (Single Instruction, Multiple Thread).

Cada bloco de 256 threads contém 8 warps. Com 6 blocos residentes por SM, temos 48 warps ativos por SM.

Esses 48 warps são distribuídos entre 4 sub-partições do SM, cada uma com seu próprio scheduler. Cada scheduler gerencia 12 warps e pode emitir uma instrução por ciclo através de 32 lanes físicas.

A beleza do design está na latency hiding. Enquanto um warp espera dados da memória (centenas de ciclos), o scheduler simplesmente executa outro warp elegível. Com 12 warps competindo por vez em cada sub-partição, quase sempre existe um warp pronto para executar.

Eligibilidade de instruções: scoreboard e stall counts

Como o scheduler decide qual warp executar em cada ciclo? Dois mecanismos principais:

1. Stall counts (latência fixa):

Para operações com latência conhecida (soma inteira, multiplicação float), o compilador ptxas codifica um contador de stall nos bits de controle da instrução. O scheduler simplesmente decrementa o contador e sabe exatamente quando o resultado estará pronto.

2. Scoreboard barriers (latência variável):

Para operações como loads de memória global — onde a latência depende de cache hits e contenção no barramento — o compilador atribui uma das 6 barreiras físicas. A instrução dependente espera até que a barreira seja liberada pelo hardware.

Veja um exemplo real do kernel vadd (soma de vetores):

// Instrução SASS com bits de controle
LDG.E R4, [R2]       // stall=4, yield=yes, set barrier B2
LDG.E R5, [R6]       // stall=4, yield=yes, set barrier B3
FADD R7, R4, R5      // stall=5, yield=no,  wait B2, B3
STG.E [R8], R7       // stall=1, yield=yes

O LDG.E (load global) seta a barreira B2 com stall mínimo de 4 e yield=yes — sinalizando que este warp pode ceder a vez. O FADD espera as barreiras B2 e B3 serem liberadas (ambos os loads completarem) antes de executar. O STG.E (store global) tem stall de apenas 1, porque não precisamos esperar o resultado — é fire-and-forget.

Esses bits de controle são a razão pela qual ptxas é tão importante. Um compilador ruim aqui significa warps travando desnecessariamente, SMs ociosos, performance destruída.

Coalescing de memória: de 32 acessos a 4 transações

Quando 32 threads de um warp executam LDG.E simultaneamente, o hardware tenta coalescer esses acessos em transações de memória otimizadas.

No caso ideal do vadd, as 32 threads acessam 32 floats consecutivos (128 bytes). A unidade de load/store agrupa isso em 4 transações de 32 bytes — perfect coalescing, zero bytes desperdiçados.

Se os acessos fossem espalhados (stride patterns, random access), cada thread poderia gerar sua própria transação de 32 bytes, desperdiçando 28 dos 32 bytes transferidos. A diferença é de 8x no throughput de memória efetivo.

A hierarquia de cache completa:

Thread → L1 Data Cache (por SM, ~128 KB)
       → L2 Cache (compartilhado, 72 MB na RTX 4090)
       → GDDR6X VRAM (24 GB, ~1 TB/s de bandwidth)

Os números que importam

O profiler ncu (Nsight Compute) revela os números reais para o kernel vadd com 1 milhão de elementos em uma RTX 4090:

O issue rate de apenas 5,17% grita “bandwidth-bound”. O kernel faz uma única FADD para cada 12 bytes transferidos — praticamente zero compute para muita memória. O DRAM throughput de quase 80% mostra que estamos saturando a banda de memória, que é exatamente o gargalo.

Tráfego real: 8,4 MB lidos em 10,78 μs ≈ 780 GB/s efetivos. A banda teórica da RTX 4090 é ~1 TB/s, então estamos aproveitando ~78% dela. Para um kernel tão simples, é um número excelente.

O que acontece quando o kernel termina

O último bloco a completar dispara um completion semaphore — campos 23-24 do QMD. O Copy Engine do device lê esse semáforo antes de iniciar a transferência DMA de volta para o host.

Um detalhe que poucos sabem: como o output de 1M de floats (4 MB) cabe inteiro no L2 cache de 72 MB da RTX 4090, o Copy Engine serve os dados diretamente do L2, sem precisar ir até a DRAM. Zero round-trips extras.

No host, cudaMemcpy espera um segundo semáforo postado pelo Copy Engine. Quando ele chega, os dados já estão na memória do host. O printf no seu código C++ imprime o resultado.

10,78 microssegundos. Dez. Do launch do kernel até o resultado na memória da GPU.

CUDA-Oxide: o futuro é Rust?

Uma novidade de 2026 que merece destaque: a NVIDIA lançou o CUDA-Oxide, um backend experimental que compila Rust diretamente para PTX. Sem passar por C++, sem nvcc.

A ideia é trazer as garantias de segurança de memória do Rust para programação GPU. Se funcionar como prometido, problemas como race conditions em shared memory e buffer overflows em kernels complexos ficam mais fáceis de detectar em tempo de compilação.

Ainda é experimental, mas o sinal é claro: o futuro da programação GPU não é necessariamente C++.

Por que isso importa para o seu código

Entender esse pipeline muda a forma como você escreve kernels:

1. Registradores são ouro. Cada registrador a mais por thread reduz o número de blocos residentes por SM. Um kernel com 32 registradores/thread pode ter metade da ocupância de um com 16. Use __launch_bounds__ para controlar isso.

2. Coalescing não é opcional. A diferença entre acesso coalescido e não-coalescido é 8x em throughput. Estruture seus dados para acessos sequenciais dentro de cada warp.

3. O compilador faz scheduling. Aqueles bits de controle no SASS determinam como warps são agendados. ptxas precisa de kernels bem escritos para gerar bons stall counts. Divergência de branches dentro de um warp confunde o scheduler e destroi a performance.

4. O L2 cache de 72 MB é seu aliado. Dados que cabem no L2 são servidos sem tocar a DRAM. Para operações sobre datasets pequenos, o throughput efetivo pode ser muito maior que a banda de memória nominal.

Se você programa CUDA e quer realmente otimizar seus kernels, o ncu (Nsight Compute) é obrigatório. Ele mostra exatamente onde estão os gargalos — registradores, bandwidth, ocupância, stalls. Chutar valores de blockDim e gridDim sem profiling é como dirigir com os olhos vendados num autódromo.

—

Fonte de inspiração: What happens when you run a CUDA kernel? por Fergus Finn