CUDA Rust

NVIDIA's CUDA Rust (announced Sep 8, 2026) compiles GPU kernels written in Rust natively to PTX, instead of writing kernels in C++ and only launching them from Rust. Two tracks mirror CUDA's two models. Both are early-stage and not production-ready; APIs will move.

	SIMT: cuda-oxide	Tile: cutile-rs (crate cutile)
Model	Write what one thread does	Write what one tile (sub-tensor) does; compiler picks threads
Compile	rustc codegen backend: MIR → Pliron IR (NVIDIA GPU dialects) → LLVM IR → PTX, ahead of time	Macro embeds kernel AST; JIT via CUDA Tile IR on first launch per specialization
Toolchain	Pinned nightly (nightly-2026-04-03), clang + libclang headers, CUDA toolkit 12.x+	Stable Rust ≥ 1.89, CUDA 13.3, no LLVM
Platform	Linux, compute capability ≥ 8.0	Linux, compute capability ≥ 8.0
Output exclusivity	DisjointSlice<T>: one element per thread	host .partition() + &mut Tensor: one sub-tensor per tile
Geometry safety	#[launch_contract] + prepare_<k>() token	derived from partition; can't mismatch
Aliasing check	per launch call (E0502)	ownership moves across launch (E0382) — stronger
You control	indexing, block size, shared memory (unsafe today)	nothing about threads/shared memory
Maturity	early alpha	further along; on crates.io; used by HuggingFace Grout, mistral.rs

Rules for using this skill

• Confidence discipline. Public material covers 1-D elementwise kernels only. Prefer API marked Confirmed below. When using anything Inferred (2-D/3-D indexing, scalar args, reductions, math functions, tuple accessors beyond .first()), tell the user in one line which identifiers to verify in the cuda-oxide book / cuTile Rust docs. Never invent versions, flags, or features. With web access, check the repos for anything version-sensitive (nightly pin, CUDA version).
• Name collision: an unrelated, apparently unmaintained cuda-oxide crate exists on crates.io. Install from NVlabs git; never cargo add cuda-oxide.
• Response shapes. Write a kernel: one-line track choice (if not given) → complete single-file program (imports, host + device, verification) → explanation of safety-relevant lines → Inferred identifiers to verify. Port: compact concept map → Rust code → behavioral differences. Debug: identify track, quote key error, explain the invariant, give the fix. "Which one?": recommend, don't survey.

Choosing a track

Default to Tile (NVIDIA's recommendation): the compiler maps tiles to each architecture, so source doesn't encode architecture choices; stable Rust; no threads to race.

Choose SIMT when the user needs explicit per-thread indexing (irregular access, custom stencils), block shared memory, block-geometry control, a structure-preserving port of CUDA C++/numba-cuda, or AOT kernels. Also: CUDA stuck at 12.x → only cuda-oxide; locked to stable Rust → cutile-rs; need maturity → cutile-rs.

Language and model are separate choices: CUDA Tile also exists in C++ and Python; NVIDIA plans cross-language interop.

Rules both tracks share

1. Inputs shared, output has one writer. Read-only: &[T] / &Tensor. Writable: DisjointSlice<T> / partitioned &mut Tensor. &mut [T] can't work for SIMT: every thread would hold the same mutable borrow.
2. Never pass one buffer as input and output of a launch — rejected at compile time whether or not the kernel would race. Use a separate output, or (SIMT) make it the DisjointSlice and read through your own element. Don't defeat the check with unsafe.
3. Bounds are a branch. SIMT get_mut(idx) → Option. Tile load_tile_like aligns input with the output tile.
4. SIMT indices are typed; .get() gives the raw usize for reading inputs only.
5. Always add #[launch_contract] (SIMT); without it only raw unsafe launchers exist.
6. Tile is lazy until .sync_on(&stream); SIMT to_host_vec(&stream) copies and syncs. Time only after sync and warm-up (Tile JIT).
7. Tile launchers consume all tensors and return them as a tuple.

---

SIMT track: cuda-oxide

Setup

rustup toolchain install nightly-2026-04-03
cargo +nightly-2026-04-03 install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide
cargo oxide new vecadd_demo && cd vecadd_demo
cargo oxide doctor   # checks GPU, CUDA, clang/libclang, nightly, optional system LLVM
cargo oxide run      # first run builds the codegen backend: slow; cached after

Host and device code live in one file, one crate, one build command. Expected template output: PASSED: all 1024 elements correct.

Confirmed API

Item	Crate	Role
#[cuda_module] mod m {..}	cuda_host	Device module; generates m::load, prepare_<k>, safe <k> launcher
#[kernel]	cuda_device	GPU entry point
#[launch_bounds(N)]	cuda_device	Max threads per block (register budget)
#[launch_contract(domain = 1, block = (256, 1, 1))]	cuda_device	Indexing dims + block shape
thread::index_1d(), idx.get()	cuda_device	Typed thread index; raw usize
DisjointSlice<T>, .get_mut(idx) -> Option<&mut T>	cuda_device	Per-thread exclusive output
CudaContext::new(0)?, ctx.default_stream()	cuda_core	Context, stream
DeviceBuffer::from_host(&stream, &v)?, DeviceBuffer::<T>::zeroed(&stream, n)?	cuda_core	Upload / zeroed alloc
buf.to_host_vec(&stream)?	cuda_core	Copy down + synchronize
unsafe { m::load(&ctx)? }	generated	Load embedded device bundle
LaunchConfig1D::new(blocks, threads_per_block, shared_mem_bytes)	cuda_core	Geometry
module.prepare_<k>(cfg)?	generated	Validates vs contract + live device limits → proof token
module.<k>(&stream, &prepared, &in.., &mut out)?	generated	Safe launch

Template

use cuda_device::{kernel, launch_bounds, launch_contract, thread, DisjointSlice};
use cuda_host::cuda_module;
use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig1D};

#[cuda_module]
mod kernels {
    use super::*;

    #[kernel]
    #[launch_bounds(256)]
    #[launch_contract(domain = 1, block = (256, 1, 1))]
    pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
        let idx = thread::index_1d();
        let i = idx.get();
        if let Some(c_elem) = c.get_mut(idx) { // None for over-launched tail threads
            *c_elem = a[i] + b[i];
        }
    }
}

fn main() -> Result<(), Box<dyn std::error::Error>> {
    let ctx = CudaContext::new(0)?;
    let stream = ctx.default_stream();
    const N: usize = 1024;
    let a_host: Vec<f32> = (0..N).map(|i| i as f32).collect();
    let b_host: Vec<f32> = (0..N).map(|i| (i * 2) as f32).collect();
    let a_dev = DeviceBuffer::from_host(&stream, &a_host)?;
    let b_dev = DeviceBuffer::from_host(&stream, &b_host)?;
    let mut c_dev = DeviceBuffer::<f32>::zeroed(&stream, N)?;

    // SAFETY: this package owns the embedded device bundle produced for `kernels`.
    let module = unsafe { kernels::load(&ctx)? };
    let prepared = module.prepare_vecadd(LaunchConfig1D::new((N as u32).div_ceil(256), 256, 0))?;
    module.vecadd(&stream, &prepared, &a_dev, &b_dev, &mut c_dev)?;

    let c_host = c_dev.to_host_vec(&stream)?; // syncs
    let errors = (0..N).filter(|&i| (c_host[i] - (a_host[i] + b_host[i])).abs() > 1e-5).count();
    if errors == 0 { println!("PASSED: all {N} elements correct"); }
    else { eprintln!("FAILED: {errors} errors"); std::process::exit(1); }
    Ok(())
}

Notes: keep launch_bounds, contract block size, and threads-per-block equal. div_ceil over-launches the tail; get_mut's Option makes that safe. Guard i < a.len() if inputs can be shorter than the output. Prepare once and reuse the token for the same geometry (Inferred). load is unsafe — keep a // SAFETY: comment.

Patterns

• Elementwise map (Confirmed shape): change the body, e.g. ReLU *o = if x[i] > 0.0 { x[i] } else { 0.0 };.
• In-place update: make the array the DisjointSlice and read your own element: *yi = alpha[0] * x[i] + *yi;. Race-free. Passing a from_host buffer as &mut is Inferred allowed.
• Scalars (Inferred): try alpha: f32; confirmed-API fallback is a 1-element DeviceBuffer read as alpha[0].
• 2-D/3-D (Inferred: domain = 2, index_2d, LaunchConfig2D): safe fallback is 1-D over w*h with (i % w, i / w).
• Gather out[i] = src[idx[i] as usize]: fine (only out written); bounds-check idx[i].
• Scatter out[idx[i]] = ..: not expressible safely by design. Use sort-then-segment, per-thread slots, or documented unsafe.
• Reductions: multi-pass halving (out[i] = in[2i] + in[2i+1] into a half-size buffer), or prefer Tile. Block-local accumulation needs shared memory.
• Shared memory: sized by the third LaunchConfig1D arg; access is unsafe today (safe path is active work). Don't invent the accessor; point to the book.
• Perf: keep data on device across launches; set launch_bounds to the real max; benchmark after sync + warm-up.

---

Tile track: cutile-rs

Setup

cargo new vecadd_demo && cd vecadd_demo
cargo add cutile
cargo run
# repo examples: clone cutile-rs, then
cargo run -p cutile-examples --example hello_world

Confirmed API

Item	Role
use cutile::prelude::*; (host), use cutile::core::*; (inside kernel module)	Imports
#[cutile::module] mod m {..}	Captures AST; generates host launchers
#[cutile::entry()]	Kernel entry
Tensor<T, { [B] }> with const B: i32	Static dim; each B is a specialization
Tensor<T, { [-1] }>	Dynamic dim read at launch; no recompile
load_tile_like(x, z)	Tile of x aligned with current sub-tensor of z
z.store(expr); tx + ty	Write tile; elementwise tile arithmetic
Device::new(0)?, device.new_stream()?	Device, stream
api::ones::<T>(&[n]), api::zeros::<T>(&[n])	Lazy constructors
.partition([w])	Exclusivity + grid (n/w tiles) + supplies B
m::entry(z, x, y)	Launcher: takes ownership, returns tuple lazily
.first(), .unpartition(), .to_host_vec(), .sync_on(&stream)?	Select output, drop wrapper (no data moves), record copy, execute
Error	Result error type

Template

use cutile::prelude::*;

#[cutile::module]
mod kernel {
    use cutile::core::*;

    #[cutile::entry()]
    fn add<const B: i32>(
        z: &mut Tensor<f32, { [B] }>, // exclusive output tile
        x: &Tensor<f32, { [-1] }>,    // shared, dynamic length
        y: &Tensor<f32, { [-1] }>,
    ) {
        let tx = load_tile_like(x, z); // runs once per tile, one logical thread
        let ty = load_tile_like(y, z);
        z.store(tx + ty);
    }
}

fn main() -> Result<(), Error> {
    let device = Device::new(0)?;
    let stream = device.new_stream()?;
    let x = api::ones::<f32>(&[1024]);                   // lazy
    let y = api::ones::<f32>(&[1024]);
    let z = api::zeros::<f32>(&[1024]).partition([128]); // 8 tiles, B = 128

    let c: Vec<f32> = kernel::add(z, x, y)
        .first()
        .unpartition()
        .to_host_vec()
        .sync_on(&stream)?;                              // everything runs here

    let errors = c.iter().filter(|&&v| (v - 2.0).abs() > 1e-5).count();
    if errors == 0 { println!("PASSED: all {} elements correct", c.len()); }
    else { eprintln!("FAILED: {errors} errors"); std::process::exit(1); }
    Ok(())
}

Notes: only mutable tensors are partitioned, and a &mut output must be. Shapes are literals in the confirmed code; the integer type for variable shapes is undocumented. Ragged tails (length not divisible by width) are Inferred/unknown: pad, or choose a dividing width. Errors surface at JIT/sync_on, so insert an earlier sync_on to isolate. Warm up before benchmarking.

Patterns

• Elementwise: change the stored expression. -, *, / likely; math functions (exp, sqrt, maximum) Inferred.
• More inputs: extra &Tensor<T, {[-1]}> params, each load_tile_like(_, z).
• Scalars (Inferred): plain alpha: f32, 1-element tensor, or const generic (recompiles per value — avoid if it varies).
• In-place: not with the same tensor (E0382); allocate a new partitioned output and rebind (lazy, so cheap to record).
• Chaining: pass .first().unpartition() output into the next launcher; one sync_on at the end. Other tuple accessors are Inferred.
• Multi-dim, reductions, attention-style: Tile IR targets these, but the Rust surface isn't shown; study cutile-examples, mistral.rs, and Grout first.

---

Safety model and compile errors

Thousands of threads hit the same buffers in no fixed order; races rarely reproduce and pass tests before failing in production. Both tracks encode "shared inputs, single-writer output" in types (paper and RustConf 2026 talk: Fearless Concurrency on the GPU).

SIMT E0502: module.vecadd(&stream, &prepared, &c_dev, &b_dev, &mut c_dev)? → cannot borrow c_dev as mutable because it is also borrowed as immutable. Fix: distinct output, or read-through-DisjointSlice for true in-place.

Tile E0382: kernel::add(z.partition([128]), z, y) → use of moved value: z. Fix: separate output. Don't .clone() a handle as a workaround unless docs confirm clone semantics — it may re-alias.

Tile's check is stronger because ownership follows tensors across the asynchronous launch; oxide's borrows end when the call returns.

Not guaranteed: numerics/overflow/NaN, OOB reads of shorter inputs (SIMT), correctness of unsafe shared memory, bundle match in load, performance, API stability.

Reviewing unsafe (oxide): require // SAFETY: stating bundle ownership (load); why a raw launch skips contract/prepare; for shared memory, which threads write which regions and why none overlap. Unsafe used only to alias input and output → remove it.

---

Porting

CUDA C++ / numba-cuda → cuda-oxide

Source	cuda-oxide
__global__ void k / @cuda.jit	#[kernel] pub fn k in #[cuda_module]
const float* in	&[f32]
float* out	mut out: DisjointSlice<f32>
int n	usually drop; slices carry length
blockIdx.x*blockDim.x+threadIdx.x / cuda.grid(1)	thread::index_1d() + .get()
if (i < n)	if let Some(o) = out.get_mut(idx)
cudaMalloc+memcpy / cuda.to_device	DeviceBuffer::from_host / zeroed
k<<<b, t, shm>>> / k[b, t]	prepare_k(LaunchConfig1D::new(b, t, shm)) + k(&stream, &p, ..)
D2H + synchronize / copy_to_host	to_host_vec(&stream)
cudaFree / error codes	Drop / Result
cuda.shared.array, per-dtype JIT	unsafe shared memory; static types (generic kernels Inferred)

Triton / CUDA Tile C++/Python → cutile-rs

Triton	cutile-rs
@triton.jit	#[cutile::entry()] in #[cutile::module]
BLOCK_SIZE: tl.constexpr	const B: i32 on output shape
program_id, offsets, mask	implicit per tile / handled by load + partition
tl.load(ptr+offs, mask)	load_tile_like(x, z)
tl.store(out+offs, v, mask)	z.store(v)
grid lambda cdiv(n, BLOCK)	.partition([BLOCK])
eager launch, raw pointers	lazy until sync_on; ownership checks

Simple elementwise SIMT kernels with no shared memory can be "upgraded" to Tile: mention it.

Checklist: classify each array (read / write / read-write at same index) → map writes to exclusive outputs → drop guards types now handle (keep them for shorter inputs) → replace geometry with contract+prepare or partition → replace alloc/copy/free → redesign shared memory, scatter, reductions → verify against a CPU reference with tolerance → benchmark after warm-up + sync → list Inferred identifiers.

---

Troubleshooting

Ask for cargo oxide doctor output (oxide) plus: uname -s, nvidia-smi --query-gpu=name,compute_cap --format=csv,noheader, the CUDA Version in nvidia-smi (driver max), nvcc --version (toolkit), rustup toolchain list, rustc +stable -V, clang --version, and ldconfig -p | grep libclang.

Symptom	Cause → fix
Compute capability < 8.0 (T4 7.5, V100 7.0)	Unsupported on both → cudarc / rust-cuda, or newer GPU
Not Linux	Unsupported; WSL2 untested
cutile fails on CUDA 12.x, or driver max < 13.3	cutile needs CUDA 13.3 toolkit and driver; else use cuda-oxide
Container sees no GPU	--gpus all / NVIDIA Container Toolkit; CUDA_VISIBLE_DEVICES
First cargo oxide run very slow	Building the backend; normal, cached after
rustc_private errors / ICEs (oxide)	Wrong nightly; install with the pin; check rustup override isn't shadowing the project toolchain
libclang not found	Install clang + libclang dev package (Debian/Ubuntu clang libclang-dev); set LIBCLANG_PATH
cargo add cuda-oxide pulled the wrong crate	Remove it; use NVlabs git + cargo oxide new
Valid Rust fails in #[kernel]	Backend coverage incomplete: remove dyn, heap (Vec/Box/String), panics/formatting, recursion, iterator chains, std calls; file a minimal issue
Only unsafe launchers generated	Add #[launch_contract]
prepare_* returns Err	Threads/block must match contract, be ≤ launch_bounds and device max (1024)
Unpartitioned &mut output error (Tile)	.partition([w]) every mutable tensor
Wrong tuple element / type (Tile)	Launcher returns all args; select the output's position
Nothing happens / zero timings (Tile)	Lazy until sync_on
Wrong values, no error	Ragged partition, missing input guard, float tolerance; shrink N to 16 and compare with CPU
Slow first Tile launch	JIT per specialization; changing width/static shapes recompiles

Minimal repro: start from the template, change one thing at a time, shrink N. Report GPU + compute capability, driver, toolkit, rustc -Vv, doctor output, full error.

---

Ecosystem and status

Rust already runs much of the AI systems layer; NVIDIA uses it in the Nova driver, Dynamo's core, and NVTX bindings. CUDA Rust fills the missing piece: the kernel itself. Non-NVIDIA project descriptions reflect early-to-mid 2026; the cuda-oxide book's ecosystem appendix is authoritative.

Project	What	Pick when
cutile-rs	Tile frontend, JIT, NVIDIA	Default for new NVIDIA-only Rust kernels on stable
cuda-oxide	SIMT codegen backend, NVIDIA	Need thread-level control; accept nightly/alpha
cudarc	Safe host bindings: driver, cuBLAS, cuDNN, NVRTC (kernels usually C via NVRTC)	Host-side libraries, older GPUs, stability; complements CUDA Rust
rust-cuda	Community Rust → NVIDIA via NVVM (NVIDIA collaborating)	Existing projects, broader GPU generations
Rust-GPU	Rust → SPIR-V	Vulkan, cross-vendor
CubeCL	Embedded GPU compute language (Burn), multi-backend JIT	Portability across CUDA/ROCm/WGPU

Portability needed → CubeCL / Rust-GPU. Pre-Ampere → cudarc / rust-cuda. Production today → cudarc + C++ kernels, evaluating cutile-rs in parallel. Buffer interop between these and CUDA Rust is undocumented; flag it.

Roadmap signals: NVIDIA is growing CUDA Rust into 2027+. Stated goals are removing the nightly pin, safe SIMT shared memory, and cross-language interop.

Resources (don't fabricate other URLs): github.com/NVlabs/cuda-oxide and the cuda-oxide book; the cutile-rs repo, cuTile Rust docs, and cutile on crates.io; GitHub Discussions on both; the cuda-oxide Discord; the paper Fearless Concurrency on the GPU; NVIDIA Technical Blog, "Introducing CUDA Rust: Two Tracks for Writing GPU Kernels" (Sep 8, 2026).