O que acontece internamente quando você executa um kernel CUDA

• Mesmo um programa CUDA simples de soma de vetores passa por pipeline de compilação, chamadas de driver, fila de comandos da GPU, escalonamento de warps, hierarquia de memória e semáforos de conclusão até produzir o resultado 2.000000
• O nvcc separa código de host e de device, gera PTX com cicc, SASS com ptxas, empacota cubin e PTX em um fatbin e o insere no executável Linux
• A sintaxe de launch vadd<<<4096, 256>>> é transformada em um host launch stub, e os argumentos da, db, dc, n são passados ao driver via runtime CUDA e libcuda.so.1
• A execução na GPU começa com QMD, pushbuffer, GPFIFO, GP_PUT e uma escrita MMIO no doorbell, enquanto os 128 SMs da RTX 4090 executam a configuração de 4096 blocos e 256 threads em unidades de warp
• Esse kernel é limitado por largura de banda de memória por causa da baixa intensidade aritmética, exigindo 12 bytes de transferência por soma de float; no Nsight Compute ele mostra 10.78μs, 79.65% do pico de DRAM e 5.17% de warp issue

---

Kernel de exemplo e escopo da observação

• O programa de exemplo usa o kernel CUDA vadd para somar dois arrays de float e armazenar o resultado em um terceiro array
  • n = 1 << 20, processando 1.048.576 floats
  • A configuração de launch é vadd<<<4096, 256>>>(da, db, dc, n), usando 4096 * 256 = n threads
• Compilando e executando para uma RTX 4090 com nvcc -arch=sm_89, a saída é c[0]=2.000000 c[n-1]=2.000000
• Mesmo essa única linha de resultado envolve dezenas de milhões de instruções de CPU, arquivos de device, cerca de 900 ioctls e registradores doorbell mapeados em memória

Como o nvcc cria o executável

• Com nvcc --keep, é possível inspecionar diretamente os artefatos do pipeline de compilação
  • vadd.ptx: PTX do código de device gerado por cicc
  • vadd.sm_89.cubin: SASS do código de device gerado por ptxas
  • vadd.fatbin: fatbin que agrupa cubin e PTX
  • vadd.cudafe1.stub.c: host launch stub e código de registro do kernel
  • vadd.o: objeto final de host contendo o fatbin
• O código de host é processado pelo compilador de host, enquanto o kernel de device vadd passa pelas etapas cicc e ptxas
• PTX é uma ISA virtual, usa registradores virtuais infinitos com tipo e não reflete diretamente a quantidade real de registradores do hardware
  • O PTX do exemplo inclui o cálculo de blockIdx.x * blockDim.x + threadIdx.x, checagem de limite, global load, soma de float e global store
  • Ponteiros CUDA são, por padrão, generic pointers, então cvta.to.global os converte para endereços globais antes do uso de ld.global
  • mul.wide.s32 transforma o índice em um offset de 4 bytes, equivalente a sizeof(float), e o expande de 32 para 64 bits
• SASS são instruções reais específicas da arquitetura, e no output para RTX 4090 aparece de forma mais compacta que o PTX
  • S2R copia registradores especiais como SR_CTAID.X e SR_TID.X para registradores gerais
  • A combinação de mul.wide e add em PTX é fundida em IMAD.WIDE no SASS
  • A conversão cvta é absorvida no processo de endereçamento
• O operando c[0x0][...] aponta para o constant bank 0 gerenciado pelo driver
  • Os ponteiros a, b, c ficam em 0x160, 0x168, 0x170
  • n fica em 0x178
  • Geometria de launch, como blockDim.x, e valores de ABI também ficam no mesmo bank
• O cubin é um arquivo ELF, o mesmo formato de contêiner de executáveis Linux
  • O fatbinary empacota cubin e PTX juntos
  • Nesta RTX 4090, o SASS é o que de fato executa, mas o PTX é incluído como fallback para JIT em outras arquiteturas
  • Como o PTX é texto puro verboso, o nvcc o comprime por padrão

Como o código de host prepara o launch

• O frontend do compilador cudafe++ insere um constructor oculto executado antes de main
  • Esse constructor registra o fatbinary embutido no runtime CUDA
  • Ele conecta o ponteiro de função de host vadd ao nome mangled do kernel de device dentro do fatbin
• A sintaxe vadd<<<4096, 256>>>(da, db, dc, n) é transformada no host launch stub gerado
  • da, db, dc, n são gravados com alinhamento em offsets 0, 8, 16, 24 em um buffer de argumentos na memória do host
  • Esses offsets correspondem às posições 0x160, 0x168, 0x170, 0x178 que o SASS lê no constant bank 0
• O stub chama __cudaLaunch, passando o endereço da função dummy vadd do lado do host
  • Esse endereço não é o de uma função a ser executada pela CPU, mas sim uma chave para consultar a tabela de registro do runtime
  • O runtime localiza o nome simbólico de device correspondente e então passa para o driver user-mode de código fechado libcuda.so.1
• Na primeira chamada à GPU, o runtime CUDA abre dinamicamente libcuda.so.1 e cria um contexto
  • Em strace, é possível ver a abertura de /lib/x86_64-linux-gnu/libcuda.so.1
  • O contexto inclui um channel pelo qual a CPU se comunica com a GPU
• Desde o CUDA 12.2, o carregamento de módulos é lazy por padrão
  • O upload do cubin SASS é adiado até o primeiro launch de um kernel específico
  • Isso pode ser controlado com CUDA_MODULE_LOADING

A fila de comandos que entrega trabalho à GPU

• A GPU não recebe chamadas de função e não faz jump para um entry point como a CPU
  • Ela lê, pela PCIe, um command stream do driver armazenado na memória do host
  • cuLaunchKernel coloca o comando de launch completo nesse stream e notifica a GPU
• Na primeira execução, o driver copia o SASS do kernel para a memória da GPU
  • Ele aloca um code buffer e copia o SASS para lá
• O channel contém duas estruturas principais na RAM do host
  • pushbuffer: região de memória onde o driver escreve os methods que são comandos da GPU
  • GPFIFO: ring buffer de ponteiros para trechos do pushbuffer
• Uma entrada do GPFIFO consiste em duas words de 32 bits que representam (base, length) de um trecho do pushbuffer
• GPU e driver acompanham consumo e produção com dois cursores
  • GP_GET: até onde a GPU já consumiu
  • GP_PUT: até onde o driver já produziu
  • Ambos ficam em uma estrutura por channel chamada USERD
• Ao lançar um kernel, o driver escreve methods em um trecho do pushbuffer, faz uma entrada do GPFIFO apontar para ele e avança GP_PUT
• Em GPUs modernas, o host engine não monitora o cursor continuamente, então é preciso um doorbell
  • A GPU mapeia para o processo uma pequena janela de registradores
  • O driver escreve no registrador doorbell o token de submissão de trabalho do channel
  • Ao receber o doorbell, o host engine lê GP_PUT e busca por DMA a entrada do GPFIFO e o trecho correspondente do pushbuffer

As informações de execução no QMD

• O launch começa com um burst de methods SET_INLINE_QMD_ADDRESS_A/B e LOAD_INLINE_QMD_DATA
• QMD (Queue Meta Data) é o descritor de launch do grid de compute
  • Inclui os tamanhos de grid e bloco, 4096 e 256
  • Inclui número de registradores por thread e exigência de memória compartilhada
  • Inclui o endereço inicial do programa e o endereço do constant bank com os argumentos do kernel
  • Também inclui onde sinalizar a conclusão
• Os argumentos empacotados pelo host stub são copiados pelo driver para o constant bank, e o endereço desse bank é gravado no QMD
• O QMD informa à GPU onde está o SASS, como configurar o programa paralelo e onde emitir o sinal de conclusão
• cuLaunchKernel retorna no momento em que o doorbell toca
  • Como a chamada é assíncrona, a CPU pode continuar executando enquanto o trabalho da GPU avança

SM, warps e ocupação

• O host engine entrega o QMD ao compute work distributor
  • Esse componente existe uma vez por GPU inteira
  • Ele distribui o fluxo linear de instruções SASS entre os SMs para execução como programa paralelo
• A GPU alvo, a GeForce RTX 4090, usa 128 SMs
  • O launch consiste em 4096 blocos com 256 threads por bloco
• Cada SM tem cache local de instruções, e cada warp ativo mantém um program counter
  • Desde Volta, existe o modelo Independent Thread Scheduling, com program counter e call stack por thread
  • Mesmo assim, o issue ainda acontece por warp
• Nesse kernel de exemplo, o limite de recursos determina a residência de blocos
  • 256 threads = 8 warps por bloco
  • O ptxas reserva 16 registradores por thread
  • Pelo critério de registradores, seriam possíveis 16 blocos por SM
  • Mas a capacidade de threads é de 1.536 threads ativas por SM, então 1536 / 256 = 6 blocos no máximo
  • Portanto, o limite real é de 6 blocos por SM, ou 48 warps residentes
• O SM é dividido em 4 processing blocks, ou sub-partitions
  • Os 48 warps residentes são distribuídos igualmente entre as 4 sub-partitions
  • Cada warp scheduler gerencia 12 warps ativos quando está cheio
  • A cada ciclo, ele escolhe um warp elegível e despacha a próxima instrução para 32 lanes

Quando um warp se torna elegível

• A GPU não extrai grandes quantidades de dependência dinâmica dentro de uma thread única como uma CPU com execução out-of-order
  • Em vez disso, mantém muitos warps residentes e, quando um estagna, troca para outro para esconder latência
  • O compilador agenda o que tem timing previsível, e o scoreboard de hardware resolve o que não é previsível
• Cada instrução SASS de 128 bits contém um control-code payload escrito pelo ptxas
  • Instruções de latência fixa carregam uma contagem estática de stall
  • A dica de yield informa se deve ceder prioridade ao scheduler
  • Operações de latência variável usam 6 barreiras físicas de scoreboard por warp
• No trecho SASS de exemplo, os dois LDG.E configuram a mesma barreira de scoreboard B2
  • O FADD espera por B2
  • Até os dois loads retornarem e a barreira ser liberada, o warp fica inelegível
  • Nesse intervalo, o scheduler escolhe outros warps da mesma sub-partition
• O trecho de FADD até STG.E é tratado como latência fixa
  • FADD tem stall=5, e o warp fica estacionado por alguns ciclos até o resultado em R9 estar pronto
  • Não é necessária uma barreira separada
• Esse payload de controle fica oculto no output padrão do nvdisasm
  • Ele aparece na codificação crua de 128 bits do cuobjdump -sass, na segunda word de 64 bits
  • O layout não é documentado oficialmente; foi reconstruído por microbenchmarking

Acesso à memória e medição de desempenho

• Quando um warp executa LDG.E, as 32 threads calculam seus endereços individualmente
  • No exemplo, o acesso é a um array de floats consecutivos, então o warp inteiro solicita um bloco contínuo de 32 * 4 = 128 bytes
• A unidade de load/store do SM faz request coalescing
  • Ela combina 32 requisições de 4 bytes em 4 requisições de setor de 32 bytes
  • Se o acesso não fosse consecutivo, poderia ser necessário ler mais dados do que o estritamente necessário
• A requisição coalescida verifica primeiro o L1 Data Cache local do SM
  • Em caso de miss, segue pela interconexão crossbar até slices do L2 Cache de 72MB
  • Se também houver miss no L2, ela vai ao controlador de memória e ao barramento até a VRAM GDDR6X
• O store STG.E segue, em princípio, o mesmo caminho no sentido inverso
• As medições do Nsight Compute mostram que esse kernel é memory-bound
  • launch__grid_size: 4.096
  • launch__block_size: 256
  • launch__registers_per_thread: 16
  • launch__waves_per_multiprocessor: 5.33
  • sm__warps_active.avg.pct_of_peak: 82.77%
  • smsp__issue_active.avg.pct_of_peak: 5.17%
  • dram__throughput.avg.pct_of_peak: 79.65%
  • gpu__time_duration.sum: 10.78μs
• O kernel tem intensidade aritmética muito baixa
  • Ele faz 1 soma de float para 12 bytes transferidos: dois loads de 4 bytes e um store de 4 bytes
  • Pelo lado de leitura da DRAM, são 8,4MB lidos em 10,78μs, cerca de 780GB/s, algo em torno de 4/5 do pico
  • A saída c de 4MB cabe no L2 de 72MB, então não é descarregada para DRAM até que a cópia device-to-host precise lê-la

Como o resultado volta para a CPU

• O launch do kernel retorna à CPU quando o doorbell é acionado, então a GPU precisa sinalizar a conclusão separadamente
• Quando todos os 4096 blocos se aposentam, a GPU publica o completion semaphore contido no QMD
  • O campo de fence do QMD está nas words 23–24
• No default stream, cudaMemcpy(c, dc, ...) fica enfileirado após o kernel
  • O copy engine da GPU permanece bloqueado até o semaphore ser acionado
  • Como c ainda está dirty no L2 de 72MB, a leitura do copy engine é atendida a partir do L2, sem ida e volta à DRAM
  • Os dados atravessam a PCIe e vão para a memória do host
• Quando a cópia termina, o copy engine publica seu próprio semaphore
  • A espera do cudaMemcpy no host termina
  • c volta a ser memória comum do host
  • printf lê c[0] e c[n-1] da RAM e imprime em stdout

Como olhar dentro do launch

• Ler apenas os open kernel modules não basta, porque libcuda é de código fechado e parte do comportamento não pode ser verificada diretamente
• Escritas de method não passam por syscall; elas vão direto para um buffer write-combined já mapeado, então para ver o pushbuffer é preciso ler memória
• Um shim com LD_PRELOAD pode interceptar mmap e registrar as regiões mapeadas a partir de /dev/nvidia*
  • Se o programa de teste chamar a função de dump do shim logo após o launch, é possível imprimir o pushbuffer mapeado
  • O dump procura o burst de methods correspondente a SET_INLINE_QMD_ADDRESS_A
• O cabeçalho de um method no pushbuffer codifica opcode, quantidade de payload, índice de subchannel e offset de registrador em campos de bits
  • 0x0318 é SET_INLINE_QMD_ADDRESS_A
  • 0x0320 + i * 4 é LOAD_INLINE_QMD_DATA(i)
  • No dump aparece um burst de increasing-method com count 66, contendo 2 words de endereço e 64 words de QMD, ou seja, 256 bytes de QMD inline
  • Dentro do QMD, a word 12 é 0x1000 e a word 18 é 0x100, correspondendo a 4096 e 256 do launch
• A configuração do driver ocorre via ioctl
  • Em um programa com um único kernel, o strace registra 948 ioctls
  • A maioria é setup executado uma única vez
  • Os principais file descriptors são /dev/nvidiactl e /dev/nvidia-uvm
  • O byte mágico dos ioctls do NVIDIA resource manager é 0x46, ou 'F'
  • O número de comando 0x2A é interpretado como NV_ESC_RM_CONTROL, e 0x2B como NV_ESC_RM_ALLOC
• Em vadd.cudafe1.stub.c, gerado por nvcc --keep, também dá para ver o código de registro na inicialização
  • Uma função com __attribute__((__constructor__)) roda antes de main
  • __cudaRegisterBinary e __cudaRegisterEntry conectam o ponteiro de função de host vadd ao entry point de device _Z4vaddPKfS0_Pfi

[GN⁺]

Comentários do Hacker News

• Foi um texto interessante, e a explicação sobre o semáforo do stream padrão também foi divertida.
  Gosto do fato de o CUDA lidar implicitamente com a sincronização de comandos e permitir o uso seletivo de comandos paralelos por meio de streams.
  Isso contrasta com o Vulkan, que joga toda a complexidade de sincronização para o usuário desde o início.

• Do lado do hardware, há alguma documentação pública.
  Não é preciso necessariamente ler o código-fonte do kernel para encontrar a documentação de métodos ou o formato QMD.
  Veja https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

• Foi muito útil.
  Em especial, a parte sobre doorbell e QMD foi a mais útil por mostrar como a sintaxe de execução do CUDA se conecta ao que de fato é enviado à GPU.
  A maioria das explicações para perto de kernels, blocos e warps, mas este texto torna o caminho CPU→driver→GPU muito mais fácil de acompanhar.

• O código de controle é um pouco mais complexo do que o descrito no texto.
  Na prática, é mais parecido com uma consulta a tabela do que com bits dentro de uma palavra de controle.

• Hoje há empresas cujo trabalho principal é otimizar kernels para executá-los mais rápido.
  Fico imaginando se um dia essas empresas serão superadas por uma biblioteca open source que faça isso muito bem.
  A Nvidia parece ser capaz de lançar algo assim a qualquer momento.
  Ou talvez isso acabe dando ainda mais certo se grandes provedores adquirirem essas empresas para usar maior velocidade de inferência como moat.

  • No curto prazo, aquisições do tipo acqui-hire parecem bem plausíveis.
    Ainda assim, vendo o avanço dos modelos em benchmarks relacionados, como o kernelbench, acho que soluções mais generalizadas inevitavelmente também vão aparecer.
    O problema é que, a cada nova geração de hardware, surgem com frequência restrições ou recursos que os modelos existentes nunca viram.
    Por exemplo, o tcgen05 do Blackwell já foi, em certo momento, um caso fora da distribuição.
    Se os modelos começarem a generalizar melhor, talvez isso não seja uma barreira fatal, mas por enquanto ainda é um obstáculo.
    [1] https://kernelbench.com/
  • Rodando CUDA em larga escala, gasta-se uma quantidade absurdamente desagradável de tempo de engenharia lidando com bugs dos drivers e bibliotecas da Nvidia.
    Não vejo muita gente ansiosa para depender ainda mais das bibliotecas da Nvidia.
  • Provavelmente não.
    Porque os detalhes da carga de trabalho — isto é, os parâmetros exatos, a representação dos dados na memória, os intervalos de valores etc. — fazem as estratégias de otimização divergirem bastante.

• Acabei de concluir um mestrado em HPC e fiz disciplinas de CUDA, MPI+CUDA e OpenCL; teria sido muito mais útil se eu tivesse lido um texto desses antes das aulas.
  Gostei especialmente do contexto em torno da parte sobre o que significa um warp estar executável.

• Primeiro, é um bom texto que explora bem vários cantos do assunto.
  Dito isso, se você não passar pela runtime API do CUDA, grande parte do vodu em espaço de usuário desaparece.
  Usando a API de driver e compilando o código-fonte do kernel como string com o compilador em tempo de execução da NVIDIA, dá para ver melhor o que acontece.
  Nem tudo, mas bastante coisa fica transparente.
  Uma versão mais “primitiva” está aqui:
  https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
  Para ver a mesma coisa em uma API C++ moderna, muito mais legível e ainda assim completamente transparente, veja isto:
  https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
  É um programa de exemplo da minha biblioteca header-only CUDA API wrappers.

  • A API de driver é legal porque permite tratar kernels CUDA como shaders recarregáveis a quente.
    É divertido poder desenvolver alterando o código durante a execução.

• Em bare metal?