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

Originally published at adityarawas.in

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.

Nvidia's "two tracks" strategy refers to two distinct ways developers can now write GPU kernels in Rust:

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.

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.

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 shaderscudarc 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.

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

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).

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

rustup toolchain install nightly
rustup component add rust-src --toolchain nightly

rustup target add nvptx64-nvidia-cuda

nvcc --version
gpu-vector-add/
├── Cargo.toml
├── build.rs
├── src/
│   ├── main.rs        # host code
│   └── kernel.rs       # device code compiled to PTX
// 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];
    }
}
php
// 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.

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.

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);
}
[dependencies]
cust = "0.3"

[build-dependencies]

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

nsight-compute before assuming parity.__shared__ semantics via nvptx intrinsics, but the ergonomics are rougher — expect more unsafe blocks than idiomatic Rust elsewhere.

ncu --set full ./target/release/gpu-vector-add
nsys profile --stats=true ./target/release/gpu-vector-add

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.

#![no_std] in kernel codestd types.CUDA_ERROR_INVALID_PTX.__global__-equivalent functions require stream.synchronize() calls will silently read stale device buffers, a bug that's just as easy to introduce in Rust as C++.cust crate plus official Nvidia tooling replaces fragmented community solutions like bindgen FFI bridges.