cuda-oxide: NVIDIA’s Experimental Compiler Enables Rust-Based GPU Kernel Development

Introduction: A New Era for GPU Programming with Rust

NVIDIA AI researchers have unveiled cuda-oxide, an experimental compiler backend that lets developers write CUDA SIMT (Single Instruction, Multiple Threads) GPU kernels using standard Rust code. Unlike existing approaches that require domain-specific languages, foreign function interface bindings, or C/C++ snippets, cuda-oxide compiles Rust directly to PTX (Parallel Thread Execution) — the low-level intermediate representation used by CUDA to target NVIDIA GPUs. This breakthrough promises to bring Rust’s safety guarantees and expressive power to the world of high-performance GPU computing.

Why This Matters: Simplifying GPU Kernel Authoring

Traditionally, writing GPU kernels means either mastering CUDA C++ or relying on Python abstractions like Triton that generate CUDA code under the hood. For Rust developers, the options have been limited and often cumbersome. cuda-oxide changes that by allowing kernel authoring, device intrinsics, and the SIMT execution model to be expressed natively in safe Rust. As the project’s stated design goal puts it: “bringing CUDA into Rust” — making the experience as close as possible to writing a __global__ function in C++, but with Rust’s memory safety and concurrency benefits.

Comparing Existing Solutions

The Rust GPU ecosystem already includes several projects, each with its own trade-offs:

• Rust-GPU targets SPIR-V for Vulkan/graphics compute workloads.
• rust-cuda uses a rustc codegen backend targeting NVVM IR, focusing on bringing Rust ergonomics (async/.await, parts of std) to NVIDIA GPUs.
• CubeCL employs an embedded DSL with a JIT runtime that cross-compiles to CUDA, ROCm, and WGPU.
• std::offload leverages LLVM’s implicit offload path for heterogeneous computing.

cuda-oxide occupies a unique niche: it prioritizes mapping CUDA concepts directly onto Rust, rather than abstracting them away. The NVlabs team has been coordinating with rust-cuda maintainers, viewing both projects as complementary.

How cuda-oxide Works: The Compilation Pipeline

At its heart, cuda-oxide is a custom rustc codegen backend — the part of the Rust compiler that normally emits native CPU machine code. Instead, the rustc-codegen-cuda crate intercepts the compilation at the CodegenBackend::codegen_crate() entry point and runs a dedicated pipeline for device code:

1. Rust source is processed by the standard rustc frontend.
2. The frontend produces Stable MIR (via the rustc_public API), a versioned and stable representation that avoids breakage across nightly updates.
3. This MIR is then transformed through several custom dialects built on Pliron, a Rust-native framework similar to MLIR but built entirely in Rust (cargo-based, no C++ or CMake required).
4. The Pliron dialects include dialect-mir (modeling MIR semantics with places and projections), mem2reg (promoting memory to registers), and dialect-llvm (lowering to LLVM IR).
5. Finally, LLVM IR is emitted and converted to PTX assembler.

Key Components: Stable MIR, Pliron, and Dialects

Two elements make this pipeline robust and developer-friendly:

• Stable MIR (rustc_public): The raw internal MIR of rustc changes unpredictably between nightly versions. By using the official stable API, cuda-oxide ensures its backend remains compatible without constant updates.
• Pliron: A lightweight, pure-Rust IR framework that avoids the heavy dependencies of upstream MLIR. The three custom dialects — dialect-mir, mem2reg, and dialect-llvm — handle each stage of lowering from Rust semantics to GPU code.

Design Philosophy: CUDA in Rust vs. Rust on GPUs

cuda-oxide’s philosophy contrasts sharply with some prior work. While rust-cuda aims to “bring Rust to NVIDIA GPUs” by abstracting over CUDA concepts (e.g., using async/.await for streams), cuda-oxide deliberately exposes the CUDA programming model — threads, blocks, shared memory — as native Rust constructs. This makes it easier for developers already familiar with CUDA C++ to transition, while also enabling fine-grained control over performance-critical details.

The result is a closer resemblance to writing a traditional kernel in C++ than to writing a generic Rust function that happens to execute on a GPU. For performance engineers and systems programmers, this direct mapping reduces cognitive overhead and allows for predictable optimization.

Conclusion: A Step Toward Safe, High-Performance GPU Compute

cuda-oxide is still experimental, but it represents a significant milestone in bridging Rust’s safety guarantees with NVIDIA’s GPU ecosystem. By compiling Rust directly to PTX without DSLs or FFI, it opens the door to writing robust, maintainable GPU kernels in one of the most beloved modern languages. As the team continues to refine the compiler and coordinate with projects like rust-cuda, the future of Rust on GPUs looks brighter than ever.