CUDA-oxide: o compilador Rust-to-CUDA oficial da Nvidia

 (nvlabs.github.io)
1 pontos por GN⁺ 2 시간 전 | 1 comentários | Compartilhar no WhatsApp
  • cuda-oxide é um compilador experimental para escrever kernels GPU SIMT em Rust idiomático, próximo da segurança, e compilar código Rust padrão diretamente para PTX
  • Usa apenas Rust, sem DSLs nem bindings para linguagens estrangeiras, e pressupõe entendimento de ownership, traits e generics; a seção sobre async também exige conhecimento de .await
  • A v0.1.0 é uma release alfa inicial, então é preciso esperar bugs, recursos inacabados e mudanças incompatíveis de API
  • O exemplo é executado com cargo oxide run vecadd, e a função #[kernel] dentro de #[cuda_module] realiza soma de vetores com thread::index_1d()
  • #[cuda_module] incorpora artefatos de dispositivo ao binário host e gera um loader tipado e métodos de execução por kernel

Como usar e código gerado

  • Início rápido

    • Depois de atender aos pré-requisitos de instalação, compile e execute o exemplo com cargo oxide run vecadd
    • As instruções de instalação estão em prerequisites
    • O exemplo define a função #[kernel] vecadd dentro de um módulo #[cuda_module], obtém o índice com thread::index_1d() e grava a[i] + b[i] em DisjointSlice<f32>
    • No lado do host, usa CudaContext::new(0), o stream padrão, kernels::load(&ctx), e executa o kernel com DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed e LaunchConfig::for_num_elems(1024)
    • O resultado é trazido com c.to_host_vec(&stream) e confirma-se que result[0] == 3.0

  • Funcionamento de #[cuda_module]

    • #[cuda_module] incorpora os artefatos de dispositivo gerados ao binário host
    • Gera a função tipada kernels::load e métodos de execução por kernel
    • Quando for necessário carregar artefatos sidecar específicos ou criar código de execução customizado, ainda é possível usar as APIs de nível mais baixo load_kernel_module e cuda_launch!

Premissas e direção

  • O cuda-oxide tem como objetivo escrever kernels GPU com o sistema de tipos e o modelo de ownership do Rust, colocando a segurança como objetivo de primeira classe
  • GPUs têm detalhes sutis, então é necessário ler the safety model
  • Não é uma DSL, e sim um backend customizado de geração de código do rustc que compila Rust puro para PTX
  • Suporta execução assíncrona, compondo o trabalho de GPU como um grafo de DeviceOperation executado de forma lazy, agendado em um pool de streams, e aguardando os resultados com .await
  • Parte do pressuposto de familiaridade com ownership, traits e generics do Rust; os capítulos posteriores sobre programação async em GPU também exigem conhecimento de async/.await e de runtimes como o tokio
  • Como material de referência, são fornecidos The Rust Programming Language, Rust by Example e Async Book
  • A release v0.1.0 está em estágio alfa inicial, então espere bugs, recursos inacabados e mudanças incompatíveis de API

Leituras relacionadas

1 comentários

 
GN⁺ 2 시간 전
Comentários do Hacker News
  • Realmente impressionante. Já faz tempo que uso kernels CUDA customizados e https://crates.io/crates/cudarc, e isso parece que pode virar quase um substituto drop-in
    Estou especialmente curioso para ver como o tempo de build se compara. A maioria dos crates Rust para CUDA depende de CMake ou de chamadas ao nvcc, então a compilação pode ficar dolorosamente lenta
    Coincidentemente, na semana passada eu estava fazendo profiling de tempos de build e vi que ferramentas como sccache podem reduzir bastante o tempo de rebuild com cache de artefatos, mas o custo das chamadas customizadas ao nvcc continua existindo. Por exemplo, o candle da Hugging Face também chama comandos nvcc customizados na compilação dos kernels: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc é realmente muito bom
      Pessoalmente, não senti tanto essa parte de que a maioria dos crates Rust para CUDA chama CMake ou nvcc e por isso compila devagar. Se você olhar o crate cuda_setup, feito para lidar com scripts de build, ele é só um build.rs simples, então recompila apenas quando os arquivos mudam, e o tempo de compilação é bem pequeno comparado ao código Rust do lado da CPU
    • Fico curioso se outras pessoas também acham que o cuda-oxide parece quase um substituto drop-in para o cudarc
      Seria ótimo se fosse, mas pessoalmente acho que provavelmente está mais para um complemento. Também fico curioso sobre o que diferencia o cuda-oxide, além do fato de a NVIDIA ter controle total sobre ele
  • Parece estranho ir “direto para PTX”. O NVIDIA MLIR recente também é bem bom e rápido. Ou então poderiam ter mirado no Tile IR [1], que é mais fácil e está mais em alta, como o CuTile usa
    O Tile IR é um pouco mais de alto nível, então é bem mais fácil de usar como alvo, e você só perde em coisas como fusão de epílogo
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • Fiquei bastante curioso sobre como adaptaram o modelo de memória do Rust à semântica do CUDA. Também queria saber o que muda em relação ao CUDA C++ e se o sistema de tipos do Rust realmente pode trazer mais segurança ao CUDA
    Acho que escrever kernels para GPU é inerentemente inseguro. Pelo jeito como o hardware funciona e pela necessidade constante de otimização extrema, é difícil demais criar uma linguagem segura para isso
    • Há quatro diferenças grandes visíveis. Primeiro, lida com use-after-free e com a semântica de drop, em vez de exigir chamadas manuais a cudaFree
      Segundo, enquanto os argumentos void* do C++ são um array de ponteiros e só validam a quantidade, aqui os argumentos do kernel são forçados com cuda_launch!
      Terceiro, há a questão de aliasing em escritas mutáveis. Em C++, um código em que duas ou mais threads escrevem em out[i] com o mesmo i ainda compila, mas DisjointSlice e ThreadIndex não têm construtores públicos, e só se usa as APIs https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... index_1d, index_2d, index_2d_runtime
      Quarto, em C++ dá para fazer cuda memcpy de std::string ou praticamente qualquer POD e corromper o estado, mas aqui só aceita DisjointSlice, escalares e closures https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      Os detalhes estão em https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... e https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a.... Claro, isso não captura tudo, mas parece oferecer muito mais guardrails contra comportamento indefinido do que arquivos .cu raw
    • Vale lembrar que o modelo de memória do Rust é intencionalmente quase exatamente igual ao do C++. As operações atômicas também são iguais, e existem conceitos como proveniência (provenance)
      Ainda está para ver se é uma linguagem conveniente para programação em GPU, mas não me surpreenderia se desse para criar uma API parecida com DSL razoável para escrever código seguro enquanto ainda aproveita todas as esquisitices específicas de GPU. No fim, CUDA já não é algo assim?
    • A documentação entra em bastante detalhe. Existem uma camada segura, uma camada majoritariamente segura e uma camada insegura
      Para trabalhos paralelos que são seguros, mas difíceis de encaixar no modelo Send/Sync do Rust, foi preciso aceitar um pouco de aspereza
    • Acho que depende do objetivo. Do ponto de vista de quem escreve aplicações em Rust e quer usar computação em GPU ocasionalmente dentro delas, honestamente isso não me importa tanto
      Se eu puder aproveitar o modelo de memória ou de ownership com pouco atrito, ótimo. Mas se isso piorar muito a experiência de uso, então não quero seguir por esse caminho
      Para mim, a linha de base é o que o Cudarc faz hoje. Não há tanta interferência no gerenciamento de memória, só uma sintaxe imperativa envolvendo FFI e algumas linhas de script de build que chamam nvcc quando o kernel muda
  • Fico curioso sobre o que isso significa para o Slang[0]. A ideia central parece ser que as pessoas querem programar GPU em uma linguagem mais moderna, e agora parece que basta usar Rust
    Dito isso, gosto bastante de Slang
    [0]: https://shader-slang.org/
    • Escrever shaders é, pelo menos por enquanto, substancialmente diferente de escrever kernels CUDA. Shaders são ao mesmo tempo mais de alto nível e mais de baixo nível, e têm várias peculiaridades por terem sido projetados para um conjunto específico e limitado de recursos de driver/GPU
      Por exemplo, descriptor sets, registradores de recursos, limites de dispatch e afins
    • O alvo é diferente. O pessoal do Slang está mais interessado em programação gráfica do que em algoritmos de IA
      Linguagens de shading também são mais amigáveis em termos de recursos. Além disso, a NVIDIA já usa Slang em produção, e esse pessoal não vai reescrever os pipelines de shader em Rust
  • Já que o assunto é Rust e linguagens de programação “seguras”, fico curioso se alguém sabe mais detalhes sobre como a NVIDIA usa Spark/Ada
    Tudo o que consigo encontrar é isto
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • A partir da frase “sem DSL, sem bindings em linguagem estrangeira, só Rust”, já dá a impressão de que, apesar de ser uma porta oficial do CUDA, ninguém teve muito cuidado nem com o parágrafo de introdução
    Mesmo assim, tentei ignorar isso e ler a documentação, e no momento em que um IR customizado começou a parecer interessante, me deparei com uma frase no estilo “a implementação em MLIR é C++ com um pouco de TableGen, um sistema de build que exige compilar o LLVM inteiro e sessões de debug que fazem você questionar suas escolhas de carreira”, e aí ficou difícil levar esse setor a sério
    • O codebase inteiro parece ter sido basicamente escrito por IA
    • Se eles não tivessem usado IA na página, a reação teria sido “por que a NVIDIA não escreve seu site e sua documentação com IA? Será que eles mesmos não acreditam nessa história de funcionário gerenciando fábrica de IA e milhares de agentes?”
      Isso parece exatamente o tipo de dogfooding que se esperaria de uma empresa movida por hype de IA
    • Até o nome CUDA-oxide mostra que não sabem que a origem do nome Rust não é oxidação, e sim fungos
    • Não entendi exatamente qual é a reclamação. É porque alguém observou que MLIR é muito complexo e dependente de LLVM?
  • Coisas como TileLang https://github.com/tile-ai/tilelang e Tile Kernels https://github.com/deepseek-ai/TileKernels um dia vão tornar o CUDA obsoleto
    • CUDA já tem quase 20 anos e não vai desaparecer tão cedo
    • É uma afirmação bem grande para ter tão pouca base
  • Pelo documento https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo..., kernels de GPU rodam em milhares de threads vendo a mesma memória ao mesmo tempo; no CPU, o Rust evita data races com ownership e borrowing, mas na GPU há 2048 threads por SM começando na mesma função e apontando para o mesmo buffer de saída, então o borrow checker não foi projetado para isso
    O cuda-oxide torna estruturalmente seguro o caso comum de “uma thread escreve um elemento”, e para casos menos comuns como memória compartilhada, warp shuffles e intrínsecos de hardware, exige unsafe com contratos documentados; já recursos de ponta como TMA, tensor cores e comunicação em nível de cluster ficam totalmente manuais, em linha com a complexidade do hardware
    Mas isso não parece muito Rust-like. Em Rust, quando as abstrações existentes não se encaixam bem no problema, criam-se novas abstrações seguras. Rust for Linux é um exemplo disso
    Se não é seguro, então qual é a motivação para usar Rust? Tudo bem fornecer APIs unsafe para quem precisa extrair a última gota de performance, mas isso não deveria ser o padrão
    Isso lembra bibliotecas de espaço de usuário para APIs como io_uring ou Vulkan. Projetar APIs seguras para isso é bem difícil, e já houve tentativas que na prática não eram sound
  • Alguém sabe se isso vai permitir compartilhar structs entre host e device? Esse sempre foi um grande pedaço que faltava nos workflows Rust/CUDA até agora
    O mesmo vale para a barreira de serialização/bytes entre os dois lados
  • Uma coisa que sempre me deixou receoso ao usar Rust com CUDA é que o Rust adiciona um pouco de overhead que normalmente dá para ignorar, mas que aqui pode importar
    Por exemplo, fico me perguntando se verificações de limite de array podem causar uso extra de registradores e reduzir a concorrência do kernel