Skip to content

cuda-oxide v0.3.0

Choose a tag to compare

View all tags
@nihalpasham nihalpasham released this 31 Oct 08:07
· 38 commits to main since this release
cuda-oxide/v0.3.0
99e571f
This commit was signed with the committer’s verified signature.
nihalpasham
SSH Key Fingerprint: +YJiAbnKxwPXavUrGsw05MwwqQWWSRH9ay8sc0sHJTA
Verified
Learn about vigilant mode.

cuda-oxide v0.3.0 is our largest release to date (nearly 600 commits from 45+ contributors since v0.2.1). It moves cuda-oxide onto official stable Rust (1.99.1), unifies the host runtime with cutile-rs inside the NVIDIA/cuda-rust monorepo, and ships an experimental preview of cute-rs (CUTLASS CuTe in Rust) running end-to-end Qwen3-8B (BF16 + NVFP4) and Qwen3-TTS on Blackwell GPUs.

                        NVIDIA/cuda-rust Monorepo (Stable Rust 1.99.1)
                                              │
        ┌─────────────────────────────────────┴─────────────────────────────────────┐
        ▼                                                                           ▼
Shared Host Runtime (v0.5.0)                                        Device Kernel Compilers (v0.3.0)
• cuda-core / cuda-async / cuda-bindings                            • Native SIMT (`cuda-device` -> LLVM -> PTX)
• One DeviceBuffer, Stream, Event, CudaGraph,                       • CuTe in Rust (`cute-rs-preview` -> CUTLASS 4.7)
  VMM & async DeviceOp across Tile and SIMT                         • Tile Libraries (`#[tile_fn]` SIMT-from-Tile FFI)

Executive Summary (What Landed)

Area What Changed in v0.3.0
Stable Rust 1.99.1 Dropped the daily nightly requirement; the full stack runs on stable Rust.
Unified Host Runtime Single cuda-core + cuda-async runtime shared across cutile-rs and SIMT.
cute-rs-preview CuTe layout algebra, TMA pipelines, and Blackwell MMAs in Rust's type system.
End-to-End Models Full 36-layer Qwen3-8B (BF16 & NVFP4) and Qwen3-TTS 1.7B running in Rust.
Blackwell SOTA Codegen 32-bit addrspace(3) SMEM pointers, 7-stage tcgen05 GEMM, value-SSA wgmma.
Hardware Coverage 1,240+ catalog-verified device intrinsics (sm_75 through sm_100a/sm_120a).

1. Official Stable Rust (1.99.1)

cuda-oxide no longer requires pinning a rolling nightly Rust compiler:

  • By pairing stable rustup's rustc-dev component (librustc_driver.so + rustc_public) with cargo-oxide, the entire monorepo (shared host crates, cutile-rs, and cuda-oxide device kernels) builds on official stable Rust 1.99.1.
  • Eliminates toolchain churn for production teams adopting Rust on GPUs.

2. Unified NVIDIA/cuda-rust Monorepo & Shared Host Runtime

cuda-oxide (SIMT + CuTe) and cutile-rs (Tile IR) now live together in NVIDIA/cuda-rust on top of a single host runtime (cuda-bindings, cuda-core, cuda-core-derive, cuda-async v0.5.0):

Before: Separate SIMT and Tile host wrappers (`cuda_async::simt` vs. Tile runtime)
After:  One `DeviceBuffer<T>`, `Stream`, `Event`, `CudaGraph`, and `DeviceOp` across both
  • Zero-copy interoperability: Pass the same DeviceBuffer, stream, or CUDA Graph node seamlessly between cutile-rs tile kernels and cuda-oxide SIMT/CuTe kernels.
  • Tile Libraries (#[tile_fn]): Export Rust SIMT device functions as Tile IR libraries (oxide-tile-export + oxide-tilebc) callable directly from Tile kernels.
  • Advanced host APIs: Virtual Memory Management (VMM), multi-GPU peer access, NVLink Switch multicast (multicast), stream priorities, and non-blocking stream completion queries.

3. cute-rs-preview: CuTe in Rust + End-to-End Blackwell LLMs

v0.3.0 introduces cute-rs-preview (cute-rs, ctm-rs, cute-layout, dialect-cute), bringing CUTLASS CuTe 4.7 to Rust for Blackwell (sm_100a B200 and sm_120a RTX 5090) with zero Python runtime:

Rust Types (`cute-rs` + `ctm-rs`)  ──►  `dialect-cute` (Verified IR)  ──►  `cutlass-mlir` (Embedded Cubin)
  • Compile-time layout & pipeline verification: Layouts (RowMajor, Swizzle, Sm1xxBlockScaleKMajor), TMA descriptors (TensorMap), producer/consumer pipelines (PipelineTmaAsync, PipelineTmaUmma, PipelineUmmaAsync), Tensor Memory (TmemAllocator), and Tensor Core MMAs (Sm100TiledMma, Sm100BlockScaledTiledMma, Mxf4TiledMma) are encoded and checked in Rust's type system.
  • 14 standalone kernel examples: Covering elementwise, RoPE, RMSNorm, SwiGLU, gated FFN, FP16/BF16 GEMM, block-scaled NVFP4 GEMM/GEMV, SM120a FMHA prefill/decode, and SM100a FMHA prefill & MLA decode.
  • 2 end-to-end generative models:
    • qwen3_8b_cute: Full 36-layer Qwen/Qwen3-8B inference in BF16 and hardware-scaled NVFP4 (3.56x smaller memory footprint).
    • qwen_decoder_cute: End-to-end Qwen3-TTS 1.7B neural audio codec and speech decoder.

4. Blackwell & Hopper Performance & Compiler Optimizations

Major codegen improvements in mir-transforms, mir-lower, and llvm-export close the gap with hand-tuned CUDA C++:

  • 32-bit Shared Memory Addressing (addrspace(3)): Lowering shared-memory pointers as 32-bit addrspace(3) (p3:32:32:32 with --nvptx-short-ptr) folds constant byte offsets directly into ld.shared, st.shared, ldmatrix, and tcgen05 instructions with zero 64-bit address overhead.
  • 7-Stage Blackwell tcgen05 & Value-Threaded Hopper wgmma:
    • gemm_sol_final adds a 7-stage TMA + tcgen05 TMEM pipeline and wave-tail tile balancing on B200 (sm_100a).
    • wgmma accumulators are threaded in SSA registers across counted K-loops (2-slot and 3-slot pipelines) with zero local-memory spills.
  • First-Class Loop Unroll Controls (mir-transforms): Added #[unroll], #[unroll(N)], and #[unroll(never)] / #[nounroll] attributes backed by loop induction analysis and LLVM loop metadata propagation.
  • Native PTX Dialect & Scheduler (dialect-ptx, ptx-parse, ptx-schedule): Parses PTX into a verified CFG dialect for instruction-scheduling perturbation and optimization campaigns.

5. 1,240+ Verified Hardware Intrinsics & Complete Device Math

  • Catalog-generated intrinsics (cuda-intrinsics-gen): Over 1,240 device intrinsics verified against hardware PTX ISA floors (sm_75 through sm_100a/sm_120a), including:
    • Blackwell tcgen05 (narrow aliases, typed FP8/FP6/FP4 descriptors, TMEM alloc/ld/st/wait),
    • Block-scaled mma.sync.aligned.block_scale (MXFP8, MXFP6, NVFP4) and Ada sm_89 sparse FP8 MMA,
    • TMA CTA-local tile loads, tensor reductions, and multicast,
    • redux.sync (f32, s32, u32), DPX integer min/max chains, and packed bf16x2 / f16x2 / fp8x2 conversions.
  • 120 libm math functions: Calls to libm (sin, cos, erf, lgamma, cbrt, hypot, etc.) automatically lower to NVIDIA libdevice.
  • Cooperative Groups & Atomics: Added u64/i64 collective reductions and scans, partial-warp collectives, warp-aggregated constant FP atomic pass, and scoped AtomicPtr / f16 atomics.

6. Developer Experience, cuda-gdb Debugging & CI Rigor

  • Full cuda-gdb DWARF Debugging: --device-debug and --lineinfo preserve source locations, inlined call scopes, closures, Rust enum variants, slices, addrspace(3) shared statics, and addrspace(1) globals, validated against cuda-gdb in CI.
  • Register & Spill Diagnostics: Native finalization surfaces ptxas register/shared/spill counts and warns when a launch-bounded kernel spills registers.
  • In-Kernel Event Tracing (iket): Semantic trace instrumentation for GPU kernels (dialect-iket + iket-lower).
  • Expanded cargo oxide CLI: Added inspect, list, clean, update, toolchain, and enhanced doctor (validates cuda-oxide.toml, pinned toolchain, and compute-sanitizer).
  • 241 GPU Smoke Tests + 16 Automated Repo Guards: Every example crate, compile-fail diagnostic, intrinsic ABI, license header, and book code snippet is continuously verified in CI.
Assets 2
Loading