Arquivo da categoria: CUDA

Compute Capability: O “detalhe” de build CUDA que pode evitar até ≈ +3300% de prejuízo em performance.

CUDA, Dicas Linux

Em projetos de IA, é comum a equipe de TI usar uma GPU “padrão de homologação” na nuvem (por exemplo, NVIDIA T4 no ambiente de teste) para compilar, testar e validar o pipeline. O problema aparece quando esse mesmo binário é executado em GPUs mais novas no ambiente produtivo e o desempenho fica absurdamente abaixo do que o hardware poderia entregar (e ninguém percebe por diversos fatores).

O motivo costuma ser simples: o binário foi compilado mirando a Compute Capability errada.

A Compute Capability (CC) é o identificador (X.Y) que representa o conjunto de recursos e instruções suportados por uma geração de GPUs NVIDIA. Na tabela oficial da NVIDIA, por exemplo, a T4 é CC 7.5 e a RTX 4060 Ti é CC 8.9. A seguir, vamos entender o porquê isso importa tanto e usar o seu exemplo real para mostrar o impacto: 4,30 ms vs 0,13 ms um ganho de ~33× (≈ +3208%, frequentemente arredondado como “3300%”). https://developer.nvidia.com/cuda/gpus

O que é Compute Capability e por que ela manda no desempenho

Compute Capability não é “marketing” é um contrato técnico entre o seu código e a microarquitetura da GPU. Ela define:

  • Quais instruções existem (ISA/SASS disponível naquele SM)
  • Quais otimizações o compilador pode aplicar
  • Quais recursos de hardware podem ser usados (ex.: instruções especializadas, caminhos de memória, operações assíncronas, etc.)
  • Qual PTX (código intermediário) pode ser gerado sem “assumir” features que não existem na GPU-alvo

A NVIDIA lista CC por GPU e deixa explícito que diferentes modelos mapeiam para CC diferentes (T4=7.5; RTX 4060 Ti=8.9). Em termos práticos: compilar para CC antiga em uma GPU nova pode “capar” o compilador, porque ele precisa gerar um código que só usa features garantidas naquela CC antiga.

Como o nvcc decide o que vai rodar: PTX, SASS e “fatbinary”

O nvcc normalmente gera:

  • PTX (intermediário, “assembly virtual”)
  • CUBIN/SASS (binário nativo para um SM específico, ex.: sm_75, sm_89)
  • Tudo isso pode ser empacotado num fatbinary dentro do executável

Na documentação do nvcc, a NVIDIA explica que o executável carrega imagens de código e, em tempo de execução, o runtime/driver escolhe a imagem mais apropriada para a GPU presente.

E aqui mora o ponto-chave:

  • Se você inclui SASS nativo para a GPU (sm_89 na 4060 Ti), ela roda direto com o melhor código possível.
  • Se você não inclui SASS para aquela GPU, mas inclui PTX, o driver pode fazer JIT (compilar na hora) porém limitado ao PTX e às features assumidas naquele PTX.
  • Se você inclui apenas código para uma GPU mais antiga e/ou PTX “antigo”, você pode até rodar, mas deixa performance na mesa.

O artigo técnico da NASA sobre compilação CUDA resume bem o modelo em 2 estágios (PTX → CUBIN) e descreve o papel de compute_XY vs sm_XY, além do JIT e do fatbinary como estratégia de portabilidade. Este texto da NASA foi inspiração deste post: https://www.nas.nasa.gov/hecc/support/kb/compiling-cuda-applications-for-different-generations-of-nvidia-gpus-at-nas_700.html

O caso real: time compila no ambiente de homologação T4 (CC 7.5) e executa na RTX 4060 Ti (CC 8.9)

Compilação “errada” (mirando T4 / compute_75)

Um equipe compila na instância com T4 (muito usada em homologação e testes) com:

$ nvcc matmul.cu -o matmul_wrong -arch=compute_75 -gencode arch=compute_75,code=sm_75

Executando esse binário numa RTX 4060 Ti (CC 8.9), o resultado medido foi:

$ ./matmul_wrong
Tempo de execução: 4.30 ms

Compilação “certa” (mirando a GPU real do alvo / compute_89)

Agora compilando especificamente para a RTX 4060 Ti:

$ nvcc matmul.cu -o matmul_right -arch=compute_89 -gencode arch=compute_89,code=sm_89

Resultado:

$ ./matmul_right
Tempo de execução: 0.13 ms

O “ganho de 3300%” (matemática do impacto)

  • Speedup (vezes mais rápido): 4,30 / 0,13 ≈ 33,08×
  • “% mais rápido” (comparando contra o tempo menor): (4,30 − 0,13) / 0,13 × 100 ≈ 3207,7%
  • Redução do tempo: (4,30 − 0,13) / 4,30 × 100 ≈ 97,0%

Ou seja: chamar isso de “~3300%” é um arredondamento comum o número exato pelo seu benchmark dá ~3208%.

Por que a Compute Capability errada pode causar um abismo desses?

Mesmo que o código-fonte CUDA seja igual, o código gerado pode mudar drasticamente quando você troca o alvo (sm_75 → sm_89), porque:

  1. Instruções e agendamentos diferentes
    • O backend (ptxas) gera SASS específico por arquitetura. Mirar sm_89 permite escolher instruções, latências e scheduling adequados ao hardware.
  2. Limitação por “conjunto mínimo de features”
    • Quando você gera PTX assumindo compute_75, você está dizendo ao compilador: “use apenas o que é garantido em CC 7.5”.
    • Mesmo rodando numa GPU CC 8.9, o caminho pode ficar preso a decisões conservadoras (ou depender mais de JIT).
  3. JIT não é “milagre”: ele compila o que você deu
    • Se você entrega PTX antigo, o driver compila, mas não inventa features que o PTX/arquitetura-alvo não permitiram assumir. A própria doc do nvcc descreve que PTX pode ser embutido e compilado dinamicamente quando não há binário nativo adequado.

Homologação na nuvem ≠ alvo final

  • No exemplo, a T4 (CC 7.5) é ótima para CI e testes, mas não pode ditar o “target único” se você entrega para desktops Ada (CC 8.9) e isso vale muito em pipelines de IA (inferencia, kernels custom, extensões CUDA, etc.).

A melhor prática para times: compile uma vez e rode bem em T4 e em 4060 Ti

Em vez de escolher “T4 ou 4060 Ti”, faça o executável carregar ambos:

  • SASS nativo para T4 (sm_75) → sem JIT, ótimo no ambiente de teste/homologação
  • SASS nativo para 4060 Ti (sm_89) → máximo desempenho no desktop
  • PTX para forward-compat (opcional, mas recomendado quando você não controla todos os alvos)

A própria documentação do nvcc mostra como --generate-code/-gencode permite repetir arquiteturas e também como incluir PTX (usando code=compute_XX) para JIT quando não existir SASS correspondente.

Exemplo de “fat binary” bem pragmático para seu cenário:

$ nvcc matmul.cu -o matmul_fat \
--generate-code arch=compute_75,code=sm_75 \
--generate-code arch=compute_89,code=sm_89 \
--generate-code arch=compute_89,code=compute_89

Assim:

  • Na T4, ele usa sm_75.
  • Na RTX 4060 Ti, ele usa sm_89.
  • Em GPUs futuras (acima de 8.9), se não houver SASS exato, ainda há PTX 8.9 para o driver compilar via JIT.

Benchmark honesto: duas dicas rápidas para não “se enganar”

Seu micro benchmark já é excelente para ilustrar o ponto, mas em artigo eu sempre recomendo mencionar:

  • Warm-up: rode o kernel algumas vezes antes de medir (evita pegar JIT/cache/boost “frio”).
  • Várias iterações: meça média/mediana de N execuções (evita ruído de clock/boost).

Isso não muda a mensagem principal (CC correta importa), mas dá credibilidade técnica.

O código do exemplo (para referência no artigo)

Abaixo está o kernel e a medição por cudaEvent exatamente como você forneceu (um microbenchmark simples para evidenciar diferenças de codegen/arquitetura):

#include <stdio.h>
#include <cuda.h>

__global__ void matmul_kernel(float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx < N) {
        float sum = 0.0f;
        for (int i = 0; i < N; i++) {
            sum += A[idx * N + i] * B[i * N + idx];
        }
        C[idx] = sum;
    }
}

int main() {
    const int N = 1024;
    const int size = N * N * sizeof(float);

    float *hA = (float*)malloc(size);
    float *hB = (float*)malloc(size);
    float *hC = (float*)malloc(N * sizeof(float));

    for(int i = 0; i < N*N; i++){
        hA[i] = 1.0f;
        hB[i] = 1.0f;
    }

    float *dA, *dB, *dC;
    cudaMalloc(&dA, size);
    cudaMalloc(&dB, size);
    cudaMalloc(&dC, N * sizeof(float));

    cudaMemcpy(dA, hA, size, cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB, size, cudaMemcpyHostToDevice);

    dim3 block(256);
    dim3 grid((N + block.x - 1) / block.x);

    // Medição
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);
    matmul_kernel<<<grid, block>>>(dA, dB, dC, N);
    cudaEventRecord(stop);

    cudaDeviceSynchronize();

    float ms = 0.0f;
    cudaEventElapsedTime(&ms, start, stop);

    printf("Tempo de execução: %.2f ms\n", ms);

    cudaMemcpy(hC, dC, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(dA); cudaFree(dB); cudaFree(dC);
    free(hA); free(hB); free(hC);

    return 0;
}

Fechando: a regra de ouro

Se o seu time como no exemplo aqui, compila em T4 (CC 7.5) para homologação, mas o alvo também inclui Ada (CC 8.9), não escolha um único target.

  • Compute Capability certa = performance “de graça”
  • Compute Capability errada = gargalo invisível que pode custar ordens de grandeza (como seus 33×)

E o melhor: corrigir isso geralmente é só ajustar flags de build e de quebra você padroniza uma entrega que roda bem em cloud e local.

Checklist final para o time de TI (cloud + desktop)

  • Não compile “só para a GPU da de testes” se o binário roda fora dela
  • Gere fat binary com -gencode para todos os alvos relevantes
  • Inclua PTX (ex.: code=compute_XX) quando você não controla 100% dos alvos
  • Padronize isso no CMake/CI (não deixe o dev “lembrar” na mão)
  • Documente o CC alvo por ambiente (homologação, dev, produção)

Conclusão

Compute Capability não é detalhe: é uma decisão de engenharia de build que pode te entregar performance “de graça” ou te fazer carregar um gargalo invisível por meses.