NVIDIA opens two Rust paths for CUDA kernels — SIMT thread model vs tile-level programming — What does it mean?
The news. On September 8, 2026, NVIDIA introduced two open-source projects for writing CUDA kernels in Rust. cuda-oxide exposes the familiar SIMT thread model and compiles Rust kernels to PTX through a custom
rustcbackend. cutile-rs exposes tile-level programming instead: its macro embeds the kernel's parsed form — its AST, or abstract syntax tree — in the host binary, which is then JIT-compiled through CUDA Tile IR. NVIDIA says both projects are early-stage and neither is production-ready. Read the announcement →
Picture a crew of painters in front of a blank wall. In the SIMT track you brief them one at a time: painter 417 covers this square inch, painter 418 the next. When you shout a stroke command, every painter in a row follows it together — and if two painters need different instructions, one of them stands idle while the other's stroke is made. Nothing gets painted twice only because you worked out in advance that no two painters were sent to the same spot. The hardware has always worked this way; what NVIDIA changed is that you can now write those per-painter orders in Rust instead of CUDA C++.
In the tile track you never talk to a painter at all. You chalk the wall into panels, hand each panel to a crew, and let the foreman decide how many people work it and in what order they stroke. You describe what happens to one panel; the compiler works out the rest. cutile-rs does exactly this — it partitions a mutable tensor into non-overlapping tiles and lets the compiler pick the physical thread mapping and the memory layout, including whether neighbouring threads end up reading neighbouring addresses.
The difference is not primarily about speed. It is about who is responsible for the mapping — and therefore who has to be right about it.
GPU Abstraction Stack
Write Python ops — runs immediately, one kernel per op
Annotate with @torch.compile — compiler fuses & optimizes for you
Write tile-level kernels in Python — full control, no C++ needed
Write thread-level kernels — maximum control, maximum complexity
That ladder already had both rungs before this week. CUDA C++ sits on the thread rung; Triton sits on the tile rung. The two Rust projects land on exactly those same two rungs — cuda-oxide beside CUDA C++, cutile-rs beside Triton — so what is new here is the language, not the model. If you have written a Triton kernel, you have already written tile code. If you have ever typed blockIdx.x * blockDim.x + threadIdx.x, you have already written SIMT code, and you have already met the expertise gap that the tile rung exists to close.
The two tracks also take different routes down to the GPU. cuda-oxide pushes a function marked as a kernel through Rust's own MIR, then Pliron — a compiler-infrastructure layer — then LLVM IR, and finally out as PTX — a build-time path that ends in the same virtual assembly the CUDA compiler emits. cutile-rs does something less usual: its macro embeds the kernel's AST — the parsed shape of the code — in the host binary, and the kernel is compiled on first use through CUDA Tile IR. Both paths end at machine code for your GPU; they differ in when the thread mapping is decided, and by whom.
The types are where the Rust part earns something rather than just being a different syntax. cuda-oxide carries launch contracts and a DisjointSlice type that enforce bounds and exclusive writes, so the property lives in the type rather than in the author's head. cutile-rs reaches something similar from the other direction: because its tiles are non-overlapping by construction, there is no way to express two crews on one panel.
| Track | What you write | Who maps threads | Compile path | Requires |
|---|---|---|---|---|
cuda-oxide | per-thread SIMT code | you | MIR → Pliron → LLVM IR → PTX, at build time | compute capability 8.0+, CUDA 12.x+ (source) |
cutile-rs | one tile body over non-overlapping slices | the compiler | embedded AST → CUDA Tile IR, JIT on first use | compute capability 8.0+, CUDA 13.3, stable Rust 1.89+ (source) |
Take the release's own example: 1,024 elements, partitioned into eight 128-element tiles.
On the SIMT track you launch one thread per element. Say you pick 256 threads per block (illustrative — the block size is your choice, not the release's): that is 4 blocks, and at 32 threads to a warp, 32 warps. None of those three numbers is the one that costs you anything, though. What you actually hand over is a single index formula, blockIdx.x * blockDim.x + threadIdx.x, and it has to be right for every one of the 1,024 values it produces. Here 1,024 divides evenly by 256, so there is no ragged end; at a size that is not a multiple of the block size you would also owe a bounds check for the threads whose indices land past the array.
On the tile track you write one tile body and declare eight 128-element tiles instead. 0–127, 128–255, and so on up to 896–1,023; the compiler picks the thread count and the memory layout underneath them. The unit you are responsible for changes from one formula that has to hold across 1,024 values to 8 ranges that do not overlap by construction. That is the trade — less control over the physical mapping, in exchange for a boundary you can see.
Neither project is production-ready — NVIDIA says so plainly — so the durable thing to take from this release is the distinction, not the tooling. Rust may or may not become a normal way to write kernels. The two models underneath it are not going anywhere, and knowing which rung a piece of GPU code sits on tells you immediately what it is your job to get right.
Goes deeper in: GPU & CUDA → Execution Model → Threads, Blocks, and Grids
Related explainers
- AsmEvo — correctness-gated assembly optimization — what the stack looks like one layer below PTX.
- NVIDIA spec decode — tile-aligned draft length — a tile boundary showing up as a serving decision rather than a kernel one.