CUDA for Rust: A Practical Guide to Nvidia's Native GPU Programming Support

Nvidia just gave systems programmers a reason to care about Rust beyond web servers and CLI tools. The announcement of native GPU programming support in Rust — dubbed the “two tracks” approach for writing CUDA kernels — signals that Nvidia is done treating Rust as a third-party curiosity and is now investing in first-class tooling. If you’ve ever fought with unsafe blocks in rust-cuda community crates or dealt with brittle FFI bindings to C++ CUDA code, this changes the calculus significantly.

This is a big deal for anyone building GPU-accelerated infrastructure — ML training pipelines, data processing engines, or custom inference servers — who wants memory safety without giving up raw throughput. Let’s break down what actually changed, how it compares to the existing C++/CUDA workflow, and how to get a kernel running today.

What Nvidia Actually Announced

Nvidia’s “two tracks” strategy refers to two distinct ways developers can now write GPU kernels in Rust:

1. Track One — CUDA Rust (low-level): A near-1:1 mapping to CUDA C++ semantics, giving you direct control over thread blocks, shared memory, warps, and memory coalescing. Think of this as “Rust wearing a CUDA C++ trench coat” — same mental model, safer syntax.

2. Track Two — High-level GPU abstractions: A more ergonomic, iterator-style API (similar to rayon for CPU parallelism) that compiles down to efficient kernels without requiring you to manually manage grid/block dimensions for every operation.

This dual approach mirrors how Rust itself handles systems programming: you can drop into unsafe for full control, or stay in safe, ergonomic Rust for 90% of your code. Nvidia is explicitly targeting both the performance-obsessed kernel author and the application developer who just wants GPU acceleration without becoming a CUDA architecture expert.

Why This Matters Now

For the last decade, GPU programming in Rust meant relying on community projects:

• rust-cuda (via ptx-builder and nvptx64-nvidia-cuda target)
• wgpu for cross-platform compute shaders
• cudarc for safer FFI bindings to the CUDA driver API

These worked, but none had Nvidia’s official backing, meaning no guaranteed compatibility with new CUDA toolkit releases, no first-party debugging tools (cuda-gdb, nsight), and constant risk of breakage across driver updates. Official support means Rust kernels get the same tooling maturity C++ has enjoyed since CUDA’s inception in 2007.

CUDA C++ vs CUDA Rust: A Practical Comparison

Before jumping into code, it’s worth understanding where Rust actually helps and where it doesn’t.

Aspect	CUDA C++	CUDA Rust
Memory safety	Manual, no compiler guarantees	Borrow checker prevents data races in host code
Kernel launch syntax	<<<blocks, threads>>> macro syntax	Explicit function calls with typed launch configs
Build tooling	nvcc + Makefiles/CMake	cargo + rustc with PTX backend
Error handling	Manual cudaError_t checks	Result<T, CudaError> with ? operator
Package management	vcpkg/Conan (fragmented)	Cargo (unified, mature ecosystem)
Debugging tools	cuda-gdb, Nsight (mature)	Nsight support in progress, PTX-level debugging works
Learning curve	Steep — manual memory management everywhere	Moderate — safety rails reduce common bugs
FFI interop with C++ CUDA libs	Native	Requires bindgen or cxx crate
Community crates ecosystem	Massive (cuDNN, cuBLAS, Thrust)	Growing, some gaps remain
Compile-time kernel verification	Limited	Stronger — type system catches more at compile time

The takeaway: Rust doesn’t make your kernels magically faster. The GPU doesn’t care what language emitted the PTX. What Rust buys you is fewer footguns on the host side — the code managing memory allocation, kernel launches, and data transfer between host and device, which is historically where most CUDA bugs live (use-after-free on device pointers, mismatched grid dimensions, forgotten cudaFree calls).

Setting Up Your First CUDA Rust Kernel

Here’s a minimal walkthrough assuming Nvidia’s tooling is installed alongside the standard CUDA Toolkit.

Prerequisites

# Install the nightly toolchain (required for GPU codegen features)
rustup toolchain install nightly
rustup component add rust-src --toolchain nightly

# Add the CUDA target
rustup target add nvptx64-nvidia-cuda

# Verify CUDA toolkit is present
nvcc --version

Project Structure

gpu-vector-add/
├── Cargo.toml
├── build.rs
├── src/
│   ├── main.rs        # host code
│   └── kernel.rs       # device code compiled to PTX

Writing the Kernel (Track One — Low-Level)

// src/kernel.rs
#![no_std]
#![feature(abi_ptx)]

use core::arch::nvptx;

#[no_mangle]
pub unsafe extern "ptx-kernel" fn vector_add(
    a: *const f32,
    b: *const f32,
    c: *mut f32,
    n: i32,
) {
    let idx = nvptx::_thread_idx_x() + nvptx::_block_idx_x() * nvptx::_block_dim_x();
    if idx < n {
        let i = idx as usize;
        *c.add(i) = *a.add(i) + *b.add(i);
    }
}

This looks almost identical to the equivalent CUDA C++ kernel — that’s intentional. Track One prioritizes familiarity for engineers porting existing CUDA codebases.

// Equivalent CUDA C++ for comparison
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

Host Code to Launch the Kernel

// src/main.rs
use cust::prelude::*;
use std::error::Error;

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

    let ptx = include_str!(concat!(env!("OUT_DIR"), "/kernel.ptx"));
    let module = Module::from_ptx(ptx, &[])?;
    let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?;

    let n = 1_000_000;
    let a: Vec<f32> = (0..n).map(|x| x as f32).collect();
    let b: Vec<f32> = (0..n).map(|x| (x * 2) as f32).collect();

    let d_a = a.as_slice().as_dbuf()?;
    let d_b = b.as_slice().as_dbuf()?;
    let mut d_c = DeviceBuffer::<f32>::zeroed(n)?;

    let func = module.get_function("vector_add")?;
    let (grid, block) = (256, 1024);

    unsafe {
        launch!(func<<<grid, block, 0, stream>>>(
            d_a.as_device_ptr(),
            d_b.as_device_ptr(),
            d_c.as_device_ptr(),
            n as i32
        ))?;
    }

    stream.synchronize()?;

    let mut result = vec![0f32; n];
    d_c.copy_to(&mut result)?;

    println!("First 5 results: {:?}", &result[..5]);
    Ok(())
}

Notice the Result propagation with ? — every CUDA API call that could fail returns a typed error instead of a raw cudaError_t you have to remember to check. This alone eliminates a huge class of silent failures in production ML pipelines.

Track Two — The High-Level Abstraction

For teams that don’t need warp-level control, Nvidia’s ergonomic API looks closer to this:

use gpu_compute::prelude::*;

fn main() -> GpuResult<()> {
    let device = GpuDevice::default()?;

    let a = device.upload(&vec![1.0f32; 1_000_000])?;
    let b = device.upload(&vec![2.0f32; 1_000_000])?;

    let c = a.zip(b).map(|(x, y)| x + y).collect(&device)?;

    println!("Sum computed on GPU, first value: {}", c[0]);
    Ok(())
}

No grid/block math, no manual PTX compilation step, no explicit stream synchronization. This is the track most application developers will actually use — think of it as rayon, but the work executes on the GPU instead of CPU threads.

Build Configuration

Cross-compiling to PTX requires a build script to invoke the nightly compiler with the correct target:

// build.rs
use std::process::Command;

fn main() {
    let out_dir = std::env::var("OUT_DIR").unwrap();

    let status = Command::new("cargo")
        .args([
            "+nightly", "rustc",
            "--release",
            "--target", "nvptx64-nvidia-cuda",
            "-p", "kernel",
            "--", "-Z", "build-std=core",
        ])
        .status()
        .expect("failed to build PTX kernel");

    assert!(status.success());
    println!("cargo:rerun-if-changed=src/kernel.rs");
    println!("cargo:rustc-env=OUT_DIR={}", out_dir);
}

# Cargo.toml
[dependencies]
cust = "0.3"

[build-dependencies]

Performance Considerations

Rust’s zero-cost abstractions theoretically hold on GPU targets, but there are real caveats:

• Register pressure: Rust’s iterator chains in Track Two can generate more instructions than hand-tuned CUDA C++ if the compiler fails to fully inline and unroll loops. Always profile with nsight-compute before assuming parity.
• Shared memory access: Track One gives you the same __shared__ semantics via nvptx intrinsics, but the ergonomics are rougher — expect more unsafe blocks than idiomatic Rust elsewhere.
• Warp divergence: Same rules apply as C++ CUDA. Rust’s type system won’t save you from writing branch-heavy kernels that stall warps.
• Occupancy tuning: Grid/block sizing still requires manual tuning based on your GPU’s SM count and register file size — no language abstracts this away yet.

# Profile your kernel exactly like you would with CUDA C++
ncu --set full ./target/release/gpu-vector-add
nsys profile --stats=true ./target/release/gpu-vector-add

Where This Fits in a Real Stack

If you’re running inference servers or data pipelines in Go or Node.js, the practical integration pattern looks like this:

// Go service calling into a Rust CUDA library via FFI
package main

/*
#cgo LDFLAGS: -L./target/release -lgpu_kernels
#include "gpu_kernels.h"
*/
import "C"
import "fmt"

func main() {
    result := C.run_vector_add()
    fmt.Println("GPU computation triggered from Go:", result)
}

Compile the Rust crate as a cdylib, expose a C ABI with #[no_mangle] extern "C", and you get GPU acceleration in services that otherwise have no business touching CUDA directly. This is the same pattern teams have used for years with C++ CUDA libraries — now with a safer implementation underneath.

Common Pitfalls

• Forgetting #![no_std] in kernel code — the device target has no OS, no heap allocator by default, and no standard library support for most std types.
• Mismatched PTX versions — kernels compiled with a newer nvptx64-nvidia-cuda target than your installed driver supports will fail at runtime with an opaque CUDA_ERROR_INVALID_PTX.
• Assuming safe Rust everywhere — kernel entry points and raw pointer arithmetic inside __global__-equivalent functions require unsafe. The safety guarantees mostly benefit host-side orchestration code.
• Ignoring stream synchronization — async kernel launches without proper stream.synchronize() calls will silently read stale device buffers, a bug that’s just as easy to introduce in Rust as C++.

Key Takeaways

• Nvidia’s “two tracks” approach gives you low-level CUDA C++-equivalent control (Track One) and a high-level ergonomic API (Track Two) — pick based on whether you need warp-level tuning or just GPU-accelerated data pipelines.
• CUDA Rust doesn’t make kernels inherently faster than C++; PTX is PTX regardless of source language. The win is safer host-side memory and error handling.
• The cust crate plus official Nvidia tooling replaces fragmented community solutions like rust-cuda and manual bindgen FFI bridges.
• You still need `unsafe