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

Nvidia has announced native GPU programming support for Rust through a "two tracks" approach, offering both a low-level CUDA Rust path that maps closely to CUDA C++ semantics and a high-level, iterator-style abstraction for more ergonomic kernel development. The move gives Rust developers first-party tooling and compatibility with new CUDA toolkit releases, replacing reliance on community crates such as rust-cuda, wgpu, and cudarc. A developer notes that Rust does not make kernels faster but reduces host-side footguns, with the borrow checker preventing data races in host code.

Originally published at adityarawas.in https://adityarawas.in/blog/cuda-for-rust-a-practical-guide-to-nvidias-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. 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 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. 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 . Here's a minimal walkthrough assuming Nvidia's tooling is installed alongside the standard CUDA Toolkit. 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 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. js // 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: php 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: js // 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 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. 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 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 code std 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.