For twenty years, if you wanted a GPU kernel that actually ran fast on Nvidia hardware, you wrote CUDA C++. Rust could call into it through FFI, wrap it, bind it — but the kernel itself, the code that runs on the device, was C++'s territory. That's the whole reason unsafe shows up everywhere in Rust-on-GPU crates today: you're trusting a foreign toolchain you can't verify.
That changed on September 8, when Nvidia shipped CUDA Rust — two open-source projects that compile Rust straight to PTX. Not a wrapper. Not codegen bolted onto nvcc. A native path. It hit Hacker News at 943 points and 395 comments in a single day — this isn't a niche tooling release, it's a shot at one of the deepest moats in tech.
Two tracks, two philosophies
Nvidia didn't pick a lane. CUDA already has two mental models for GPU programming, and they mapped each one to Rust separately.
cuda-oxide — SIMT, explicit control
cuda-oxide is the classic "you own every thread" model that CUDA C++ developers already know. It's a custom rustc codegen backend that routes #[kernel] functions through Rust MIR, into a Pliron IR (an MLIR-like framework, written in Rust), through LLVM, and out as PTX.
#[kernel]
fn vec_add(a: &[f32], b: &[f32], c: DisjointSlice<f32>) {
let i = thread::index();
if let Some(elem) = c.get_mut(i) {
*elem = a[i] + b[i];
}
}
DisjointSlice<T> is the trick: it statically guarantees each thread only ever touches its own element. Try to pass the same buffer as both input and mutable output — a textbook GPU race condition — and you don't get a Heisenbug three weeks into production. You get a compile error:
error[E0502]: cannot borrow `c_dev` as mutable because it is also borrowed as immutable
That's the actual pitch. Rust's borrow checker, which normally polices CPU memory, is now catching GPU race conditions before the kernel ever launches. Shared memory and raw launch configs still need unsafe — Nvidia is upfront that it's "safe(ish)" — but the common failure mode of CUDA (aliased buffers, out-of-bounds thread indexing) moves from a 3am pager alert to a cargo build failure.
Status: early alpha, pinned nightly Rust, CUDA 13.0+, LLVM 21+, Linux only. 3.5k GitHub stars, 39 open issues. Expect breakage.
cutile-rs — Tile, the compiler drives
cutile-rs is the more interesting bet long-term. Instead of thinking in threads, you think in tiles — sub-tensors of data. Each tile block runs your kernel body once, as a single logical unit, over one chunk of the tensor. No thread indexing, no manual shared memory management, because the compiler owns both.
#[cutile::module]
fn vec_add(a: Tile<f32>, b: Tile<f32>) -> Tile<f32> {
a + b
}
That's not a toy example missing boilerplate — that's close to the real thing. The compiler figures out thread mapping, synchronization, and hardware targeting. It runs on stable Rust 1.89+, no nightly toolchain, no custom LLVM. cargo add cutile and you're compiling GPU kernels.
This is already past the toy-demo stage: it's running in production inside Hugging Face's Grout inference engine and in mistral.rs. That's the detail that matters more than the GitHub stars — someone is already trusting this in a serving path.
The part that should worry AMD and Intel
Nvidia explicitly says the two tracks are a response to the same pressure everyone's been talking about for two years: Triton hits 90–105% of hand-tuned CUDA performance while cutting kernel dev time from three days to four hours, and Mojo and ThunderKittens have been chipping at the "you must write raw CUDA C++ to get max perf" assumption. cuTile reads like a direct answer to Triton — except it's Nvidia's own compiler stack, not a third party's, which means it gets first-class support for every future architecture on day one instead of playing catch-up.
That's the actual strategic move here. The CUDA moat was never really "Nvidia has fast GPUs" — AMD's hardware is competitive on paper. The moat is the twenty years of libraries, docs, Stack Overflow answers, and trained engineers that only work if you write C++. By opening a first-class, memory-safe Rust front end, Nvidia isn't weakening that moat — it's widening the front door while keeping the walls exactly as high. Every Rust systems engineer who was previously locked out of GPU work because "GPU means C++" is now a potential CUDA developer. The lock-in migrates from the language to the platform.
Nvidia's roadmap backs this up: they're explicitly planning interop between CUDA Rust, CUDA C++, and CUDA Python, so picking Rust doesn't strand you outside the existing ecosystem. You get memory safety without giving up any of the twenty years of tooling.
What this doesn't do
Don't mistake this for "GPU programming is now safe by default." Shared memory tiling and raw launch configs in cuda-oxide still require unsafe, and that's where most real CUDA performance work lives. Both projects say plainly they are not production-ready and APIs will break. This is Linux-only, needs a pinned nightly for cuda-oxide, and the compile times on first build are rough because the codegen backend itself has to build.
And this doesn't touch AMD or Intel GPUs. It's not a portable Rust GPU story — it's Nvidia extending its own stack, on its own silicon, on its own terms. If you were hoping for wgpu-style write-once-run-anywhere, this isn't it.
Should you try it
If you're already deep in CUDA C++ and shipping kernels that need every ounce of shared-memory control, cuda-oxide gets you memory safety without losing that control — worth watching, not worth migrating production code to yet given "early alpha" is doing a lot of work in that phrase.
If you're writing new GPU code and don't need SIMT-level control, cutile-rs is the one to actually try this week. Stable Rust, no nightly pin, no custom LLVM, and it's already running in production inference engines. cargo add cutile is a genuinely low-cost way to find out if the tile model fits your problem.
Either way, the headline isn't "Rust can now do GPUs" — crates like rust-cuda and wgpu already let you do that. The headline is that Nvidia itself built the compiler, which means Rust just went from "community project tolerated by CUDA" to "first-class citizen funded by the company that owns the hardware." That's the kind of signal that decides which language wins the next decade of systems programming on GPUs.
Sources: Nvidia's official announcement · cuda-oxide on GitHub · cutile-rs docs · HN discussion
Top comments (0)