Fundamentos de CUDA

O que é CUDA (e por que engenheiros de ML se importam)

CUDA (Arquitetura Unificada de Dispositivos de Computação (Compute Unified Device Architecture)) é a plataforma de programação da NVIDIA para executar código de propósito geral em GPUs NVIDIA. Em stacks modernos de ML, você raramente escreve CUDA diretamente — mas os conceitos de CUDA explicam por que certo código de modelo é rápido (ou lento) e como frameworks como PyTorch/JAX/TensorFlow mapeiam suas operações com tensores no hardware da GPU.

Em alto nível:

  • Uma GPU expõe paralelismo massivo: milhares de threads leves executam em grupos em sincronia.
  • O desempenho muitas vezes é limitado por movimentação de memória, não por aritmética. Isso se conecta diretamente a Memória e Largura de Banda.
  • A maioria das cargas de trabalho de aprendizado profundo é dominada por um pequeno conjunto de kernels: multiplicações de matrizes, convoluções, primitivas de atenção, normalização e operações elementwise — tipicamente via bibliotecas altamente otimizadas como cuBLAS e cuDNN.
  • Compilação e fusão de kernels importam muito, o que se relaciona com Compiladores e Runtimes.
  • O treinamento multi-GPU adiciona restrições de comunicação governadas por Interconexões (NVLink/InfiniBand) e bibliotecas coletivas (NCCL).

CUDA é, ao mesmo tempo:

  • um modelo de programação (threads/blocos/grids, espaços de memória, sincronização), e
  • uma cadeia de ferramentas/runtime (compilador nvcc, código intermediário PTX, APIs do driver/runtime CUDA, profilers).

Modelo de execução de GPU: Threads, Warps, Blocos e Grids

Programas CUDA executam “kernels” na GPU. Um kernel é uma função executada por muitas threads em paralelo.

Grid → Blocos → Threads

Ao lançar um kernel, você especifica:

  • um grid: todo o lançamento
  • composto por blocos de threads
  • cada bloco contém múltiplas threads

Cada thread tem índices:

  • threadIdx (dentro de um bloco)
  • blockIdx (dentro do grid)
  • blockDim, gridDim (formas)

Esse modelo hierárquico é central: threads no mesmo bloco podem cooperar eficientemente (memória compartilhada rápida + sincronização). Threads em blocos diferentes geralmente não conseguem sincronizar dentro de um único lançamento de kernel.

Warps: a unidade “real” de escalonamento

Dentro da GPU, as threads são executadas em grupos de geralmente 32 threads chamados warp. Implicações principais:

  • Se threads em um warp tomam desvios diferentes (um if em que metade vai para um lado e metade para o outro), o warp serializa os caminhos. Isso se chama divergência de warp (warp divergence).
  • Muitas regras de desempenho são sobre manter warps ocupados e alinhados nos padrões de acesso à memória.

Um exemplo mínimo de kernel (soma de vetores)

// CUDA C/C++ example
__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

void launch(const float* a, const float* b, float* c, int n) {
    int threads = 256;
    int blocks  = (n + threads - 1) / threads;
    vec_add<<<blocks, threads>>>(a, b, c, n);
}

Mesmo para essa operação simples, o desempenho pode variar bastante dependendo dos padrões de acesso à memória, do overhead de lançamento e de a operação ser fundida com operações adjacentes por um compilador.

Hierarquia de memória: o núcleo do desempenho em CUDA

Em ML, GPUs são rápidas porque têm throughput aritmético enorme — mas kernels frequentemente são limitados por largura de banda de memória e reutilização de dados. Entender a hierarquia de memória do CUDA ajuda você a raciocinar sobre gargalos.

Os principais espaços de memória

Do mais rápido/menor para o mais lento/maior (conceitualmente):

  • Registradores: por thread, mais rápidos, muito limitados. Uso excessivo (“pressão de registradores”) pode reduzir a ocupação.
  • Memória compartilhada: por bloco, rápida, scratchpad explicitamente gerenciado (geralmente dezenas de KB por bloco). Ótima para tiling e reutilização de dados.
  • Caches L1/L2: cache gerenciado por hardware para acessos à memória global.
  • Memória global (HBM/GDDR): grande, alta largura de banda, maior latência. É onde seus tensores normalmente vivem.
  • Memória constante / de textura: caminhos especializados com cache somente leitura (menos centrais para kernels típicos de DL hoje, mas ainda relevantes em alguns casos).

Para mais sobre por que memória domina o tempo de execução, veja Memória e Largura de Banda.

Coalescência: como tornar a memória global rápida

A memória global se torna eficiente quando threads em um warp acessam endereços contíguos. Isso é acesso coalescido (coalesced access).

  • Bom: thread 0 carrega x[0], thread 1 carrega x[1], … thread 31 carrega x[31]
  • Ruim: threads acessam x[stride * tid] com um stride grande → muitas transações de memória

Muitas decisões de layout de tensores (NCHW vs NHWC, tensores contíguos vs não contíguos, custos de transposição) se resumem a coalescência e localidade.

Tiling em memória compartilhada (por que GEMM é rápido)

O desempenho de multiplicação de matrizes depende de reutilizar valores. Uma estratégia clássica:

  1. Carregar um tile de A e B da memória global para a memória compartilhada.
  2. Sincronizar o bloco.
  3. Calcular produtos parciais usando memória compartilhada rápida.
  4. Repetir para os tiles.

Um esboço simplificado:

__global__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];

    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;

    float acc = 0.0f;
    for (int t = 0; t < N; t += 16) {
        As[threadIdx.y][threadIdx.x] = A[row * N + (t + threadIdx.x)];
        Bs[threadIdx.y][threadIdx.x] = B[(t + threadIdx.y) * N + col];
        __syncthreads();

        #pragma unroll
        for (int k = 0; k < 16; k++) {
            acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = acc;
}

Kernels reais de ML são significativamente mais complexos (cargas vetorizadas, prefetching, Tensor Cores, pipelining), mas o princípio — mover dados uma vez, reutilizar muitas vezes — é o mesmo.

Ocupação e ocultação de latência

GPUs ocultam a latência de memória alternando entre warps prontos enquanto outros aguardam a memória.

Ocupação (occupancy) é aproximadamente: “quantos warps podem estar ativos em um SM (Multiprocessador de Streaming (Streaming Multiprocessor)) ao mesmo tempo?”

A ocupação é influenciada por:

  • registradores usados por thread
  • memória compartilhada usada por bloco
  • threads por bloco
  • limites de hardware (varia conforme a “capacidade de computação (compute capability)” da GPU)

Nuance importante: ocupação máxima nem sempre é desempenho máximo. Às vezes, usar mais registradores (menor ocupação) reduz recomputação e vence.

Conclusão prática para ML:

  • Kernels com baixa reutilização de dados ficam limitados por largura de banda; aumentar a ocupação não ajuda muito.
  • Kernels eficientes buscam equilibrar throughput de instruções, throughput de memória e concorrência suficiente para ocultar latência.

Sincronização e comunicação dentro de uma GPU

Sincronização no nível de bloco

  • __syncthreads() sincroniza threads dentro de um bloco.
  • É essencial ao usar memória compartilhada de forma colaborativa.

Atômicas

Operações atômicas (atomics) fornecem atualizações concorrentes seguras (por exemplo, atomicAdd), mas podem virar gargalos por contenção. Elas aparecem em:

  • algumas reduções
  • atualizações esparsas
  • operações tipo histograma

GPUs modernas melhoraram o desempenho de atômicas, mas o desenho de algoritmos ainda tenta evitar contenção pesada.

Primitivas no nível de warp

Como warps executam em sincronia, CUDA oferece intrínsecos rápidos para comunicação dentro do warp (por exemplo, operações de shuffle) que podem implementar reduções eficientemente sem memória compartilhada.

Kernels de frameworks frequentemente usam isso para reduções de softmax, estatísticas de normalização de camada, etc.

Streams, assincronia e sobreposição de trabalho

CUDA expõe streams, que são sequências de operações que executam em ordem. Streams diferentes podem se sobrepor se os recursos de hardware permitirem.

Isso importa em ML porque você quer sobrepor:

  • transferências host↔device com computação,
  • kernels independentes entre si,
  • comunicação com computação no treinamento multi-GPU.

Exemplo prático: sobrepor cópia H2D com computação

Pseudo-código usando a API de runtime CUDA:

cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);

// pinned host memory is important for fast async transfers
cudaHostAlloc(&h_buf, bytes, cudaHostAllocDefault);

cudaMemcpyAsync(d_in, h_buf, bytes, cudaMemcpyHostToDevice, s1);
my_kernel<<<grid, block, 0, s2>>>(d_in, d_out);

cudaStreamSynchronize(s1);
cudaStreamSynchronize(s2);

Duas ideias-chave:

  • cudaMemcpyAsync só é realmente assíncrono com memória do host fixada (pinned/page-locked).
  • A sobreposição depende dos motores de cópia da GPU e do escalonamento.

Dataloaders e pipelines de entrada de aprendizado profundo frequentemente tentam usar esses princípios (mesmo que não diretamente via chamadas CUDA).

Precisão, Tensor Cores e por que FP16/BF16 são rápidos

GPUs modernas da NVIDIA incluem Tensor Cores, unidades especializadas para multiplicação-acumulação de matrizes. Elas entregam throughput enorme para:

  • FP16 (half precision)
  • BF16
  • TF32 (em Ampere e posteriores; usado para acelerar treinamento tipo FP32 para matmuls)
  • INT8/INT4 em contextos de inferência

Para muitos modelos, os ganhos de velocidade de precisão mista vêm em grande parte de:

  • aceleração por Tensor Cores de GEMMs e convoluções
  • menor pressão de largura de banda de memória (ativações/pesos com metade do tamanho)

Isso se conecta diretamente a Quantização para inferência e a técnicas de treinamento com precisão mista em Redes Neurais.

Nota prática:

  • Você normalmente não “usa Tensor Cores” manualmente; você habilita AMP (precisão mista automática (automatic mixed precision)) e confia em cuBLAS/cuDNN e em kernels gerados pelo compilador para selecionar caminhos de Tensor Cores.

A cadeia de ferramentas CUDA em uma figura

Mesmo que você nunca compile CUDA por conta própria, ajuda conhecer as camadas:

  • CUDA C++ / kernels: seu código-fonte (ou código gerado)
  • NVCC: driver de compilação que produz:
    • PTX (ISA virtual, um tanto portátil entre gerações de GPU)
    • cubin/SASS (código de máquina específico do hardware)
  • Driver/runtime CUDA carrega kernels compilados e os lança

Frameworks de ML adicionam mais camadas:

  • PyTorch: operações eager despacham para kernels CUDA, cuDNN/cuBLAS; compilação opcional via torch.compile
  • JAX: XLA compila grafos de computação em kernels de GPU fundidos
  • TensorFlow: grafo + XLA, ou TensorRT para inferência

Para entender por que compilação/fusão muda o desempenho drasticamente, veja Compiladores e Runtimes.

Bibliotecas CUDA que impulsionam ML

A maioria das cargas de trabalho de DL de alto desempenho depende fortemente de bibliotecas do fornecedor em vez de kernels customizados:

  • cuBLAS / cuBLASLt: GEMM (multiplicação de matrizes) e álgebra linear relacionada
  • cuDNN: convoluções, RNNs (legado), normalização, ativação, etc.
  • NCCL: coletivas multi-GPU (all-reduce, all-gather) críticas para treinamento distribuído
  • cuSPARSE: álgebra linear esparsa
  • Thrust / CUB: primitivas paralelas (scan, reduce, sort) frequentemente usadas dentro de operações customizadas

Se você vê um modelo ficar rápido depois de “usar atenção fundida” ou “flash attention”, isso tipicamente é uma história sobre:

  • melhor desenho de kernel (tiling, memória compartilhada, menos leituras/escritas)
  • menos lançamentos de kernel (fusão)
  • melhor uso de Tensor Cores

Overhead de lançamento de kernel, fusão e por que “muitas operações pequenas” são lentas

O lançamento de um kernel de GPU tem overhead não trivial (trabalho do driver/runtime, escalonamento). Em frameworks eager, muitas operações elementwise minúsculas podem ser dominadas pelo overhead de lançamento e pelo tráfego de memória.

Padrões comuns de otimização:

  • Fusão de kernel (kernel fusion): combinar múltiplas operações elementwise em um único kernel para que os dados sejam lidos/escritos uma vez.
  • Fusão de operador (operator fusion) para sequências como bias+ativação+dropout, layernorm+residual, etc.
  • Atenção fundida (fused attention): reduz tráfego de memória ao computar softmax e somas ponderadas em uma única passagem.

Essa é uma grande motivação por trás de compiladores modernos e sistemas de captura de grafo (novamente: Compiladores e Runtimes).

Noções básicas de multi-GPU: de CUDA a NCCL e interconexões

O CUDA em si é focado em nó único, mas sustenta a execução multi-GPU:

  • GPUs se comunicam via PCIe e/ou NVLink.
  • O treinamento distribuído depende fortemente de comunicação coletiva (collective communication) (all-reduce para gradientes).
  • O desempenho depende da topologia e de largura de banda/latência, cobertas em Interconexões (NVLink/InfiniBand).

Implicações práticas para ML:

  • Se o tempo por passo não melhora com mais GPUs, você pode estar limitado por comunicação (communication-bound).
  • Sobrepor comunicação com computação (por exemplo, all-reduce de gradientes durante o backprop) é fundamental.

Preocupações em nível de cluster como disponibilidade de GPU e justiça de jobs são cobertas em Escalonamento de GPU e Filas de Cluster.

Depuração e profiling: como você de fato fica mais rápido

Trabalho de desempenho em CUDA é empírico. Dois kernels que “parecem similares” podem diferir devido a acesso à memória, ocupação ou mistura de instruções.

Ferramentas comuns:

  • Nsight Systems: visão de linha do tempo (CPU, lançamentos de kernel, memcopies, sobreposições)
  • Nsight Compute: métricas no nível do kernel (ocupação, throughput de memória, warp stalls)
  • Profilers de frameworks:
    • PyTorch Profiler (frequentemente integra com Nsight)
    • TensorFlow profiler
    • Ferramentas de profiling do JAX/XLA

O que observar em cargas de trabalho de ML:

  • Você está limitado por largura de banda (high memory throughput, baixa utilização de computação)?
  • Muitos kernels pequenos demais (overhead de lançamento dominando)?
  • Cópias host↔device inesperadas (sincronizações, memória não fixada, conversões de forma)?
  • Tensores não contíguos causando transposições/cópias ocultas?
  • Baixa utilização de Tensor Cores quando você espera precisão mista?

Orientação prática para engenheiros de ML (sem escrever CUDA)

Você pode aplicar raciocínio de CUDA mesmo se nunca escrever um kernel:

  • Prefira layouts de tensores contíguos para operações pesadas; evite transposições desnecessárias.
  • Use precisão mista quando apropriado para explorar Tensor Cores.
  • Reduza overhead do lado Python agrupando trabalho; prefira operações fundidas quando disponíveis.
  • Fique atento a pontos acidentais de sincronização CPU/GPU (por exemplo, .item(), imprimir tensores, fluxo de controle dependente de forma).
  • Use ferramentas de compilação/fusão (torch.compile, XLA) quando ajudarem, e meça.
  • Ao escalar, trate comunicação como custo de primeira classe (topologia importa).

Para um contexto mais amplo de onde GPUs se encaixam em sistemas de IA, veja Introdução a Hardware. Para efeitos de armazenamento e pipeline de entrada (muitas vezes o gargalo oculto), veja Armazenamento.

Quando você *de fato* precisa de CUDA customizado

Kernels CUDA customizados são mais justificáveis quando:

  • você precisa de uma operação não fornecida por bibliotecas existentes,
  • você precisa de um kernel fundido para reduzir tráfego de memória,
  • você está implementando uma primitiva especializada de atenção/normalização/esparsidade,
  • você está otimizando para uma restrição específica de implantação (latência, pegada de memória).

Em muitos casos, opções de nível mais alto são mais fáceis primeiro:

  • use a interface de operações customizadas do framework
  • use Triton (DSL em Python que compila para kernels de GPU) se disponível na sua stack
  • confie na fusão do compilador e em otimizações de grafo

Mas, ainda assim, os fundamentos de CUDA continuam úteis: eles dizem por que um desenho específico de kernel fundido supera uma composição ingênua de operações.

Conceitos-chave para lembrar

  • CUDA organiza paralelismo como threads → warps → blocos → grids.
  • O desempenho frequentemente é dominado por movimentação de memória, não por matemática.
  • Acesso coalescido à memória global e reutilização via memória compartilhada são fundamentais.
  • Ocupação ajuda a ocultar latência, mas não é o único objetivo.
  • Streams habilitam concorrência e sobreposição, críticas para throughput.
  • Tensor Cores tornam precisão mista rápida; frameworks normalmente as acessam por meio de bibliotecas.
  • Profiling (Nsight + ferramentas do framework) é essencial para identificar gargalos reais.

CUDA é a “física” por baixo do ML acelerado por GPU: você não precisa dela para toda tarefa, mas é a forma mais rápida de construir intuição sobre desempenho, escalabilidade e trade-offs de projeto de sistemas.