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.

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/