NVIDIA AI Just Released cuda-oxide: An Experimental Rust-to-CUDA Compiler Backend that Compiles SIMT GPU Kernels Directly to PTX

NVIDIA AI researchers just lately launched cuda-oxide, an experimental compiler that permits builders to write CUDA SIMT (Single Instruction, Multiple Threads) GPU kernels in commonplace Rust code. The undertaking compiles Rust straight to PTX (Parallel Thread Execution) — the assembly-like intermediate illustration that CUDA makes use of to goal NVIDIA GPUs — with out requiring domain-specific languages, international perform interface bindings, or C/C++ code.

How This Makes a Change

Writing GPU kernels immediately sometimes means writing C++ and utilizing the CUDA programming mannequin straight, or counting on Python-level abstractions like Triton that generate CUDA underneath the hood. The Rust GPU ecosystem has had tasks trying to bridge this hole — Rust-GPU targets SPIR-V for Vulkan/graphics compute, rust-cuda makes use of a rustc codegen backend concentrating on NVVM IR, CubeCL makes use of an embedded DSL with a JIT runtime that cross-compiles to CUDA/ROCm/WGPU, and std::offload makes use of LLVM’s implicit offload path.

cuda-oxide occupies a particular place on this house. Its said design heart is “bringing CUDA into Rust” — kernel authoring, machine intrinsics, the SIMT execution mannequin, and the CUDA programming mannequin expressed natively in secure Rust — nearer in spirit to writing a __global__ perform in C++ than to writing a generic Rust perform that occurs to run on a GPU. By distinction, the closest neighbor, rust-cuda, focuses on “bringing Rust to NVIDIA GPUs”: Rust ergonomics like async/.await, components of the usual library operating on-device, and a Rust-first programming mannequin that abstracts over CUDA ideas. The NVlabs crew notes it has been coordinating with rust-cuda maintainers and considers the 2 tasks complementary.

The Compilation Pipeline

At the core of cuda-oxide is a customized rustc codegen backend — the layer within the Rust compiler liable for producing machine code. Instead of emitting native CPU code, the rustc-codegen-cuda crate intercepts the compiler on the CodegenBackend::codegen_crate() entry level and runs a separate pipeline for machine code:

Rust Source → rustc frontend → rustc_public (Stable MIR) → dialect-mir → mem2reg → dialect-llvm → LLVM IR (.ll) → PTX (.ptx)

Here are some necessary parts:

Why rustc_public? The uncooked inside MIR illustration in rustc modifications between nightly variations with no stability ensures. cuda-oxide makes use of rustc_public — often known as Stable MIR — which is Rust’s official versioned, steady API over the compiler’s internals. This lets the backend learn MIR with out breaking on each nightly replace.

What is Pliron? The center levels use Pliron, a Rust-native MLIR-like IR framework written completely in Rust. Choosing Pliron as a substitute of upstream MLIR means your entire compiler builds with cargo — no C++ toolchain, no CMake, no tablegen. cuda-oxide defines three customized Pliron dialects: dialect-mir (modeling Rust MIR semantics — locations, projections, rvalues, terminators), dialect-llvm (modeling LLVM IR with textual .ll export), and dialect-nvvm (NVIDIA GPU intrinsics like thread indexing, boundaries, and TMA).

What does llc do? After the dialect-llvm printer serializes the IR right into a textual .ll file, the exterior llc binary (the LLVM static compiler with NVPTX backend) compiles it to PTX meeting. This is the one stage exterior pure Rust. The ensuing .ptx file is written subsequent to the host binary — for instance, goal/debug/vecadd.ptx — and loaded by the CUDA driver at runtime.

You as a developer can observe every stage with:

cargo oxide pipeline vecadd

This prints the complete hint from Rust MIR by every dialect down to PTX output.

Single-Source Compilation and the Host/Device Split

Host and machine code reside in the identical .rs supply file. cargo oxide units -Z codegen-backend=librustc_codegen_cuda.so, which routes code era by cuda-oxide’s backend. The backend then scans compiled code for monomorphized features whose names carry the reserved cuda_oxide_kernel_<hash>_<title> prefix — the namespace that the #[kernel] proc macro creates. Functions matching that prefix undergo the cuda-oxide pipeline to produce PTX; all different host code is delegated to rustc’s commonplace LLVM backend. The results of a single cargo oxide construct is a number binary plus a .ptx file.

cargo oxide run vecadd
cargo oxide debug vecadd --tui    # debug with cuda-gdb

Device code from library dependencies is compiled lazily: the backend reads their Stable MIR from .rlib metadata on demand, solely compiling features a kernel truly calls.

What You Can Write in a Kernel

cuda-oxide helps a significant subset of Rust in GPU kernel features, marked with the #[kernel] attribute macro. This contains:

• Generic features with monomorphization — fn scale<T: Copy>(...) is compiled to a concrete PTX kernel per kind used on the name web site.
• Closures with captures — closures handed from the host are scalarized and handed as PTX kernel parameters routinely.
• User-defined structs and enums — commonplace Rust information buildings work inside kernels.
• Pattern matching — match, if let, and associated constructs work in machine code.
• Full GPU intrinsics — the cuda-device crate supplies wrappers for thread indexing, warp operations (shfl_sync, ballot_sync, and so on.), shared reminiscence, boundaries, TMA (Tensor Memory Accelerator), Thread Block Clusters, and scoped atomics (6 sorts × 3 scopes × 5 orderings).

One necessary GPU-specific compiler element: rustc’s SoarThreading MIR optimization — which duplicates perform calls into each branches of an if-statement — is disabled for machine code in cuda-oxide. On CPUs it is a secure optimization, however on GPUs it breaks barrier semantics: all threads in a block should converge on the similar bar.sync instruction, and duplicating it throughout branches violates that requirement. Additionally, sync primitives are marked convergent within the emitted LLVM IR so that LLVM’s optimization passes can’t transfer or duplicate them throughout management stream.

How to Use NVIDIA Star Elastic

nvcc --version

# Install the pinned nightly toolchain
rustup toolchain set up nightly-2026-04-03

# Add required elements
rustup element add rust-src rustc-dev 
  --toolchain nightly-2026-04-03

# Confirm the toolchain is energetic
rustup present

# Ubuntu/Debian
sudo apt set up llvm-21

# Verify the NVPTX backend is current
llc-21 --version | grep nvptx

# Pin to a particular llc binary
export CUDA_OXIDE_LLC=/usr/bin/llc-21

# Install the complete clang-21 bundle (contains useful resource headers)
sudo apt set up clang-21

# Alternatively, the -dev header bundle additionally works
sudo apt set up libclang-common-21-dev

git clone https://github.com/NVlabs/cuda-oxide.git
cd cuda-oxide

# cargo oxide works out of the field through a workspace alias
cargo oxide run vecadd

# Install globally from the git supply
cargo set up 
  --git https://github.com/NVlabs/cuda-oxide.git 
  cargo-oxide

# On first run, cargo-oxide fetches and builds the codegen backend

cargo oxide physician

# Build and run end-to-end
cargo oxide run vecadd

✓ SUCCESS: All 1024 parts right!

# Print the complete Rust MIR — dialect-mir — mem2reg — dialect-llvm — LLVM IR — PTX hint
cargo oxide pipeline vecadd

cargo oxide debug vecadd --tui

use cuda_device::{kernel, thread, DisjointSlice};

// Tier 1 security: race-free by development, no `unsafe` wanted.
// DisjointSlice::get_mut() solely accepts a ThreadIndex —
// a hardware-derived opaque kind guaranteeing distinctive writes per thread.
#[kernel]
pub fn scale(enter: &[f32], issue: f32, mut out: DisjointSlice<f32>) {
    let idx = thread::index_1d();
    if let Some(elem) = out.get_mut(idx) {
        *elem = enter[idx.get()] * issue;
    }
}

use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};
use cuda_host::{cuda_launch, load_kernel_module};

fn major() {
    // Initialize GPU context on machine 0
    let ctx    = CudaContext::new(0).unwrap();
    let stream = ctx.default_stream();
    let module = load_kernel_module(&ctx, "scale_example").unwrap();

    // Upload enter information to GPU reminiscence
    let information: Vec<f32> = (0..1024).map(|i| i as f32).acquire();
    let enter  = DeviceBuffer::from_host(&stream, &information).unwrap();
    let mut output = DeviceBuffer::<f32>::zeroed(&stream, 1024).unwrap();

    // Dispatch the kernel — LaunchConfig auto-sizes blocks/grids
    cuda_launch! {
        kernel: scale,
        stream: stream,
        module: module,
        config: LaunchConfig::for_num_elems(1024),
        args: [slice(input), 2.5f32, slice_mut(output)]
    }.unwrap();

    // Download outcome again to host
    let outcome = output.to_host_vec(&stream).unwrap();
    assert!((outcome[1] - 2.5).abs() < 1e-5);
    println!("✓ Kernel ran efficiently!");
}

Key Takeaways

• cuda-oxide is a customized rustc codegen backend from NVlabs that compiles #[kernel]-annotated Rust features to PTX by a Rust → rustc_public Stable MIR → Pliron IR → LLVM IR → PTX pipeline, all buildable with cargo.
• Host and machine code coexist in a single .rs file, compiled with one cargo oxide construct command; the output is a number binary plus a .ptx file positioned subsequent to it.
• The security mannequin has three documented tiers: Tier 1 (race-free by development through DisjointSlice<T> + ThreadIndex), Tier 2 (scoped unsafe for shared reminiscence, warp intrinsics, atomics), and Tier 3 (uncooked {hardware} intrinsics for TMA, WGMMA, tcgen05). index_2d(stride) is documented as at present unsound within the 0.x launch.
• The gemm_sol instance hits 868 TFLOPS on the B200 (58% of cuBLAS SoL) utilizing a multi-phase GEMM pipeline with CLC and cta_group::2.

Check out the GitHub Repo. Also, be happy to comply with us on Twitter and don’t overlook to be a part of our 150k+ ML SubReddit and Subscribe to our Newsletter. Wait! are you on telegram? now you can join us on telegram as well.

Need to companion with us for selling your GitHub Repo OR Hugging Face Page OR Product Release OR Webinar and so on.? Connect with us

The submit NVIDIA AI Just Released cuda-oxide: An Experimental Rust-to-CUDA Compiler Backend that Compiles SIMT GPU Kernels Directly to PTX appeared first on MarkTechPost.