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.