CUDA Rust: Write GPU Kernels in Rust Today

CUDA Rust GPU programming is no longer a hobbyist experiment or a wrapper hack. Nvidia has shipped official, first-party support for writing CUDA kernels in Rust, and if you write GPU code for machine learning infrastructure or high-performance compute, this changes your options in a meaningful way.

This post cuts past the announcement hype. We'll look at both programming tracks Nvidia is offering, show real kernel code for each, compare them honestly against C++ CUDA on the dimensions that matter, and tell you which one to reach for depending on your situation.

The Two Tracks: cuda-std vs cuda-core

Nvidia's CUDA Rust support ships in two distinct crates, and conflating them is the first mistake most engineers make when reading the docs.

cuda-std is the higher-level crate. It gives you a Rust-idiomatic API surface that maps onto the CUDA programming model, with thread indexing, shared memory, and synchronization primitives exposed as safe Rust abstractions where possible.

cuda-core is the lower-level crate. It gives you direct, largely unsafe access to the raw CUDA device API. Think of it as a thin Rust skin over the PTX and CUDA C++ runtime.

The choice between them is not about performance. It's about how much control you need and how much undefined behavior you're willing to reason about manually. More on that in a moment.

What the Code Actually Looks Like

A Vector Addition Kernel in cuda-std

The canonical "hello world" of GPU programming is vector addition. Here's what it looks like using cuda-std, pulled directly from Nvidia's cuda-std repository:

#![cfg_attr(target_os = "cuda", no_std)]
#![feature(abi_ptx)]
#![no_main]

use cuda_std::prelude::*;

#[kernel]
pub unsafe fn add(a: &[f32], b: &[f32], c: *mut f32) {
    let idx = thread::index_1d() as usize;
    if idx < a.len() {
        let elem = unsafe { c.add(idx) };
        unsafe { *elem = a[idx] + b[idx] };
    }
}

A few things to call out here. The #[kernel] attribute macro is doing real work: it sets up the PTX function signature and thread dispatch boilerplate that you'd write by hand in CUDA C++. thread::index_1d() maps directly to threadIdx.x + blockIdx.x * blockDim.x. The bounds check on idx < a.len() is something CUDA C++ programmers add manually and frequently forget. Here it's a natural Rust pattern.

The unsafe on the function signature is required because GPU kernels operate outside the normal Rust borrow checker's reach. You're still responsible for ensuring the pointers passed from host code are valid and correctly sized.

The Same Kernel in cuda-core

use cuda_core::prelude::*;

unsafe fn add_kernel(
    a: *const f32,
    b: *const f32,
    c: *mut f32,
    n: u32,
) {
    let idx = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize;
    if idx < n as usize {
        *c.add(idx) = *a.add(idx) + *b.add(idx);
    }
}

This version is almost identical to what you'd write in C++. No macro magic. Every pointer dereference is explicit unsafe. The bounds check is still there, but it's manual. If you're migrating existing CUDA C++ code and want a line-by-line translation path, cuda-core is the track.

Host-Side Launch Code (Shared Pattern)

Regardless of which kernel crate you use, the host side looks like this:

use cust::prelude::*;

fn main() -> Result<(), Box<dyn std::error::Error>> {
    let _ctx = cust::quick_init()?;

    let n = 1024u32;
    let a = vec![1.0f32; n as usize];
    let b = vec![2.0f32; n as usize];

    let a_gpu = DeviceBuffer::from_slice(&a)?;
    let b_gpu = DeviceBuffer::from_slice(&b)?;
    let mut c_gpu = DeviceBuffer::zeroed(n as usize)?;

    let module = Module::from_ptx(PTX, &[])?;
    let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?;
    let kernel = module.get_function("add")?;

    let (grid, block) = kernel.suggested_launch_configuration(n, 0.into())?;

    unsafe {
        launch!(kernel<<<grid, block, 0, stream>>>(
            a_gpu.as_device_ptr(),
            b_gpu.as_device_ptr(),
            c_gpu.as_device_ptr(),
            n
        ))?;
    }

    stream.synchronize()?;

    let c: Vec<f32> = c_gpu.as_host_vec()?;
    println!("c[0] = {}", c[0]); // 3.0
    Ok(())
}

The cust crate handles the host-side CUDA runtime calls. suggested_launch_configuration is a nice quality-of-life wrapper around the occupancy API that C++ developers have to call manually.

Memory Safety: What You Actually Get

This is where the Rust pitch gets complicated, and I want to be direct about the limits.

The borrow checker does not follow your data onto the GPU. Once a pointer crosses the host-device boundary, Rust's ownership system has no visibility into what happens to it. The unsafe blocks in kernel code are not a formality, they're a real boundary.

What Rust does buy you on the GPU side:

• No use-after-free within a single kernel if you stay in safe cuda-std abstractions. The #[kernel] macro enforces that slice references passed in have valid lengths accessible via .len().
• No integer overflow panics in release mode (same as CPU Rust in release). Overflow wraps. Whether that's what you want in an index calculation is your problem to reason about.
• Compile-time type checking across the kernel/host boundary. If you declare a kernel that takes &[f32] and you try to pass a DeviceBuffer<i32>, the compiler catches it.

What Rust does not buy you:

• Race conditions between threads. Shared memory races are still fully your responsibility. __syncthreads() is still unsafe in cuda-std.
• Out-of-bounds global memory access. The GPU will silently read or write garbage, or segfault, same as C++.
• Correctness of your grid/block dimensions. Launch a kernel with the wrong block count and nothing in the type system stops you.

The honest summary: Rust eliminates a class of bugs at the host-side setup layer and gives you better type safety across the kernel boundary. Inside the kernel, you're still writing low-level concurrent code that demands the same discipline as C++.

Benchmarks vs C++ CUDA: What the Numbers Show

Based on published benchmarks from the Rust-CUDA project and community testing on Hacker News threads, the performance picture is roughly:

Workload	Rust cuda-std vs C++	Rust cuda-core vs C++
Vector addition (simple memory-bound)	Within 1-2%	Within 0.5%
Matrix multiply (compute-bound)	Within 2-5%	Within 1-2%
Reduction kernels	3-8% slower	1-3% slower
Shared memory tiling	5-10% slower	1-4% slower

The gap in cuda-std comes from the abstraction layer. The #[kernel] macro generates conservative PTX in some cases, particularly around shared memory alignment and warp-level primitives. cuda-core compiles much closer to what a C++ developer would write by hand.

For most machine learning infrastructure work, the cuda-std gap is noise. If you're writing a custom attention kernel that will run billions of times per day, reach for cuda-core or stay in C++.

One thing the benchmarks don't capture: compile times. Rust's GPU compilation pipeline is slower than nvcc today. Plan for longer CI cycles if you're integrating this into a production ML training stack.

The Migration Path from C++ CUDA

If you have existing CUDA C++ kernels and want to port them, the practical path is:

1. Start with cuda-core. The API surface maps nearly 1:1 to what you already have.
2. Replace raw pointer arithmetic with Rust's offset and add methods. Functionally identical, slightly more readable.
3. Add bounds checks explicitly. In C++ you probably have them in debug builds only. In Rust, make them real.
4. Move host-side setup code to cust. This is where you'll see the biggest ergonomic win. The CUDA runtime C++ API is verbose and error-prone. cust's Result-returning wrappers make error handling composable.
5. Only migrate to cuda-std abstractions if you're writing new kernels from scratch and want the higher-level API.

There's no automated translation tool. This is a manual port. Budget roughly 1.5x to 2x the time it took to write the original C++ code, accounting for learning the Rust GPU compilation model.

Cloud Infrastructure and AI Workloads: Why This Matters Now

The timing of Nvidia's announcement is not accidental. The AI infrastructure layer is under pressure to consolidate around safer, more maintainable code. Training runs for large models like the architectures discussed in DeepSeek-v4.1 Flash's KV cache compression work push custom CUDA kernels to their limits, and the teams writing those kernels are dealing with the same memory bugs and undefined behavior that systems programmers have dealt with for decades.

Rust's ownership model doesn't solve GPU concurrency. But it does reduce the surface area of bugs on the host side, where a surprising number of CUDA production incidents actually originate. Misconfigured buffer sizes, wrong type casts, missing synchronization before readback, these are host-side bugs, and Rust's type system catches most of them at compile time.

For cloud compute workloads where you're renting A100 or H100 time at significant cost, catching a misconfigured kernel launch before it burns a 6-hour training run has real dollar value.

cuda-std vs cuda-core: The Direct Verdict

Use cuda-std when:

• You're writing new kernels and want Rust-idiomatic code
• Your team knows Rust but is newer to GPU programming
• The workload is not performance-critical at the microsecond level
• You want the compiler to catch as many mistakes as possible at the kernel boundary

Use cuda-core when:

• You're porting existing C++ CUDA kernels
• You need maximum performance and are willing to manage unsafe manually
• You're writing warp-level primitives or heavily optimized shared memory code
• You need access to CUDA features that cuda-std hasn't abstracted yet

Don't use either when:

• You need production-grade, heavily optimized kernels today and have no Rust CUDA experience on the team. C++ CUDA with cuBLAS and cuDNN is still the right call for most ML infrastructure teams in 2026.

The C++26 standard is also worth watching here. The C++26 specification is closing some of the undefined behavior gaps that made C++ GPU code hazardous, which means the gap between C++ safety and Rust safety on the GPU side is narrowing somewhat from the C++ direction too.

Getting Started

The official entry point is the Rust-GPU project. Add these to your Cargo.toml:

[dependencies]
cust = "0.3"

[target.'cfg(target_os = "cuda")'.dependencies]
cuda-std = "0.3"
# or
cuda-core = "0.3"

You'll need the CUDA toolkit installed (11.4 minimum, 12.x recommended) and a nightly Rust toolchain. The cuda-builder crate handles the PTX compilation step and integrates with Cargo's build system.

The toolchain is real and it works. The ecosystem is young. If you're building something that needs to ship in 90 days, plan accordingly. If you're making an infrastructure bet for 2027 and beyond, this is the direction worth investing in.