NVIDIA Oficializa Rust para Programação de GPUs: O Que Muda na Prática

Se você programa GPUs, sabe que a vida inteira foi assim: CUDA com C++, aceitar os segfaults, rezar para o ponteiro estar certo e seguir em frente. Só que em setembro de 2026, a NVIDIA decidiu que isso não precisa mais ser o único caminho. A empresa anunciou oficialmente o CUDA Rust, duas ferramentas que permitem escrever kernels GPU nativamente em Rust, com segurança de memória em tempo de compilação.

E não é um experimento de garagem. O cutile-rs já está publicado no crates.io e roda no HuggingFace Grout e no mistral.rs. Isso já é produção.

Por que isso importa (de verdade)

Vou ser direto: programação GPU em Rust não é novidade como conceito. Projetos como rust-gpu, wgpu e até o falecido Rust-CUDA tentaram resolver isso antes. Todos esbarraram no mesmo problema: o backend LLVM PTX não funcionava direito, gerava PTX inválido para operações comuns do Rust, e a experiência era tão dolorosa que a maioria dos devs voltava pro C++ depois de dois dias.

A diferença agora é que a própria NVIDIA está bancando a solução. Não é um projeto open source da comunidade torcendo para funcionar. É a dona do hardware dizendo “Rust é um caminho oficial para programar nossas GPUs”. Isso muda o cálculo completamente.

Três razões concretas:

1. Segurança de memória no kernel GPU: data races em código CUDA são um pesadelo para debugar. O borrow checker do Rust agora se estende até a fronteira GPU.
2. Toolchain oficial: não precisa mais hackear backends LLVM. A NVIDIA mantém o compilador.
3. Ecossistema Rust: cargo add cutile e pronto. Sem CMake, sem nvcc, sem dor de cabeça de build system.

As duas trilhas: cuda-oxide e cutile-rs

A NVIDIA não lançou uma ferramenta. Lançou duas, cada uma para um perfil diferente de programador.

cuda-oxide: controle total, thread por thread

O cuda-oxide segue o modelo SIMT (Single Instruction, Multiple Threads), que é o modelo clássico do CUDA. Você controla cada thread individualmente, decide o layout de memória, gerencia blocos e grids na mão.

Por baixo dos panos, ele funciona assim: é um codegen backend customizado para o rustc que compila código Rust passando pelo MIR (Mid-level IR do Rust), depois pelo framework Pliron IR, LLVM, e finalmente gera PTX.

use cuda_oxide::prelude::*;

#[kernel]
fn vector_add(a: &[f32], b: &[f32], c: &mut DisjointSlice<f32>) {
    let idx = thread_idx_x() + block_idx_x() * block_dim_x();
    if idx < c.len() {
        c[idx] = a[idx] + b[idx];
    }
}

Repara no DisjointSlice. Esse tipo garante em tempo de compilação que nenhuma thread vai escrever na mesma posição de memória que outra. Se você tentar fazer aliasing, o compilador barra. Isso é algo que em C++ CUDA você descobre com um segfault às 3 da manhã.

Requisitos:

• Linux (sem Windows por enquanto)
• GPU com compute capability 8.0+ (Ampere ou mais novo)
• CUDA 12.x+
• Toolchain nightly do Rust (pinned em 2026-04-03)

cargo +nightly-2026-04-03 install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide

Status: alpha. A NVIDIA deixa claro que a API pode mudar. É para quem quer explorar e contribuir, não para colocar em produção amanhã.

cutile-rs: o caminho pragmático

Se o cuda-oxide é o bisturi, o cutile-rs é o bisturi robótico. Você descreve operações sobre blocos de dados (tiles), e o compilador decide como mapear esses tiles para threads e hardware.

use cutile::prelude::*;

#[cutile::kernel]
fn matmul(a: &Tensor<f32>, b: &Tensor<f32>, c: &mut Tensor<f32>) {
    let tile_a = a.partition([128, 64]);
    let tile_b = b.partition([64, 128]);
    let acc = cutile::zeros::<f32>([128, 128]);

    for k in 0..a.shape()[1] / 64 {
        acc += tile_a[k] @ tile_b[k];
    }
    c.store(acc);
}

O &mut Tensor passado no kernel é tratado pelo sistema de tipos como um bloco isolado com acesso exclusivo garantido. O borrow checker do Rust se estende até a GPU: se dois workers tentarem escrever no mesmo tensor, o código não compila.

Requisitos:

• Linux
• GPU com compute capability 8.0+
• CUDA 13.3
• Rust estável 1.89+

cargo add cutile

Status: mais maduro. Já está no crates.io, já roda em produção no HuggingFace Grout e no mistral.rs. Esse é o que a maioria dos devs vai querer usar primeiro.

Comparação lado a lado

O elefante na sala: performance

Eu sei o que você está pensando. “Legal, mas roda mais devagar que C++ CUDA?”

A resposta honesta é: depende da trilha.

O cutile-rs usa JIT compilation. Na primeira execução, o kernel é compilado; nas seguintes, roda do cache. Os benchmarks disponíveis até agora mostram performance comparável ao CUDA C++ para operações de GEMM (multiplicação de matrizes) e convoluções. O fato de o HuggingFace já usar no Grout para inferência é um sinal forte: se fosse significativamente mais lento, eles não trocariam.

O cuda-oxide é mais difícil de avaliar porque ainda está em alpha. A compilação passa por mais camadas (MIR, Pliron IR, LLVM, PTX), o que em teoria pode introduzir overhead. Mas como gera PTX no final, o driver NVIDIA ainda otimiza na hora de rodar, igual faz com código C++.

O ponto é: você provavelmente não vai perder performance mensurável usando cutile-rs, e vai ganhar muito em produtividade e segurança. Se precisa espremer cada nanosegundo, o cuda-oxide dá o controle necessário, embora com uma experiência mais crua.

O que deu errado antes (e por que agora é diferente)

A tentativa mais conhecida de Rust no CUDA foi o projeto Rust-CUDA (sem ligação com a NVIDIA), que surgiu em 2021. O problema era fundamental: o backend LLVM PTX não era projetado para Rust. Operações comuns como iteradores, closures e traits genéricas geravam PTX inválido ou crashes no compilador.

Outros projetos tentaram caminhos alternativos:

• rust-gpu (Embark Studios): compilava Rust para SPIR-V (Vulkan), não para CUDA. Funcionava para shaders gráficos mas não para compute pesado.
• wgpu: abstração sobre WebGPU/Vulkan/Metal/DirectX. Portável, mas sem acesso ao hardware NVIDIA de forma otimizada.
• vulkano: wrapper Rust para Vulkan. Bom para gráficos, limitado para GPGPU.

O que a NVIDIA fez de diferente foi construir a toolchain do zero especificamente para suas GPUs. O cuda-oxide tem um codegen backend próprio, não depende do backend PTX genérico do LLVM. E o cutile-rs usa o Tile IR nativo da NVIDIA, que já foi otimizado para o hardware deles.

É a diferença entre tentar fazer um carro andar em trilhos de trem vs construir um carro para a estrada certa.

Quem deveria prestar atenção agora

Nem todo mundo precisa sair correndo para reescrever código CUDA em Rust. Aqui está quem eu acho que se beneficia mais:

Quem está começando com GPU programming: se você já sabe Rust e quer entrar no mundo GPU, cutile-rs é provavelmente o caminho com menos atrito que existe hoje. cargo add cutile, escreve um kernel, roda. Sem precisar aprender a toolchain C++ do CUDA.

Times que já usam Rust no backend: se seu stack é Rust e vocês precisam de aceleração GPU (inferência de ML, processamento de sinais, simulações), agora dá para manter tudo na mesma linguagem. Menos context switching, menos bugs na fronteira entre linguagens.

Pesquisadores de ML que querem kernels customizados: em vez de escrever CUDA C++ para um operador customizado do PyTorch, você pode escrever em Rust com segurança de memória e integrar via FFI.

Quem NÃO precisa mudar agora: se você tem uma codebase CUDA C++ funcionando, com testes, em produção, sem bugs de memória: relaxa. A NVIDIA prometeu interoperabilidade entre CUDA Rust, C++ e Python. Não é um “agora ou nunca”.

O futuro: interoperabilidade e além

A NVIDIA já sinalizou que o plano é permitir misturar kernels escritos em Rust, C++ e Python no mesmo projeto. Isso significa que a migração pode ser gradual: novos kernels em Rust, kernels existentes ficam em C++, e eventualmente você migra o que faz sentido.

Também tem a questão do ecossistema. Com cutile-rs no crates.io, a expectativa é que bibliotecas de ML, processamento de imagem e computação científica comecem a oferecer backends GPU via Rust. O mesmo movimento que aconteceu com o ecossistema Python/CUDA pode se repetir com Rust, só que dessa vez com segurança de memória desde o início.

E tem o contexto maior. Em 2026, a Microsoft colocou Rust como linguagem Tier-1 para Windows e Azure. A AWS já usa Rust no Firecracker e no S3. O kernel do Linux aceita código Rust desde 2022. Agora a NVIDIA fecha o último gap: GPU programming. Rust está se tornando a linguagem que roda em todo lugar, do kernel ao chip gráfico.

Como começar (em 5 minutos)

Se você quer testar agora, o caminho mais rápido é o cutile-rs:

# Verifique sua GPU
nvidia-smi

# Crie um projeto
cargo new gpu-test && cd gpu-test

# Adicione cutile
cargo add cutile

# Escreva seu primeiro kernel (src/main.rs)

use cutile::prelude::*;

fn main() {
    let a = vec![1.0f32; 1024];
    let b = vec![2.0f32; 1024];

    let result = cutile::run(|ctx| {
        let da = ctx.upload(&a);
        let db = ctx.upload(&b);
        let dc = ctx.alloc::<f32>(1024);

        ctx.launch(vector_add, (&da, &db, &mut dc));
        ctx.download(&dc)
    });

    println!("Resultado: {:?}", &result[..5]);
}

#[cutile::kernel]
fn vector_add(a: &[f32], b: &[f32], c: &mut [f32]) {
    let i = cutile::index();
    c[i] = a[i] + b[i];
}

cargo run --release
# Resultado: [3.0, 3.0, 3.0, 3.0, 3.0]

Compare com o equivalente em CUDA C++: um CMakeLists.txt, um nvcc, headers do CUDA, cudaMalloc, cudaMemcpy, cudaFree… Acho que você entende a diferença.

Na prática: migrando um kernel real

Para dar um exemplo mais concreto, veja como fica a migração de um kernel de redução (soma de array) que é clássico em tutoriais CUDA.

CUDA C++ (versão clássica):

__global__ void reduce_sum(float* input, float* output, int n) {
    __shared__ float sdata[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? input[i] : 0.0f;
    __syncthreads();

    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }

    if (tid == 0) output[blockIdx.x] = sdata[0];
}

Problemas que você já conhece: se errar o tamanho do sdata, acessa memória inválida. Se esquecer o __syncthreads(), race condition silenciosa. Se o n não for múltiplo do tamanho do bloco, precisa do guard manual.

cutile-rs (versão Rust):

#[cutile::kernel]
fn reduce_sum(input: &Tensor<f32>, output: &mut Tensor<f32>) {
    let tile = input.partition([256]);
    let sum = tile.reduce(|a, b| a + b);
    output.store(sum);
}

Seis linhas. Sem shared memory manual, sem sync explícito, sem guard de bounds. O compilador cuida de tudo isso e garante que o acesso à memória é exclusivo por tile.

Claro, isso é o caso feliz. Kernels mais complexos (fusão de operações, acesso coalescente customizado, pipelines de memória) ainda vão exigir mais código. Mas o ponto é que o floor de complexidade caiu drasticamente. O que antes exigia 50 linhas de C++ com 3 armadilhas de memória agora exige 6 linhas de Rust sem armadilhas.

O impacto no ecossistema de ML

Não dá para falar de CUDA sem falar de machine learning. Hoje, praticamente toda framework de ML (PyTorch, JAX, TensorFlow) depende de kernels CUDA escritos em C++. Quando você chama torch.matmul(), por baixo roda um kernel cuBLAS em C++.

Com cutile-rs, a porta se abre para kernels customizados em Rust. Imagine escrever um operador de atenção customizado para seu modelo de transformer, com seguranca de memória garantida pelo compilador, sem precisar tocar em C++. Isso é especialmente relevante para pesquisadores que precisam de operadores não-padrão: sparse attention, mixture of experts, quantização em tempo de execução.

O HuggingFace já saiu na frente integrando cutile-rs no Grout, o engine de inferência deles. O mistral.rs, que roda modelos Mistral localmente, também já usa. Esses são sinais claros de que a comunidade de ML em Rust não vai esperar para adotar.

E tem um efeito cascata interessante: com kernels GPU em Rust, projetos como candle (framework de ML em Rust puro do HuggingFace) e burn (framework de deep learning em Rust) ganham acesso a GPU compute de primeira classe sem depender de bindings C++. O stack inteiro pode ser Rust, do modelo ao kernel.

O que fica na mesa

Algumas limitações que você precisa conhecer:

• Só Linux: sem suporte a Windows ou macOS ainda. Se você desenvolve no Mac, vai precisar de uma máquina remota com GPU NVIDIA.
• GPUs Ampere+: nada de GPUs mais antigas. Se sua placa é uma GTX 1080, vai ter que esperar (ou fazer upgrade).
• cuda-oxide é alpha: API instável, pode ter bugs, documentação limitada. Use por sua conta e risco.
• CUDA 13.3 para cutile-rs: se você está em CUDA 12.x, vai precisar atualizar o driver.

Essas limitações fazem sentido para um lançamento inicial. A NVIDIA claramente priorizou fazer funcionar bem no hardware moderno antes de se preocupar com retrocompatibilidade.

Quem programa GPU sabe que a maior parte do tempo não é gasto escrevendo o kernel. É debugando race conditions, entendendo por que o resultado está errado, e lutando contra o build system. Se Rust resolve dois desses três problemas desde o primeiro dia, já é uma vitória enorme.

A pergunta não é se Rust vai se tornar uma linguagem relevante para GPU programming. Isso já aconteceu. A pergunta é quanto tempo até cutile-rs ter o mesmo ecossistema de bibliotecas que o CUDA C++ construiu em 15 anos.