cuTile Rust 0.4.0

0.4.0 is built on CUDA 13.4 Tile IR. These notes cover changes since 0.3.0.

Profiling and debugging

• Dev-profile builds compile kernels with full device debug information at O0. Release builds use none. CUDA_RUST_DEBUG=none|line|full overrides the profile at build time; explicit CompileOptions take precedence (#285).
• CUDA_RUST_DEBUG=line produces optimized device code with source-line tables. Nsight Compute and Nsight Systems attribute SASS to the .rs line that produced it. Nsight Compute 2026.3 adds a Tile IR Source view.
• Helpers retain their source locations and subprograms, including calls across modules and files. The recorded call-site chain gives cuda-gdb backtraces such as assert_in_range / validate_scale_factor / your_kernel, with stops at the requested line (#238, #285, #311).
• Per-launch CompileOptions::{device_debug, lineinfo, sanitize_memcheck, opt_level} map to --device-debug, --lineinfo, --sanitize=memcheck (for compute-sanitizer --tool memcheck) and --opt-level. All these options participate in kernel cache keys (#236).

CUDA_RUST_DEBUG=line cargo build --release      # optimized, with source lines
ncu --set full -o report ./my_cutile_program     # Nsight Compute
nsys profile ./my_cutile_program                 # Nsight Systems
CUDA_RUST_DEBUG=full cargo build && cuda-gdb --args ./target/debug/my_program

The Debugging and Profiling chapter walks through each tool.

DGX Spark

cuTile Rust runs on the DGX Spark's GB10 (sm_121, aarch64) with CUDA 13.3 and 13.4 on the 580 driver. The Tile IR assembler emits native sm_121 code, and the default workspace builds on generic aarch64 (#295). The full GPU test suite passes on the Spark.

Inference on DGX Spark runs Qwen3 in 16-bit precision with Grout, the cuTile Rust inference engine (#303).

Model, 16-bit weights	Decode	Weights in memory
Qwen3-4B	29.5 tokens/s	8 GB (15 GB in use)
Qwen3-32B	3.7 tokens/s	66 GB

Both run at roughly 90% of the LPDDR5X roofline. Qwen3-32B sustains 245 GB/s against a 273 GB/s peak. The model fits in the Spark's 128 GB of unified memory. The tutorial covers toolkit selection on the 580 driver and per-op profiling with GROUT_PROFILE_OPS=1. An sm_121 tuning pass helps prefill, not decode.

Programmatic dependent launch

programmatic_dependent_launch() on a generated kernel builder lets the kernel start its prologue while the previous kernel on the same stream finishes (#298). Enable it on the consumer:

unsafe {
    kernels::producer(input.device_pointer()).grid((1, 1, 1))
        .then(|_| kernels::consumer(input.device_pointer(), output.device_pointer())
            .programmatic_dependent_launch().grid((1, 1, 1)))
}.sync_on(&stream)?;

The producer chains its stores into gdc_launch_dependents_tko(Some(stored)) when dependents may begin. The consumer calls gdc_wait_tko(None) and passes the returned token to every load that depends on the producer, via Tensor::set_token for tensors or directly for pointer operations. The launch signal permits scheduling; data is ready after the wait.

PDL requires Tile IR 13.4, an sm_90 or newer GPU and a driver with cuLaunchKernelEx. Unsupported requests fail instead of silently serializing. The unsafe per-launch contract requires work before the wait not to race the producer. See the host API reference for details. cutile-examples/examples/pdl.rs implements a producer/consumer pair and compares stream-order and PDL timings.

Also on CUDA 13.4

Tile IR 13.4 also adds insert, fpowi, fpowf, alias fencing, f8e5m3fnu on sm_107+, saturating float-to-int, explicit pointer classification, view inbounds, NearestAway rounding, and kernel return inside loop/while. The bytecode writer negotiates the assembler version and checks each feature against the target architecture before assembly. A 13.3 toolkit builds what it supports and rejects the rest with a source-located error (#298). The README lists each feature's toolkit and architecture requirement.

Since 0.3.0, in brief

cargo add cutile is sufficient for a kernel crate (#248). Kernels gained program_id(axis)/num_programs(axis), *const T raw pointers with Rust's capability model, x.load_like(z) and shape![..] (#240, #259, #261). Cache-hit launches dropped from 3.9 µs to 1.6 µs of host time. The launcher no longer allocates on that path; a four-tensor launch retains two allocations in the submission envelope. A 32-node graph replay enqueues in 1.7 µs (#265, #302, #307). The host-side crates gained cuda-oxide's SIMT surface, CUDA VMM wrappers and NVLink switch multicast (#202, #233). Toolkit discovery finds /opt/cuda, CUDA_HOME layouts and WSL2's driver shim (#190, #251, #277). Kernel-cache eviction works without the experimental autotuner (#282).

Bug fixes

• The async runtime prevents use after free from safe code by retaining tensor storage until work completes, including after discarded outputs, errors, panics and forgotten futures. Conflicting cross-stream accesses are rejected at enqueue. Graph replays reacquire their captured buffers (#275, fixes #272 and #252).
• Loads through a read-only Partition were chained on each other's tokens, serializing load-bound kernels by up to 35% (#310).
• Line info for kernels calling a trait-dispatch wrapper as a free function pointed at the module line (#311).
• f8 constants round to nearest-even instead of truncating (#275, fixes #250).
• Pointer alignment hints use the full 64-bit address (#244).
• Integer division in kernels rounds toward zero as in Rust (#225).

Breaking changes

0.4.0 is the breaking release announced in the 0.3.1 notes. Global takes a sealed atomic type (Global<AtomicI32, { [] }>). Tensor::store returns its Token, and the raw token helpers are unsafe. Latency takes an i32. Hand-written DeviceOp::execute implementations must retain what they own. Dev-profile builds default to device debug at O0. The crate family is version-pinned in lockstep. The CHANGELOG includes the 0.3.1 items.

Requirements

CUDA 13.2 or newer for the Tile compiler, 13.4 for PDL and the other 13.4 features; GPUs from sm_80 up (sm_90+ for PDL, sm_100+ for FP4 and block-scaled MMA, sm_107 for f8e5m3fnu).

cuTile Rust 0.3.1

0.3.1 makes getting started one line — cargo add cutile is all a kernel crate needs — and brings profiling support, enhanced autotuning, and a faster launch path.

Highlights

• One-crate install: generated launchers resolve compiler support through the cutile crate root, so cargo add cutile alone builds and runs kernels (#248).
• Profiling support: CompileOptions::{device_debug, lineinfo, sanitize_memcheck, opt_level} (#236) and source locations in the bytecode debug section (#238), so Nsight Compute/Systems attribute SASS to your kernel's .rs lines. See Debugging and Profiling.
• Enhanced autotuning (experimental): resumable trial logs on the engine-objective path, required incumbent configs, kernel-cache management for sweeps, and a dedicated book chapter: Autotuning (#239).
• Faster launches: cache-hit launches cost ~2.4x less host time (3.9 → 1.6 µs measured) via per-site resolution caching and allocation-free marshalling (#265).
• Kernel API: program_id(axis)/num_programs(axis) (Triton/cuTile Python vocabulary), *const T raw pointers with Rust's safety model and Tensor::as_ptr, x.load_like(z), and shape![..] (#240, #259, #261).
• Host-side crates: cuda-oxide's host surface (simt modules, #[derive(DeviceCopy)]), CUDA VMM wrappers, NVLink switch multicast (#202, #233, #251).

Breaking changes in a patch release

0.3.1 contains breaking API changes that semantically belonged in 0.4.0 — several soundness fixes (unsafe fn atomics and pointer stores, unsafe trait DType) and one silent semantics fix: integer / in kernels now rounds toward zero as in Rust, and u8/u16 arithmetic is properly unsigned, so programs that relied on the old lowering produce different results. The full list is in the CHANGELOG.

To keep this from recurring: the crate family is now version-pinned in lockstep so cargo can never mix releases (#267), CI runs cargo-semver-checks against the latest published release so an API break without the matching version bump fails the PR, and the next breaking release will bump the minor version.

Requires CUDA 13.2+.

cuTile Rust 0.3.0

This release introduces compiler optimizations that improve safe kernel performance, async launch performance improvements, persistent kernel tuning, and zero-copy interop with externally owned device memory. The front-end compiler now proves a broader class of symbolic facts about loads and stores during JIT compilation. Checks which the compiler cannot discharge outright are progressively hoisted out of control flow constructs (e.g. hot loops) and, when possible, out of the kernel entirely, re-materializing as single host-side comparisons at launch. This frees kernel code from bounds checks and bounds-related register pressure.

Highlights

• Bounds-check placement: Every partition access walks one proof ladder (axis provenance, static folding, declared preconditions, host-side launch checks) and only then pays an in-kernel assert, hoisted to the outermost provable loop preheader. Kernels indexed by loop iterands, mapped schedule components, or tile-block ids typically compile with zero in-kernel checks and no annotations, and deny_in_kernel_checks = true turns any residual in-kernel check into a compile error. PartitionMut::store joins the checked-by-default surface (store_unchecked is the explicit escape) at identical register counts. The placement optimizer is differentially validated: Every kernel in the test corpus runs both with optimized placement and with all checks forced to their access sites, and the optimized build must never accept an execution the reference build rejects. The new book chapter, Bounds-Check Placement, documents the placement rules and the verification workflow.
• Partial partition mapping: A partitioned output normally maps every block to a tile program; partition_prefix allows a partial, still-disjoint mapping — the leading blocks along each axis get tile programs, the rest are left untouched. Mapping beyond the partition remains a launch error, and partition still requires the mapping to be total.
• Zero-copy interop: Tensor::from_foreign borrows device memory owned by an external framework (cudarc, torch, VMM ranges), holding the owner alive by refcount and verifying the addressable extent at construction; Tensor::borrow_raw_parts, Device::borrow_with_owner, and Stream::borrow_with_owner cover bare handles.
• Autotuner core: declared configuration spaces, pluggable searchers, resumable grid search, and provenance-checked tuning records.
• cutile::bench for device-event kernel benchmarking, an opt-in persistent on-disk cubin cache, a process-global single-flight kernel cache with a .compile() warmup API, and async launch latency at parity with sync.

Performance results

• Fully checked kernels match their zero-check unsafe twins: Paired A/B runs on B200 (Qwen3-32B, 16K and 32K prefill) agree within 0.5% in both directions with byte-identical output, and on RTX 5090 the fully checked flash-attention prefill kernel runs within 2.5% of its checks-disabled floor. Checked and unchecked stores compile to identical register counts.
• Async launch latency is at parity with sync.

Changed

• with_bounds and Dim::new are deprecated: bounds inference and declared preconditions prove everything they proved, with the checks placed at launch instead of in the kernel. Migration notes are on the deprecations.
• Tile IR bytecode version selection reads the CUDA toolkit when discoverable and otherwise probes tileiras with a representative module (an entry with a for region and view chain), so an accepted version is one real kernels compile at.
• The persistent_gemm example is rewritten in the zero-annotation form and compiles under deny_in_kernel_checks = true.

Fixed

• An entry declared as bare #[cutile::entry], without an argument list, was silently treated as a plain function; both attribute forms are now recognized.
• Interval range facts no longer survive arithmetic that can wrap i32; a wrapped intermediate could previously discharge a check the machine value violated.
• Bounds analysis guards division and remainder against zero divisors, saturates interval arithmetic, and uses overlap analysis for eq/ne.
• Device-op Send bounds, CUDA driver flag types, and aarch64 flake support.
• Security bumps for crossbeam-epoch (RUSTSEC-2026-0204) and anyhow (RUSTSEC-2026-0190).

cuTile Rust 0.2.0

cuTile Rust 0.2.0 adds low-precision inference support and accompanies our paper, Fearless Concurrency on the GPU.

Paper: https://arxiv.org/abs/2606.15991
Artifacts: cutile-benchmarks/paper/
Announcement: NVlabs/cutile-rs#164

Highlights

• CUDA 13.3-oriented low-precision support, including NVFP4 packing and unpacking and block-scaled matrix multiply.
• Runnable NVFP4 and MXFP8 examples.
• New cutile-kernels crate with reusable inference-oriented kernels written in cuTile Rust.
• Paper reproducibility artifacts under cutile-benchmarks/paper/, including benchmark harnesses, CSV results, hardware details, and clock settings.
• Expanded compile-only regression coverage and example/test cleanup for release validation.

Performance context

On NVIDIA B200, cuTile Rust reaches 7 TB/s for element-wise operations and 2 PFlop/s for GEMM, about 91% of peak memory bandwidth and 92% of dense f16 peak, respectively. Safe Rust persistent GEMM reaches 2.07 PFlop/s at M=N=K=8192, within 0.3% of the corresponding low-level Tile IR variant.

We also evaluated Grout, a Qwen3 inference engine built with cuTile Rust in collaboration with Hugging Face. In batch-1 Qwen3 decode, Grout reaches 171 tokens/s for Qwen3-4B on NVIDIA GeForce RTX 5090 and 82 tokens/s for Qwen3-32B on B200, showing strong performance on memory-bound inference tasks, consistent with our HBM roofline analysis.

Notes

cuTile Rust remains early-stage research software. CUDA 13.3 is required for the new low-precision features, and hardware-specific features such as native FP4 require architectures that support them.

cuda-oxide v0.2.1

Patch release: correctness fixes in the device pipeline, plus a licensing cleanup in cuda-bindings that supersedes v0.2.0 for anyone consuming that crate.

Correctness

• Device functions and call sites now carry LLVM's convergent attribute, preventing illegal transforms around warp-synchronous operations (#145).
• The detected GPU architecture is treated as a hint, not a hard target, so explicitly requested --arch builds behave as asked (#145).
• NaN float literals are emitted as hex bit patterns instead of bare nan, which llc rejects (#116, ported from #63).
• Loads, stores, and allocas now carry Rust-side ABI alignment instead of LLVM defaults (#122, ported from #113).

cuda-bindings licensing cleanup

• crates/cuda-bindings ships under the NVIDIA Software License and accepts no external contributions. Two externally-authored patches that briefly landed there have been removed and their functionality re-implemented under NVIDIA authorship (#152): the CUDA_HOME toolkit fallback and the CUDA 12.8 cuEventElapsedTime_v2 compatibility shim.
• The FFI bindings themselves were never affected: they are generated at build time by bindgen from your local CUDA headers and are not checked in.
• The crate's SPDX headers now match its declared license, and CI now rejects non-NVIDIA-authored changes to the crate (#152).

Testing, CI, and docs

• 14 new mir-lower control-flow unit tests (#111, thanks @goog00).
• Error-demo examples are now classified in STATUS.md and enforced in CI (#86, thanks @ronakv).
• gemm_sol gained a live cublasLt speed-of-light baseline for perf work.
• thread::index_1d / index_2d uniqueness is documented as launch-conditional (#127).

v0.2.0 - first community release

cuda-oxide v0.2.0 — first community release

cuda-oxide is an experimental rustc backend that compiles CUDA GPU kernels written in pure Rust to PTX, with host and device code in a single source tree built through cargo oxide. It owns the SIMT side of the Rust GPU story (complementary to the Tile DSL work).

The headline is momentum. We open-sourced v0.1.0 on May 7 as a solo alpha. In the four weeks since, the project has taken 37 pull requests from 23 contributors, spanning compiler correctness, new device capabilities, host-runtime safety, packaging, and documentation. v0.2.0 is assembled almost entirely from that contributed work.

What landed, at a high level

• Reach. Kernels now embed directly in the host binary with a typed launch API, so a Rust GPU program is a single self-contained executable.
• New capability. CUDA constant memory and additional device math (max / min, atan) landed through community PRs, widening what real kernels can express.
• Trust. A class of silent miscompiles now fails loudly at build time rather than producing wrong results, and the host runtime gained explicit memory-safety contracts.
• Onboarding. A dev container and a Nix flake make first-build friction near zero, which is what is converting drive-by interest into PRs.
• Interop. A working cuTile ↔ SIMT inter-kernel example demonstrates the Tile-to-SIMT handoff on a shared stream, the integration point we own in the broader CUDA-Rust strategy.
• Upstream alignment. Migrated the LLVM dialect onto upstream pliron-llvm, retiring our in-tree fork in favour of tracking the community project.

Highlights

New device & compiler features

• CUDA constant memory via #[constant] / Constant<T> (PTX .const, addrspace 4) with generated host setters.
• New device math via libdevice: f32::max / f32::min, and atan / atan2 for f32 and f64.
• A globaltimer() device intrinsic for cross-SM timing.
• TypeId-based kernel PTX naming and a single-byval kernel-entry ABI, fixing generic-closure and aggregate-parameter launches.

Embedded kernels & a typed launch API

• A new oxide-artifacts bundle format embeds generated PTX / NVVM-IR / LTOIR / cubin directly into the host binary, so kernels load straight from the running executable with no separate .ptx file.
• A #[cuda_module] proc-macro generates a typed, per-kernel launch surface (&[T] maps to &DeviceBuffer<T>), and by-value closure launches now marshal correctly across Fn / FnMut / FnOnce.

Correctness & soundness

• Several silent miscompiles now fail loudly at build time instead of producing wrong PTX (dropped wgmma accumulate, elided destructors, swallowed statements).
• Fixed niched Option<NonZero<T>> / Option<&T> Transmutes (the step_by miscompile), tuple-returning device functions, and arithmetic-trait Output resolution on aggregate operands.
• Device-codegen panics now print a clean cuda-oxide diagnostic instead of a rustc ICE banner.

Host runtime

• A DeviceCopy safety contract on all device transfers, pinned (page-locked) host buffers with stream-ordered transfers, propagated async errors, and a raw CUmodule escape hatch for driver interop.
• Fixed a DeviceBuffer allocator mismatch and added CUDA 12.8 / 13 and aarch64 compatibility.

Compiler internals

• Migrated onto upstream pliron-llvm. The in-tree dialect-llvm crate has been retired and replaced with the community pliron-llvm dialect from pliron-org/pliron; the local crate is renamed dialect-llvm → llvm-export to reflect its narrower remaining role (textual LLVM-IR / bitcode emission for libnvvm and llc). pliron is pinned to a cuda-oxide-llvm-additions branch so we can land additions upstream incrementally while keeping the trunk aligned.
• The textual LLVM-IR exporter was refactored into a module (byte-identical output), with correctness fixes for global-symbol (addressof) export and new select emission, plus expanded mir-lower test coverage.

Build, CLI, and developer experience

• cargo oxide now respects your RUSTFLAGS, nvcc host compiler, and CUDA_HOME; resolves llc from the Rust toolchain; auto-detects the run target GPU; and invalidates a stale backend cache on upgrade.
• A reproducible dev container and a Nix flake (nix run .#new) for one-command onboarding.
• A cuTile ↔ SIMT inter-kernel interop example.

Documentation

• A complete LLVM 21 / clang install path for Ubuntu 24.04 and DGX-OS / DGX Spark, plus book and troubleshooting refreshes.
• Doctest failures and rustdoc warnings fixed and gated in CI.

Contributors

Thanks to everyone who shipped code or docs into v0.2.0:

• Alejandro Soto Franco (@alejandro-soto-franco)
• Alex Afshar (@afsharalex)
• Chronara Group (@CG-8663)
• Daniel Lindestad (@lindestad)
• Eden Gilbert Kiseka (@edengilbertus)
• Fabio Baravalle (@fbrv)
• George Amponis (@g-ampo)
• Hiroki Taniai (@nyoki-mtl)
• Ibrahim Rayamah (@iBz-04)
• Jonatan Lindh (@JonatanLindh)
• Joseph Glanville (@josephglanville)
• Magyar Máté András (@Barrow099)
• Mohamed Samir (@mohamedsamirx)
• Onyeka Obi (@MavenRain)
• Paul Zhang (@paulzhng)
• Rob Parolin (@rparolin)
• Robert Forsman (@mutantbob)
• Shaurya Mishra (@winshaurya)
• Stephen Zhang (@hsqStephenZhang)
• sunteng (@goog00)
• Tim Dortmann (@tdortman)
• ur4t (@ur4t)
• Vito Lin (@9-8-7-6)

Status

Early alpha, git-installed (not yet on crates.io), Linux only.

cuTile Rust 0.1.1

This is a small 0.1.x release for NVFP4 Tile IR support and documentation updates.

• Added NVFP4 type and op coverage across Tile IR, compiler lowering, examples, tests, and book material.
• Added versioned book build and publish support.
• Bumped the workspace crates to 0.1.1 and refreshed CI/docs.

cuTile Rust 0.1.0

cuTile Rust v0.1.0 release

cuTile Rust v0.1.0 is now available on crates.io.

We've finalized our host-side and device-side APIs. We
are not planning any further breaking changes to the kernel authoring model, tensor
launch API, DeviceOp execution model, or core device operation surface.

It is also much easier to try cuTile Rust in a normal Rust project:

cargo add cutile@0.1.0

or add it directly to Cargo.toml:

[dependencies]
cutile = "0.1.0"

That should be enough for normal kernel authoring; you should not need to
depend on the internal workspace crates directly.

Highlights in v0.1.0:

• Finalized host and device APIs for tensor kernels, DeviceOps, async
  execution, CUDA graphs, and CUDA interop.
• Device-side operations now closely track Tile IR, including tensor views,
  partition views, atomics, memory ordering, shape operations, and tile math.
• The JIT compiler now has stronger type inference, static dispatch lowering,
  type aliases, global constants, Global, else if, and better diagnostics.
• Mapped partitions support safe persistent scheduling patterns, including a
  persistent GEMM example.
• Dynamic-shape and read-only tile-like loads generate faster code in important
  cases.
• Runtime ergonomics improved with dynamic CUDA bindings, tileiras override,
  custom memory pools, and memory accounting.

The README and book have also been updated with a shorter quick-start example
and current API docs:

• Repository: https://github.com/NVlabs/cutile-rs
• Book: https://nvlabs.github.io/cutile-rs/
• Crates.io: https://crates.io/crates/cutile

Please keep sending feedback, especially on the v0.1 API surface as you start
using it from crates.io.

cuda-oxide v0.1.0

cuda-oxide v0.1.0

Initial open-source release of cuda-oxide.

cuda-oxide is an experimental rustc backend for compiling CUDA GPU kernels written in pure Rust to PTX. It supports single-source Rust GPU programs where host and device code live together and are built through cargo oxide.

This release includes:

• A custom rustc codegen backend for compiling #[kernel] functions to CUDA PTX
• Single-source host + device compilation with cargo oxide build / cargo oxide run
• Device-side Rust abstractions for indexing, shared memory, barriers, atomics, warp and cluster operations, TMA, and Blackwell-era intrinsics
• Host-side CUDA runtime crates for memory management, kernel launching, and async GPU execution
• A Rust-native compiler pipeline: Rust -> MIR -> Pliron IR -> LLVM IR -> PTX
• Example kernels covering vector add, generics, closures, atomics, clusters, async execution, MathDx interop, and GEMM

This is an early alpha release. Expect bugs, incomplete features, and API changes as the project evolves.

Quick start:

cargo install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide
cargo oxide doctor
cargo oxide run vecadd

Requirements include Linux, CUDA Toolkit 12.x+, Rust nightly with rust-src and rustc-dev, LLVM 21+ with NVPTX support, and Clang/libclang headers for host CUDA bindings.

For more details, see the project README and the cuda-oxide book: https://nvlabs.github.io/cuda-oxide/