Nvidia CUDA Rust in 2026: GPU Kernels Without C++

For twenty years, “writing a CUDA kernel” meant writing C++. Nvidia has now shipped official CUDA Rust support, and it did so with two separate compilation tracks rather than one compromise — a native rustc backend that emits NVVM IR directly, and a transpilation path that lowers Rust into CUDA C++ before handing it to nvcc. That converts a decade of community effort (Rust-CUDA, cust, cudarc, bindgen wrappers held together by hope) into a supported toolchain with a support matrix and a release cadence. If your infra org is under a memory-safety mandate and your GPU layer was the one place you kept writing raw pointers, that excuse just expired.

What’s actually new about CUDA Rust

The headline is first-class language status. Until now, CUDA officially spoke C, C++, Fortran, and Python (via the CUDA Python bindings and the newer cuda.core/numba-cuda stack). Rust was tolerated at the FFI boundary: you could call CUDA from Rust all day, but the kernel itself — the code that runs on the SM — had to be C++ compiled by nvcc, or LLVM IR you generated yourself with an unsupported fork. Nvidia’s announcement moves the kernel side into the supported column. Bug reports now have a destination, and the compiler is no longer a weekend project that breaks every rustc release.

The two-track design is the interesting engineering choice. Track one is the rustc NVVM backend: a codegen backend that takes Rust MIR through LLVM and emits NVVM IR, the same IR nvcc produces, which then goes through libNVVM to PTX and on to SASS. This descends from the Rust-CUDA project’s rustc_codegen_nvvm, now with Nvidia engineering behind it. You write #[kernel]-annotated Rust functions in a crate targeting nvptx64-nvidia-cuda, and you get a PTX or cubin artifact out. Idiomatic Rust, no C++ in the loop, and — critically — cargo as the build system rather than CMake plus a wall of nvcc flags.

Track two is the Rust-to-CUDA-C++ path, which transpiles a restricted Rust subset into CUDA C++ source and compiles it with nvcc. This looks like a hack until you consider what it buys: instant access to every nvcc-only feature — inline PTX conventions, the full CUDA template libraries, CUTLASS, cooperative groups, tensor-core intrinsics, and new Blackwell-class instructions the moment nvcc supports them, without waiting for an LLVM intrinsic to land. The NVVM backend is the “pure Rust, long-term” answer; the transpiler is the “I need TMA and tcgen05 on Blackwell this quarter” answer. Shipping both admits that neither alone covers the real workload spectrum, and that the C++ ecosystem’s surface area is too large to reimplement.

Why it matters

• Memory-safety mandates finally reach the GPU. CISA and the ONCD memory-safety push made C and C++ a compliance liability in new infrastructure code. Teams rewrote networking, parsers, and daemons in Rust, then stopped at the CUDA boundary because there was no supported option. Now there is, and “the kernel had to be C++” stops being a defensible exemption in a design review.
• One language across host and device. The nastiest bugs in GPU code sit at the boundary, not inside the kernel — a struct laid out differently on each side, a stale FFI signature, a length passed in elements when the kernel wanted bytes. Writing both halves in Rust with shared type definitions kills an entire bug class at compile time.
• Cargo replaces the build archaeology. Anyone who has debugged a CMake + nvcc + separate-compilation + -gencode matrix knows the tax. cargo build with a target triple and a compute capability flag meaningfully cuts the cost of shipping a kernel.
• Rust’s inference-engine layer gets a native floor. Candle, Burn, and the Rust side of the vLLM/TensorRT-LLM ecosystem share the same shape: safe Rust orchestration wrapping C++ or hand-written PTX kernels. A supported path to write those kernels in Rust removes the last unsafe island in otherwise-safe stacks.
• Hiring and review economics shift. A Rust systems engineer can now be productive on GPU work without first learning template-heavy CUDA C++. That widens the pool for kernel work, one of the tightest talent markets in AI infra.
• It pressures the portability layers. If Rust GPU compute is first-class on Nvidia hardware, the “write it in Rust once, run anywhere” pitch from wgpu, Rust-GPU (SPIR-V), and AMD’s HIP story has to compete against a vendor-supported native path with full hardware feature access. Expect AMD and Intel to respond.

How to use CUDA Rust today

1. Confirm the toolkit and driver. Both tracks need a recent CUDA toolkit; check what you have before debugging anything else.

   nvcc --version
   nvidia-smi --query-gpu=name,compute_cap,driver_version --format=csv

2. Install the nightly toolchain and the PTX target. The NVVM backend rides on rustc internals, so it is nightly-pinned. Pin it in the repo so CI and laptops agree.

   rustup toolchain install nightly
   rustup component add rust-src llvm-tools --toolchain nightly
   rustup target add nvptx64-nvidia-cuda --toolchain nightly

   Then commit a rust-toolchain.toml:

   [toolchain]
   channel = "nightly"
   components = ["rust-src", "llvm-tools"]
   targets = ["nvptx64-nvidia-cuda"]

3. Split the workspace into a device crate and a host crate. This is the single most important structural decision. The device crate is no_std and compiles for the GPU target; the host crate is ordinary Rust that loads the resulting PTX.

   cargo new --lib gpu-kernels
   cargo new gpu-host

   In gpu-kernels/Cargo.toml:

   [package]
   name = "gpu-kernels"
   version = "0.1.0"
   edition = "2021"
   
   [lib]
   crate-type = ["cdylib", "rlib"]

4. Write the kernel. A SAXPY, the “hello world” of writing GPU kernels in Rust. Note the shape: thread indexing looks like CUDA C++, but bounds handling and slices are Rust’s.

   #![no_std]
   #![feature(abi_ptx)]
   
   use cuda_std::prelude::*;
   
   #[kernel]
   #[allow(improper_ctypes_definitions)]
   pub unsafe fn saxpy(a: f32, x: &[f32], y: &[f32], out: *mut f32) {
       let idx = thread::index_1d() as usize;
       if idx < out_len(x, y) {
           let elem = &mut *out.add(idx);
           *elem = a * x[idx] + y[idx];
       }
   }
   
   #[inline(always)]
   fn out_len(x: &[f32], y: &[f32]) -> usize {
       if x.len() < y.len() { x.len() } else { y.len() }
   }

   The kernel entry point stays unsafe because the output pointer comes from the host — that is honest, not a wart. Everything downstream of the bounds check is safe Rust.

5. Build to PTX. Target the architecture you deploy on: sm_100 for Blackwell data-center parts, sm_90 for Hopper, sm_89 for Ada.

   cargo +nightly build \
     --target nvptx64-nvidia-cuda \
     --release \
     -Z build-std=core,alloc
   
   # inspect the generated PTX
   ls target/nvptx64-nvidia-cuda/release/*.ptx

6. Load and launch from the host. The host side uses the driver API through a safe wrapper — cust or cudarc, depending on which ecosystem your project already leans on.

   use cust::prelude::*;
   
   static PTX: &str = include_str!(concat!(env!("OUT_DIR"), "/kernels.ptx"));
   
   fn main() -> Result<(), Box<dyn std::error::Error>> {
       let _ctx = cust::quick_init()?;
       let module = Module::from_ptx(PTX, &[])?;
       let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?;
   
       let n = 1 << 20;
       let x = DeviceBuffer::from_slice(&vec![1.0f32; n])?;
       let y = DeviceBuffer::from_slice(&vec![2.0f32; n])?;
       let mut out = DeviceBuffer::<f32>::zeroed(n)?;
   
       let func = module.get_function("saxpy")?;
       let (grid, block) = (((n as u32) + 255) / 256, 256u32);
   
       unsafe {
           launch!(func<<<grid, block, 0, stream>>>(
               3.0f32, x.as_device_ptr(), x.len(),
               y.as_device_ptr(), y.len(), out.as_device_ptr()
           ))?;
       }
       stream.synchronize()?;
   
       let mut host = vec![0.0f32; n];
       out.copy_to(&mut host)?;
       assert_eq!(host[0], 5.0);
       Ok(())
   }

7. Profile it like any other kernel. Rust provenance does not change the tooling — Nsight Compute reads the SASS and does not care what front end produced the PTX.

   ncu --set full --target-processes all ./target/release/gpu-host
   nsys profile --stats=true ./target/release/gpu-host

   Do this early. The first thing you want to know about a Rust-authored kernel is whether register pressure and occupancy match the C++ version you are replacing — that number decides whether the migration is real.

How it compares: Rust CUDA vs C++ CUDA and the alternatives

Approach	Support status	Hardware feature access	Memory safety	Best for
CUDA C++ (nvcc)	Official, mature, 20 years of ecosystem	Complete — day-one Blackwell intrinsics, CUTLASS, cooperative groups	None; manual discipline	Peak-performance kernels, anything needing the newest hardware instructions
CUDA Rust — rustc NVVM backend	Official (new)	Good and growing; lags nvcc on brand-new intrinsics	Rust guarantees on the host and inside kernel bodies	Pure-Rust stacks, shared host/device types, cargo-native builds
CUDA Rust — Rust-to-CUDA-C++	Official (new)	Effectively full — inherits whatever nvcc supports	Rust-level guarantees on a restricted subset	Teams that need bleeding-edge CUDA features without writing C++ by hand
Community Rust-CUDA / cudarc	Community, unsupported	Varies; gaps around newest features	Yes, with unsafe FFI seams	Existing projects; now a migration source rather than a destination
Triton / numba-cuda (Python)	Official / widely adopted	Strong for tiled ML kernels, narrower outside that shape	N/A (managed runtime)	ML researchers iterating on attention and GEMM variants
Rust-GPU (SPIR-V) / wgpu	Community	Vulkan-class compute; no CUDA-specific tensor features	Yes	Cross-vendor portability over peak Nvidia performance

The honest read: if you are chasing the last fifteen percent on a Blackwell GEMM, you are still writing CUDA C++ or CUTLASS, and you will be for a while. If you are writing the other ninety percent of GPU code — preprocessing, custom ops, sampling, quantization, layout transforms, the glue that surrounds the hot kernel — Rust is now a legitimate default, and the maintenance math favors it.

What’s next

The near-term thing to watch is stabilization. The rustc NVVM backend depends on internal compiler APIs, which is why it lives on nightly; the path to stable Rust runs through either the codegen-backend interface stabilizing or Nvidia shipping a pinned toolchain that abstracts the churn. Until that resolves, treat CUDA Rust as production-viable but toolchain-pinned: commit your rust-toolchain.toml, cache your nightly in CI, and expect at least one painful bump per year.

The second signal is how fast intrinsic coverage tracks new silicon. GPU kernel programming in 2026 is dominated by features that did not exist three generations ago — TMA descriptors, thread-block clusters, distributed shared memory, fifth-generation tensor cores, FP4 and FP6 formats. If the NVVM backend gets those within a toolkit release or two of nvcc, the native track becomes the default and the transpiler becomes a compatibility bridge. If it lags by a year, the transpiler will carry most serious workloads and the two tracks will stay permanently split. Watch the intrinsic changelogs, not the marketing.

Third, watch the ML frameworks. The tell that this is real will be Candle or Burn landing a Rust-authored kernel in a hot path and publishing benchmark numbers against their C++ equivalent. That is the moment Rust CUDA vs C++ CUDA stops being a philosophical argument and becomes a table in a PR description. Also worth tracking: whether AMD responds with a supported Rust path for HIP/ROCm, because a Rust-first GPU ecosystem locked to one vendor is a strategically worse outcome than the C++ status quo it replaces.

Frequently Asked Questions

Does CUDA Rust make GPU kernels memory-safe?

Partially, and the distinction matters. Rust eliminates the host-side classes — use-after-free on device buffers, mismatched lifetimes, data races in the orchestration layer — and gives you bounds-checked slices inside kernel bodies. It does not make raw device-pointer arithmetic safe, and it cannot prevent a logic error in shared-memory indexing across threads. Kernel entry points remain unsafe by design. The win is that unsafety becomes a small, auditable region instead of the entire file.

Which of the two tracks should I start with?

Start with the NVVM backend if you are writing new kernels in an existing Rust codebase and your kernels are element-wise, reduction, or scan-shaped. Use the Rust-to-CUDA-C++ track when you need a CUDA feature the native backend has not exposed yet — cooperative groups, the newest tensor-core paths, or anything that requires CUTLASS templates. Many projects will end up using both, per-kernel.

Is there a performance penalty versus CUDA C++?

For straightforward kernels, expect parity or close to it — both paths converge on NVVM IR and the same libNVVM optimizer, so the SASS looks similar. Where you can lose is bounds checks in inner loops and less aggressive inlining across crate boundaries. Profile with Nsight Compute and compare register counts and occupancy against the C++ baseline before you commit to a migration. Do not accept a benchmark you did not run on your own hardware.

What happens to the community Rust-CUDA project?

Its work is the foundation of the native track — rustc_codegen_nvvm, cust, and cuda_std shaped what Nvidia shipped. Existing Rust-CUDA users have the smoothest migration path of anyone, since the programming model and much of the API surface carry over. If you have been running Rust-CUDA in production, you are early, not stranded.

Do I need Blackwell hardware to use this?

No. It targets the same compute capabilities the toolkit does, so Ampere, Ada, and Hopper all work. Blackwell matters because that is where the newest instructions live and where the intrinsic-coverage gap between the two tracks is widest — on older architectures the native backend has less catching up to do.

Can I call existing CUDA C++ libraries from Rust kernels?

From the host side, yes — cuBLAS, cuDNN, cuFFT, and friends have always been callable through FFI, and that is unchanged. Device-side calls into C++ template libraries are where the two tracks diverge sharply: the transpilation path handles it naturally because it emits CUDA C++, while the native backend requires the library to expose a linkable, non-template symbol. Plan your kernel boundaries around this constraint rather than fighting it.