From e9acc331c7d354c6e13c849eb58e34feaafb3659 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 23 Oct 2026 12:04:42 +0530 Subject: [PATCH] chore(cuda-oxide): update flake.lock for 1.99.1, scaffold config, and clean up docs - Update `rust-overlay` in `cuda-oxide/flake.lock` so `nix develop` resolves stable Rust `1.99.1`, export `RUSTC_BOOTSTRAP=1` in `flake.nix`, and use a deterministic per-user driver symlink directory instead of leaking `/tmp` directories under `nix develop --command`. - Have `cargo oxide new` scaffold `.cargo/config.toml` with `RUSTC_BOOTSTRAP = "1"` so plain `cargo check` and `rust-analyzer` work out of the box on stable `1.99.1`. - Verify in `check-toolchain-parity.sh` that the `oxalica/rust-overlay` revision locked in `cuda-oxide/flake.lock` serves the manifest for the pinned toolchain channel. - Remove the unused `Install LLVM 21` step from `.github/workflows/unit-tests.yml` (`cuda-oxide-codegen` uses the Rust sysroot's LLVM 23 `llvm-tools` component). - Update remaining nightly wording in `CONTRIBUTING.md`, `AGENTS.md`, `cuda-oxide/README.md`, `getting-started/installation.md`, `getting-started/hello-gpu.md`, and `.github/workflows/status-guard.yml`. - Remove backtick styling from book sidebar titles, section headings, and prose crate/tool names across the `cute/` and `compiler/rustc-public.md` pages, and fix right-edge bar label clipping in `03-rtx5090-benchmark-comparison.svg`. Signed-off-by: nihalpasham --- .github/workflows/status-guard.yml | 17 +++-- .github/workflows/unit-tests.yml | 15 ---- AGENTS.md | 12 ++-- CONTRIBUTING.md | 17 +++-- cuda-oxide/README.md | 2 +- .../cargo-oxide/src/commands/scaffold.rs | 13 ++++ .../crates/cargo-oxide/src/commands/tests.rs | 2 + .../cuda-oxide-book/compiler/rustc-public.md | 21 +++--- cuda-oxide/cuda-oxide-book/cute/contracts.md | 6 +- .../cute/cutedsl-comparison.md | 28 ++++---- .../cute/gemm-and-pipelines.md | 12 ++-- .../cute/layouts-and-tiling.md | 10 +-- cuda-oxide/cuda-oxide-book/cute/overview.md | 24 +++---- .../cuda-oxide-book/cute/support-matrix.md | 2 +- .../getting-started/hello-gpu.md | 4 +- .../getting-started/installation.md | 6 +- cuda-oxide/cuda-oxide-book/index.md | 2 +- .../examples/blockscale_gemm_ctm/README.md | 2 +- .../models/qwen3_8b_cute/README.md | 14 ++-- .../svg/03-rtx5090-benchmark-comparison.svg | 40 +++++------ cuda-oxide/flake.lock | 6 +- cuda-oxide/flake.nix | 17 +++-- cuda-oxide/scripts/check-toolchain-parity.sh | 72 +++++++++++++++++-- 23 files changed, 209 insertions(+), 135 deletions(-) diff --git a/.github/workflows/status-guard.yml b/.github/workflows/status-guard.yml index 1bb635e1df..56deee2e06 100644 --- a/.github/workflows/status-guard.yml +++ b/.github/workflows/status-guard.yml @@ -60,13 +60,12 @@ jobs: - name: Verify README lists every workspace member run: bash cuda-oxide/scripts/check-crate-inventory.sh - # The pin is copied into places no build compares: the nested backend pin - # (rustup reads it for the backend's own [workspace], so a mismatch there - # is a load failure rather than a warning), the scaffold `cargo oxide new` - # writes into user projects, the devcontainer's preinstalled toolchain, the - # book's quoted [toolchain] blocks, and the dated rustup/cargo commands in - # the book and READMEs. #727 added rustfmt to both real files and left both - # book quotes behind. + # The pin is copied into places no build compares: the scaffold `cargo oxide + # new` writes into user projects, the devcontainer's preinstalled toolchain, + # the book's quoted [toolchain] blocks, the explicit rustup/cargo commands in + # the book and READMEs, and the `rust-overlay` input in `cuda-oxide/flake.lock`. + # #727 added rustfmt to both real files and left both book quotes behind; + # #1609 bumped the pin to 1.99.1 and left `flake.lock` behind. toolchain-pin-parity: name: toolchain pin parity runs-on: ubuntu-latest @@ -130,9 +129,9 @@ jobs: - name: Clone repository uses: actions/checkout@v4 -# rustfmt comes from cuda-oxide/rust-toolchain.toml's component list through the + # rustfmt comes from cuda-oxide/rust-toolchain.toml's component list through the # rustup proxy, as in fmt.yml, so the order checked is the pinned - # nightly's. No cargo, no CUDA toolkit, no network beyond the toolchain. + # toolchain's. No cargo, no CUDA toolkit, no network beyond the toolchain. - name: Verify the book's Rust blocks use rustfmt's import order run: bash cuda-oxide/scripts/check-book-import-order.sh diff --git a/.github/workflows/unit-tests.yml b/.github/workflows/unit-tests.yml index f873394207..312b598531 100644 --- a/.github/workflows/unit-tests.yml +++ b/.github/workflows/unit-tests.yml @@ -122,21 +122,6 @@ jobs: sub-packages: '["nvcc", "cudart-dev"]' non-cuda-sub-packages: '["libcurand-dev"]' - - name: Install LLVM 21 (llc/opt for the NVPTX backend) - if: matrix.package == 'cuda-oxide-codegen' - run: | - set -euo pipefail - source /etc/os-release - sudo install -d -m 0755 /etc/apt/keyrings - wget -qO- https://apt.llvm.org/llvm-snapshot.gpg.key \ - | sudo tee /etc/apt/keyrings/apt-llvm-org.asc >/dev/null - echo "deb [signed-by=/etc/apt/keyrings/apt-llvm-org.asc] https://apt.llvm.org/${VERSION_CODENAME}/ llvm-toolchain-${VERSION_CODENAME}-21 main" \ - | sudo tee /etc/apt/sources.list.d/llvm-21.list - sudo apt-get update - sudo apt-get install -y --no-install-recommends llvm-21 - llc-21 --version | head -2 - opt-21 --version | head -2 - - name: Run crate unit tests env: CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }} diff --git a/AGENTS.md b/AGENTS.md index a9836f863c..85a13cdb67 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -31,8 +31,9 @@ contribution guide. There is no root `Justfile`. The SIMT recipes live in - `cuda-bindings/`, `cuda-core/`, `cuda-core-derive/`, and `cuda-async/` — shared host crates in the root workspace. Both kernel stacks depend on these by path. - `cuda-oxide/` — the SIMT model. Its `Cargo.toml` explicitly lists workspace - members under `crates/`; its `rust-toolchain.toml` pins the nightly needed for - `rustc_private`. The book is `cuda-oxide/cuda-oxide-book/`, and + members under `crates/`; its `rust-toolchain.toml` pins the stable release + (`1.99.1`) and `rustc-dev` / `llvm-tools` components needed for `rustc_private` + (unlocked via `RUSTC_BOOTSTRAP=1`). The book is `cuda-oxide/cuda-oxide-book/`, and `cuda-oxide/intrinsics/` holds the intrinsic catalog and ABI ledger. `cuda-oxide/crates/rustc-codegen-cuda/` declares its own workspace and inherits the SIMT toolchain pin. Examples and test fixtures can also declare separate @@ -84,15 +85,16 @@ nested workspaces. ## Toolchains The root `rust-toolchain.toml` and `cutile-rs/rust-toolchain.toml` carry matching -stable pins. `cuda-oxide/rust-toolchain.toml` supplies the SIMT nightly, including -for its nested backend workspace. Do not float a pin to work around a failure. +stable pins. `cuda-oxide/rust-toolchain.toml` supplies the SIMT stable pin and +compiler components (`rustc-dev`, `llvm-tools`), including for its nested backend +workspace. Do not float a pin to work around a failure. `cuda-oxide/scripts/check-toolchain-parity.sh` checks the SIMT pin against its scaffold, development container, and documentation. rustup normally resolves the toolchain from the working directory; an explicit override can change that. `--manifest-path` does not change the working directory, so pointing a root Cargo command at the SIMT manifest does not select the SIMT -nightly. Change directory first or select the required toolchain explicitly. +toolchain pin. Change directory first or select the required toolchain explicitly. Nix environments live in `cuda-oxide/flake.nix` and `cutile-rs/flake.nix`. Run `nix develop` from the relevant component. diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 5109e9308c..93c51116b7 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -142,13 +142,16 @@ Each product has its own workspace and toolchain under `cuda-oxide/` or `cutile-rs/`. The setup commands below cover cuda-oxide; the Tile checks have their own section further down. -cuda-oxide requires the Rust nightly toolchain with `rustc_private` support. -See the [cuda-oxide README](cuda-oxide/README.md) for setup instructions. - -`cuda-oxide/flake.nix` provides the CUDA 13, LLVM 22, Clang, and pinned nightly -development environment. From the repository root, run `cd cuda-oxide` and -then `nix develop` with flakes enabled. Container setup is documented in -[cuda-oxide/.devcontainer/README.md](cuda-oxide/.devcontainer/README.md). +cuda-oxide pins its stable Rust toolchain (`1.99.1`) and required `rustc-dev` +and `llvm-tools` components in `cuda-oxide/rust-toolchain.toml` (with +`RUSTC_BOOTSTRAP=1` configured via `cuda-oxide/.cargo/config.toml` and +`cargo-oxide`). See the [cuda-oxide README](cuda-oxide/README.md) for setup +instructions. + +`cuda-oxide/flake.nix` provides the CUDA 13, Clang/libclang, and pinned Rust +(with LLVM 23 `llvm-tools`) development environment. From the repository root, +run `cd cuda-oxide` and then `nix develop` with flakes enabled. Container setup +is documented in [cuda-oxide/.devcontainer/README.md](cuda-oxide/.devcontainer/README.md). ### Running the cuda-oxide checks diff --git a/cuda-oxide/README.md b/cuda-oxide/README.md index 79f8349b6b..f8dddf0e9a 100644 --- a/cuda-oxide/README.md +++ b/cuda-oxide/README.md @@ -150,7 +150,7 @@ cargo oxide update ### Requirements - **cargo-oxide** — cargo subcommand that drives the build pipeline (`cargo oxide run`, `build`, `sanitize`, `debug`, etc.) -- **Rust nightly** with `rust-src` and `rustc-dev` and `llvm-tools` components (pinned in `rust-toolchain.toml`) +- **Rust 1.99.1 (stable)** with `rust-src`, `rustc-dev`, and `llvm-tools` components (pinned in `rust-toolchain.toml`) - **CUDA Toolkit** (13.0+, including the cuRAND headers; `libcurand-dev` on Ubuntu). The shared `cuda-bindings` crate loads `libcuda` at run time and needs a CUDA 13.x driver (R580+) - **Clang + libclang dev headers** (`clang-21` / `libclang-common-21-dev`) — needed by `bindgen` when building the host `cuda-bindings` crate - **Linux** (tested on Ubuntu 24.04) diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs b/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs index 3a317d78cf..8e338985ad 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs @@ -18,6 +18,10 @@ channel = "1.99.1" components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"] "#; +const CARGO_CONFIG_TOML: &str = r#"[env] +RUSTC_BOOTSTRAP = "1" +"#; + const SCAFFOLD_GITIGNORE_EXTRA: &[&str] = &[ "/target/", "**/*.bc", // bitcode leftovers not in the clean suffix list @@ -52,6 +56,7 @@ pub(super) fn scaffold_gitignore() -> String { #[derive(Clone, Debug, Eq, PartialEq)] pub(super) struct ScaffoldFiles { pub(super) cargo_toml: String, + pub(super) cargo_config_toml: String, rust_toolchain_toml: String, pub(super) gitignore: String, pub(super) readme: String, @@ -299,6 +304,7 @@ fn main() -> Result<(), Box> { pub(super) fn scaffold_files(name: &str, async_mode: bool) -> ScaffoldFiles { ScaffoldFiles { cargo_toml: scaffold_cargo_toml(name, async_mode), + cargo_config_toml: CARGO_CONFIG_TOML.to_string(), rust_toolchain_toml: RUST_TOOLCHAIN_TOML.to_string(), gitignore: scaffold_gitignore(), readme: scaffold_readme(name, async_mode), @@ -319,10 +325,17 @@ pub fn scaffold_new(name: &str, async_mode: bool) { eprintln!("Error creating directory: {}", e); std::process::exit(1); }); + let cargo_dir = project_dir.join(".cargo"); + std::fs::create_dir_all(&cargo_dir).unwrap_or_else(|e| { + eprintln!("Error creating directory: {}", e); + std::process::exit(1); + }); let files = scaffold_files(name, async_mode); std::fs::write(project_dir.join("Cargo.toml"), files.cargo_toml) .expect("Failed to write Cargo.toml"); + std::fs::write(cargo_dir.join("config.toml"), files.cargo_config_toml) + .expect("Failed to write .cargo/config.toml"); std::fs::write( project_dir.join("rust-toolchain.toml"), files.rust_toolchain_toml, diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs b/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs index 2ddcaeb911..59647d3a78 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs @@ -5741,6 +5741,8 @@ fn scaffold_names_the_canonical_repository() { .contains("git = \"https://github.com/NVIDIA/cuda-rust.git\"") ); assert!(!files.cargo_toml.contains("NVlabs")); + assert!(files.cargo_config_toml.contains("[env]")); + assert!(files.cargo_config_toml.contains("RUSTC_BOOTSTRAP = \"1\"")); } } diff --git a/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md b/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md index d046635d95..d42ac59af3 100644 --- a/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md +++ b/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md @@ -282,7 +282,7 @@ already explicit. The `mir-importer` crate translates this into Pliron IR Even with `rustc_public` providing a stable API, the bridge layer (`rustc_public_bridge`) is compiled against `librustc_driver.so` and is tied to an exact compiler build. Instead of chasing daily `nightly` snapshots, -`cuda-oxide` pins an official **stable** Rust release via +cuda-oxide pins an official **stable** Rust release via `rust-toolchain.toml`: ```toml @@ -291,22 +291,27 @@ channel = "1.99.1" components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"] ``` +When bumping this pin, run `cuda-oxide/scripts/check-toolchain-parity.sh` to +verify that `scaffold.rs`, `.devcontainer/devcontainer.json`, the book, and the +READMEs agree, and run `nix flake update rust-overlay` in `cuda-oxide/` so +`cuda-oxide/flake.lock` carries the new release manifest. + ```{figure} images/ecosystem-and-bootstrap.svg :align: center :width: 100% -How custom Rust compiler backends ship in production, and what lives inside an official stable `rustup` release. +How custom Rust compiler backends ship in production, and what lives inside an official stable rustup release. ``` -### Why stable `rustc` already has `rustc_private` and `rustc_public` +### Why stable rustc already has rustc_private and rustc_public -Official stable `rustc` binaries are compiled from the same source tree as -nightly; nothing is stripped out of the binary. In addition, `rustup` publishes +Official stable rustc binaries are compiled from the same source tree as +nightly; nothing is stripped out of the binary. In addition, rustup publishes the `rustc-dev` component (`librustc_driver.so`, `librustc_public.rlib`, `librustc_public_bridge.rlib`, and `librustc_middle.rlib`) for every stable release. -When `cargo-oxide` invokes `cargo` (and inside the workspace via +When cargo-oxide invokes Cargo (and inside the workspace via `cuda-oxide/.cargo/config.toml`), it sets `RUSTC_BOOTSTRAP=1` in the process environment. That unlocks `#![feature(rustc_private)]` and `-Z` flags (`-Z codegen-backend=...`, `-Z mir-enable-passes=-JumpThreading`, @@ -318,10 +323,10 @@ stabilization cycle. :align: center :width: 100% -Pinning an official stable release (`1.99.1`) in `cuda-oxide/rust-toolchain.toml` with `RUSTC_BOOTSTRAP=1` managed by `cargo-oxide`. +Pinning an official stable release (`1.99.1`) in `cuda-oxide/rust-toolchain.toml` with `RUSTC_BOOTSTRAP=1` managed by cargo-oxide. ``` -### Upstreaming to `rustc_public` without emergency toolchain bumps +### Upstreaming to rustc_public without emergency toolchain bumps When a feature requires a new `rustc_public` query that we upstream to `rust-lang/rust` (such as `ArgAttributes` getters in diff --git a/cuda-oxide/cuda-oxide-book/cute/contracts.md b/cuda-oxide/cuda-oxide-book/cute/contracts.md index 55e2184702..45e07fe5d3 100644 --- a/cuda-oxide/cuda-oxide-book/cute/contracts.md +++ b/cuda-oxide/cuda-oxide-book/cute/contracts.md @@ -8,7 +8,7 @@ with the author at `unsafe` boundaries. :align: center :width: 100% -Where `ctm-rs` sits: typed carriers select instruction families, one of two compiler profiles consumes them, and the author keeps the schedule. +Where ctm-rs sits: typed carriers select instruction families, one of two compiler profiles consumes them, and the author keeps the schedule. ``` ## What is checked @@ -16,8 +16,8 @@ Where `ctm-rs` sits: typed carriers select instruction families, one of two comp | Layer | Checks | | :--- | :--- | | Rust types | layout, swizzle, operand role, element format, MMA shape, pipeline stage count, barrier ids; `TensorMut` over `DisjointSlice`; `TmemGuard` typestate; sealed traits | -| `cute-layout` proofs | cooperative copy coverage and per-thread contiguity, `ldmatrix` source validity, TMA encodability and swizzle phase alignment | -| `dialect-cute` verifier | tensor provenance, one pipeline init with complete producer and consumer state, cursor provenance, scheduler grids, MMA placement, epilogue phases, no compiler-only type reaches runtime storage | +| cute-layout proofs | cooperative copy coverage and per-thread contiguity, `ldmatrix` source validity, TMA encodability and swizzle phase alignment | +| dialect-cute verifier | tensor provenance, one pipeline init with complete producer and consumer state, cursor provenance, scheduler grids, MMA placement, epilogue phases, no compiler-only type reaches runtime storage | | Target gate | `sm_100a` or `sm_120a`; `tcgen05` only on `sm_100a`; MXF4 register MMA only on `sm_120a` | | Launch ABI | cubin argument sizes and offsets match the exported MLIR; `#[grid_constant] &TensorMap` is a 128-byte by-value argument aligned to 64 | diff --git a/cuda-oxide/cuda-oxide-book/cute/cutedsl-comparison.md b/cuda-oxide/cuda-oxide-book/cute/cutedsl-comparison.md index b8b7595a45..091ce1977c 100644 --- a/cuda-oxide/cuda-oxide-book/cute/cutedsl-comparison.md +++ b/cuda-oxide/cuda-oxide-book/cute/cutedsl-comparison.md @@ -1,6 +1,6 @@ -# From Python `CuTeDSL` to `cute-rs` +# From Python CuTeDSL to cute-rs -If you have used NVIDIA's Python **`CuTeDSL`**, the concepts in `cute-rs` will feel immediately familiar: layouts, composed swizzles, `zipped_divide`, TMA copy atoms, and tiled MMA. Both compile down to the **exact same CUTLASS `cute` and `cute_nvgpu` MLIR operations**. +If you have used NVIDIA's Python **CuTeDSL**, the concepts in cute-rs will feel immediately familiar: layouts, composed swizzles, `zipped_divide`, TMA copy atoms, and tiled MMA. Both compile down to the **exact same CUTLASS `cute` and `cute_nvgpu` MLIR operations**. The key difference is **where the static configuration lives** and **when the IR is constructed**: @@ -8,14 +8,14 @@ The key difference is **where the static configuration lives** and **when the IR :align: center :width: 100% -Python `CuTeDSL` traces Python objects wrapping live MLIR values at runtime; Rust `cute-rs` encodes layouts, swizzles, and MMA plans in Rust types and translates them ahead of time. +Python CuTeDSL traces Python objects wrapping live MLIR values at runtime; Rust cute-rs encodes layouts, swizzles, and MMA plans in Rust types and translates them ahead of time. ``` --- ## Why the Kernel Signatures Look Different -In Python **`CuTeDSL`**, `@cute.jit` traces a Python function at runtime to build an MLIR module. Because Python objects like `cute.Tensor`, `cute.CopyAtom`, `cute.TiledMma`, and `cute.ComposedLayout` directly wrap MLIR values, the host script constructs all of them and passes them as parameters into `@cute.kernel`: +In Python **CuTeDSL**, `@cute.jit` traces a Python function at runtime to build an MLIR module. Because Python objects like `cute.Tensor`, `cute.CopyAtom`, `cute.TiledMma`, and `cute.ComposedLayout` directly wrap MLIR values, the host script constructs all of them and passes them as parameters into `@cute.kernel`: ```python # Python CuTeDSL: 18 parameters wrapping MLIR types during @cute.jit tracing @@ -34,7 +34,7 @@ def kernel( ): ... ``` -In Rust **`cute-rs`**, there is no runtime tracer: `rustc` compiles the kernel ahead of time. Nothing is hidden: +In Rust **cute-rs**, there is no runtime tracer: rustc compiles the kernel ahead of time. Nothing is hidden: 1. **Kernel parameters carry only the hardware launch ABI** (for example, 128-byte hardware TMA descriptors passed by value via `#[grid_constant]`, device pointers, and scalar dimensions). 2. **Layouts, swizzles, operand roles, and MMA configurations live in Rust types** (`ASmem`, `Sm100TiledMma<...>`, `MainloopPipeline`) with zero runtime byte overhead. @@ -59,9 +59,9 @@ pub fn fp16_gemm_256x352( --- -## How `mir-importer` Translates Rust Types to CuTe IR +## How mir-importer Translates Rust Types to CuTe IR -When `mir-importer` translates a `cute-rs` call, it matches the callee's canonical `rustc` `DefId`, structurally walks its monomorphized generic type arguments (`static_config.rs`), runs `cute-layout` static validation proofs, and emits a `dialect-cute` operation combining those static attributes with the call's runtime SSA operands: +When mir-importer translates a cute-rs call, it matches the callee's canonical rustc `DefId`, structurally walks its monomorphized generic type arguments (`static_config.rs`), runs cute-layout static validation proofs, and emits a dialect-cute operation combining those static attributes with the call's runtime SSA operands: ```{figure} images/rust-types-to-cute-ir.svg :align: center @@ -74,8 +74,8 @@ Compile-time Rust types supply the static IR attributes (shape, stride, swizzle, ## Side-by-Side Concept Map -| Concept | Python `CuTeDSL` | Rust `cute-rs` | -| :------ | :--------------- | :------------- | +| Concept | Python CuTeDSL | Rust cute-rs | +| :------ | :------------- | :----------- | | **Layout definition** | `cute.make_layout((128, 64), stride=(64, 1))` | `RowMajor<128, 64>` or `Layout<(C<128>, C<64>), (C<64>, C<1>)>` | | **Swizzled SMEM layout** | `cute.make_composed_layout(Swizzle(3, 3, 3), 0, layout)` | `ComposedLayout, 0, RowMajor<128, 64>>` | | **Tensor view** | `cute.make_tensor(ptr, layout)` | `Tensor::from_slice(a)`, `TensorMut::from_disjoint_slice(&mut c)`, `Tensor::from_storage(storage, layout)`: constructors are `from_*`, the Rust form of `make_*` | @@ -92,16 +92,16 @@ Compile-time Rust types supply the static IR attributes (shape, stride, swizzle, | **Tiled MMA** | `cute.make_tiled_mma(mma_op, atom_layout, ...)` | `Mxf4TiledMma::::get_slice(lane)` or `Sm100TiledMma` | | **Tiled copy** | `thr_copy = tiled_copy.get_slice(tidx)`; `cute.copy(tiled_copy, thr_copy.partition_S(gA), thr_copy.partition_D(sA))` | `tiled_copy.get_slice(thread_idx).copy(&g_a, tile_coord, &mut s_a)`: both partitions come from the plan's layouts inside `copy_g2s`. No kernel can reach this surface yet; see the next section | | **Register vectors** | `TensorSSA` with `a + b`, `a - b`, `a * b`, `a / b`, `a // b`, `-a`, `where`, conversions and reductions | `RegisterTile` with `+`, `-`, `*`, `/` and unary `-`; `map`, `zip_map`, `fold` and `gather` for the rest | -| **Architecture primitives** | `cute.arch.thread_idx()`, `cute.arch.warp_idx()`, `cute.arch.elect_one()`, `cute.arch.cluster_arrive()`, `cute.arch.cluster_dim()`, the primitives-layer `fence_sync_restrict(MemOrder.ACQUIRE)`, ... | `cute_rs::arch::thread_idx()`, `warp_idx()`, `elect_one_sync()`, `cluster_arrive()`, `cluster_dim()`, `fence_sync_restrict::()`, ...: forwarders over the PTX-named `ctm_rs` layer, the mask passed as an argument | +| **Architecture primitives** | `cute.arch.thread_idx()`, `cute.arch.warp_idx()`, `cute.arch.elect_one()`, `cute.arch.cluster_arrive()`, `cute.arch.cluster_dim()`, the primitives-layer `fence_sync_restrict(MemOrder.ACQUIRE)`, ... | `cute_rs::arch::thread_idx()`, `warp_idx()`, `elect_one_sync()`, `cluster_arrive()`, `cluster_dim()`, `fence_sync_restrict::()`, ...: forwarders over the PTX-named ctm-rs layer, the mask passed as an argument | --- -## Where `cute-rs` Deviates, and Why +## Where cute-rs Deviates, and Why The names follow CuTeDSL 4.7. Where the Rust form differs on purpose, the difference and its reason are listed here, so a CuTeDSL author is not surprised. -| Concept | CuTeDSL 4.7 | Rust `cute-rs` | Why | -| :------ | :---------- | :------------- | :-- | +| Concept | CuTeDSL 4.7 | Rust cute-rs | Why | +| :------ | :---------- | :----------- | :-- | | **Constructors and plans** | `make_layout`, `make_tensor`, `make_tiled_mma` and `make_tiled_copy` build objects while tracing | `new` and `from_*` constructors; a compile-time plan is a type or a const generic (`RowMajor<128, 64>`, `TiledCopy`, `zipped_divide::<8>()`) | the compiler reads the plan from the generic arguments, so there is nothing for a `make_*` call to build | | **Pipeline producer step** | `producer_acquire` folds the arrive-expect into the wait, under `is_leader_cta` in `PipelineTmaUmma` | `producer_acquire` waits; the elected lane calls `producer_expect_tx`, or `producer_expect_tx_cursor` on `TwoCtaRing` | electing the lane is a convergence obligation the compiler does not prove, so it stays visible in kernel code; the pipeline value stores pointers only | | **Scheduler** | `WorkTileInfo.is_valid_tile`; `get_grid_shape(params, max_active_clusters)` returns `(m, n, persistent_clusters)`; `advance_to_next_work(advance_count=1)`; `tile_idx` is a `Coord` | `has_work()` on the scheduler; `get_grid_shape(sm_count, active_ctas_per_sm) -> Option`; `advance_to_next_work()`; `tile_idx() -> (m, n, l)` with `linear_idx()` beside it | the grid is one dimensional and has no params object; the bounds check stays outside the coordinate arithmetic; the step is always one grid of blocks; the coordinate is a Rust tuple | @@ -113,4 +113,4 @@ The names follow CuTeDSL 4.7. Where the Rust form differs on purpose, the differ | **Tiled copy** | `make_tiled_copy`, `get_slice`, `partition_S`, `partition_D`, `cute.copy` | `TiledCopy`, `get_slice` and `ThrCopy::copy` exist, but no kernel can reach them yet | `GlobalCopyTensor` has no public constructor: `from_carrier` is `pub(crate)` and the only public producers are the block-scaled tensors' `values_for_copy` and `scale_words_for_copy`; the importer's `shared_tensor_overlay` hook accepts only the block-scaled roles `Mkl` and `Nkl`; the verifier rejects a role-typed shared carrier as a `copy_g2s` destination. The kernels that need a block-wide copy call `copy_g2s` directly | | **Register vectors** | `/` on integer vectors converts both sides to f32 and divides; `//` floors; `-a` is `0 - a` | `/` on integer lanes truncates toward zero; unary `-` negates each lane, so a `+0.0` lane becomes `-0.0` | the operators are Rust's own: floating-point division matches the DSL, integer division and the sign of a negated zero differ | -`ctm-rs` keeps PTX instruction names (`cp_async_bulk_tensor_shared_cta_global`, `mbarrier_arrive_expect_tx`, `fence_sync_restrict_acquire`); the `cute.arch` concept names live in `cute_rs::arch` as forwarders over them. Dialect op, type and attribute names are compiler-internal and are not part of this surface. +ctm-rs keeps PTX instruction names (`cp_async_bulk_tensor_shared_cta_global`, `mbarrier_arrive_expect_tx`, `fence_sync_restrict_acquire`); the `cute.arch` concept names live in `cute_rs::arch` as forwarders over them. Dialect op, type and attribute names are compiler-internal and are not part of this surface. diff --git a/cuda-oxide/cuda-oxide-book/cute/gemm-and-pipelines.md b/cuda-oxide/cuda-oxide-book/cute/gemm-and-pipelines.md index 276d80e147..f2753b2588 100644 --- a/cuda-oxide/cuda-oxide-book/cute/gemm-and-pipelines.md +++ b/cuda-oxide/cuda-oxide-book/cute/gemm-and-pipelines.md @@ -1,10 +1,10 @@ -# TMA Pipelines, Tensor Cores, and `ctm-rs` Interop +# TMA Pipelines, Tensor Cores, and ctm-rs Interop -The Blackwell kernels in this preview combine multi-stage TMA pipelines, swizzled shared-memory views, and tensor-core MMA operations. Depending on the workload, you can express the entire schedule in **`cute-rs`** (`blockscale_gemm_cute`) or pair **`cute-rs`** with **`ctm-rs`** instruction primitives in the same kernel (`fp16_gemm_256x352_cute`, `sm100_fmha_prefill_cute`, `sm100_mla_decode_fp16_cute`). +The Blackwell kernels in this preview combine multi-stage TMA pipelines, swizzled shared-memory views, and tensor-core MMA operations. Depending on the workload, you can express the entire schedule in **cute-rs** (`blockscale_gemm_cute`) or pair **cute-rs** with **ctm-rs** instruction primitives in the same kernel (`fp16_gemm_256x352_cute`, `sm100_fmha_prefill_cute`, `sm100_mla_decode_fp16_cute`). --- -## 1. Block-Scaled MXFP4 GEMM in `cute-rs` (SM120a) +## 1. Block-Scaled MXFP4 GEMM in cute-rs (SM120a) In `blockscale_gemm_cute`, one producer warp (warp 8) fills a 3-stage TMA ring buffer in shared memory while eight compute warps (warps `0..7`) load swizzled FP4 fragments and `Float8E8M0FNU` scales into `Mxf4TiledMma`: @@ -44,17 +44,17 @@ consumer = consumer.advance(); --- -## 2. Combining `cute-rs` + `ctm-rs` on SM100a (2-CTA GEMM) +## 2. Combining cute-rs and ctm-rs on SM100a (2-CTA GEMM) On datacenter Blackwell (`sm_100a`), two CTAs cooperate on a `256×352` output tile (`256×192` + `256×160` N-partitions sharing the A collector) and accumulate into Tensor Memory (TMEM). -`fp16_gemm_256x352_cute` uses **`cute-rs`** for layout-checked cluster TMA, pipelines, paired MMA, and TMEM epilogue slices (`dialect-cute`), while using **`ctm-rs`** for cluster rank, barriers, and collective `TmemGuard` allocation (`dialect-nvvm`): +`fp16_gemm_256x352_cute` uses **cute-rs** for layout-checked cluster TMA, pipelines, paired MMA, and TMEM epilogue slices (dialect-cute), while using **ctm-rs** for cluster rank, barriers, and collective `TmemGuard` allocation (dialect-nvvm): ```{figure} images/sm100-two-cta-gemm.svg :align: center :width: 100% -In `fp16_gemm_256x352_cute`, `cute-rs` (teal) owns the 5-stage cluster TMA pipeline, paired 2-CTA MMA, and TMEM epilogue, while `ctm-rs` (purple) supplies cluster rank, barriers, and `TmemGuard` lifecycle in the same Pliron module. +In `fp16_gemm_256x352_cute`, cute-rs (teal) owns the 5-stage cluster TMA pipeline, paired 2-CTA MMA, and TMEM epilogue, while ctm-rs (purple) supplies cluster rank, barriers, and `TmemGuard` lifecycle in the same Pliron module. ``` ```rust diff --git a/cuda-oxide/cuda-oxide-book/cute/layouts-and-tiling.md b/cuda-oxide/cuda-oxide-book/cute/layouts-and-tiling.md index 43c3b27770..0b7885654d 100644 --- a/cuda-oxide/cuda-oxide-book/cute/layouts-and-tiling.md +++ b/cuda-oxide/cuda-oxide-book/cute/layouts-and-tiling.md @@ -1,12 +1,12 @@ # Layouts, Tiling, and Partitions by Example -`cute-layout` and `cute-rs` let you describe how logical coordinates map to physical memory—whether that is a 1D vector split across threads, a row-wise reduction across warps, or a 2D swizzled shared-memory tile shared by host TMA encoding and device MMA. +cute-layout and cute-rs let you describe how logical coordinates map to physical memory: whether that is a 1D vector split across threads, a row-wise reduction across warps, or a 2D swizzled shared-memory tile shared by host TMA encoding and device MMA. --- -## 1. Evaluating Layouts with `cute-layout` +## 1. Evaluating Layouts with cute-layout -`cute-layout` is a `no_std` crate shared by both `cute-rs` and the compiler. A `Layout` pairs a shape `IntTuple` with a stride `IntTuple` and maps logical coordinates to linear offsets: +cute-layout is a `no_std` crate shared by both cute-rs and the compiler. A `Layout` pairs a shape `IntTuple` with a stride `IntTuple` and maps logical coordinates to linear offsets: ```rust use cute_layout::{IntTuple, Layout}; @@ -20,7 +20,7 @@ In kernel signatures, you write static layouts as Rust type markers (`C`, Rus --- -## 2. 1D Tiling and Ragged Tails (`zipped_divide` + `slice`) +## 2. 1D Tiling and Ragged Tails (zipped_divide and slice) `zipped_divide::()` restructures a flat 1D tensor into a 2D view indexed by `(offset_inside_tile, tile_number)`. Calling `.slice(tid)` selects the tile owned by thread `tid`. `local_tile::(tile)` is the one-call form, equal to `zipped_divide::().slice(tile)`: CuTe's `local_tile` for a one-dimensional tiler. @@ -89,4 +89,4 @@ pub fn gemm_kernel(#[grid_constant] a: &TensorMap) { } ``` -If a kernel attempts to pass `&TensorMap` into a copy or MMA expecting a different layout or element type, `rustc` rejects the program at compile time before device codegen even runs. +If a kernel attempts to pass `&TensorMap` into a copy or MMA expecting a different layout or element type, rustc rejects the program at compile time before device codegen even runs. diff --git a/cuda-oxide/cuda-oxide-book/cute/overview.md b/cuda-oxide/cuda-oxide-book/cute/overview.md index 755b083b26..4fae407277 100644 --- a/cuda-oxide/cuda-oxide-book/cute/overview.md +++ b/cuda-oxide/cuda-oxide-book/cute/overview.md @@ -1,43 +1,43 @@ -# Overview: `cute-rs` and `ctm-rs` (preview) +# Overview: cute-rs and ctm-rs (preview) :::{admonition} Preview :class: warning -`cute-rs` compiles for two targets, `sm_100a` and `sm_120a`, through the pinned CUTLASS 4.7 compiler, and covers a validated subset of CuTeDSL 4.7 on them. The [support matrix](support-matrix.md) lists what is supported, what is partial and what is missing. Expect API changes between releases. +cute-rs compiles for two targets, `sm_100a` and `sm_120a`, through the pinned CUTLASS 4.7 compiler, and covers a validated subset of CuTeDSL 4.7 on them. The [support matrix](support-matrix.md) lists what is supported, what is partial and what is missing. Expect API changes between releases. ::: -**`cute-rs`** brings CUTLASS CuTe's layout algebra, partitions, TMA transfers, pipelines, and tensor-core MMA operations into Rust's type system for two Blackwell targets, without a Python interpreter or runtime JIT tracer. +**cute-rs** brings CUTLASS CuTe's layout algebra, partitions, TMA transfers, pipelines, and tensor-core MMA operations into Rust's type system for two Blackwell targets, without a Python interpreter or runtime JIT tracer. ```{figure} images/cute-rs-preview-stack-and-flow.svg :align: center :width: 100% -The `cute-rs-preview` stack: kernels author tensor computation in `cute-rs` over `ctm-rs` instruction primitives, import into one Pliron module (`dialect-cute` + `dialect-mir` + `dialect-nvvm`), pass whole-module verification, and compile through CUTLASS 4.7 MLIR into an embedded cubin. +The cute-rs-preview stack: kernels author tensor computation in cute-rs over ctm-rs instruction primitives, import into one Pliron module (dialect-cute, dialect-mir, and dialect-nvvm), pass whole-module verification, and compile through CUTLASS 4.7 MLIR into an embedded cubin. ``` --- ## The Two-Layer Kernel Stack -Every CuTe kernel in `cuda-oxide` is built on two complementary device crates plus a shared layout algebra crate: +Every CuTe kernel in cuda-oxide is built on two complementary device crates plus a shared layout algebra crate: | Crate | Layer | What it owns | | :---- | :---- | :----------- | -| **`cute-rs`** | Semantic CuTe layer | `Tensor` / `TensorMut`, `zipped_divide`, `slice`, `local_tile`, cooperative/tiled/TMA copies, producer/consumer pipelines, reductions, and tiled MMA (`mma::bf16`, `mma::block_scaled`, `sm100`, `attention::sm100`). | -| **`ctm-rs`** | Instruction & ABI layer | `#[cuda_module]`, `#[kernel]`, launch contracts, thread/warp/cluster coordinates, `mbarrier` state machines, `tcgen05::tmem::TmemGuard`, and packed register math over `cuda-device`. | -| **`cute-layout`** | Shared layout algebra | `IntTuple`, `Layout`, `ComposedLayout`, `Swizzle`, and static compile-time proofs for cooperative copies, `ldmatrix`, and TMA swizzle phase alignment. | +| **cute-rs** | Semantic CuTe layer | `Tensor` / `TensorMut`, `zipped_divide`, `slice`, `local_tile`, cooperative/tiled/TMA copies, producer/consumer pipelines, reductions, and tiled MMA (`mma::bf16`, `mma::block_scaled`, `sm100`, `attention::sm100`). | +| **ctm-rs** | Instruction & ABI layer | `#[cuda_module]`, `#[kernel]`, launch contracts, thread/warp/cluster coordinates, `mbarrier` state machines, `tcgen05::tmem::TmemGuard`, and packed register math over cuda-device. | +| **cute-layout** | Shared layout algebra | `IntTuple`, `Layout`, `ComposedLayout`, `Swizzle`, and static compile-time proofs for cooperative copies, `ldmatrix`, and TMA swizzle phase alignment. | ```{figure} images/three-tiers-comparison.svg :align: center :width: 100% -When to use `cute-rs`, `ctm-rs`, and `cuda-device`. +When to use cute-rs, ctm-rs, and cuda-device. ``` --- -## Your First `cute-rs` Kernel +## Your First cute-rs Kernel -In `cute-rs`, a `Tensor` pairs storage with a logical coordinate layout. Splitting a vector into 4-element (16-byte) tiles and slicing the tile owned by `tid` takes three lines: +In cute-rs, a `Tensor` pairs storage with a logical coordinate layout. Splitting a vector into 4-element (16-byte) tiles and slicing the tile owned by `tid` takes three lines: ```rust use ctm_rs::memory::DisjointSlice; @@ -90,5 +90,5 @@ CUDA_OXIDE_MLIR_OUTPUT="$PWD/module.mlir" \ ``` :::{note} -`dialect-cute`, `dialect-mir`, and `dialect-nvvm` coexist inside a **single Pliron module** rather than acting as separate sequential passes. Before exporting MLIR, `dialect-cute`'s whole-module verifier checks tensor origins, pipeline state transitions, and TMA/MMA layout and collector contracts. +dialect-cute, dialect-mir, and dialect-nvvm coexist inside a **single Pliron module** rather than acting as separate sequential passes. Before exporting MLIR, dialect-cute's whole-module verifier checks tensor origins, pipeline state transitions, and TMA/MMA layout and collector contracts. ::: diff --git a/cuda-oxide/cuda-oxide-book/cute/support-matrix.md b/cuda-oxide/cuda-oxide-book/cute/support-matrix.md index 086ac8512f..63c1f2ae5d 100644 --- a/cuda-oxide/cuda-oxide-book/cute/support-matrix.md +++ b/cuda-oxide/cuda-oxide-book/cute/support-matrix.md @@ -1,6 +1,6 @@ # Support Matrix (preview) -`cute-rs` is a **preview**. It compiles Rust CuTe kernels ahead of time through the pinned CUTLASS 4.7 compiler for two targets, `sm_100a` and `sm_120a`, and covers a validated subset of what CuTeDSL 4.7 offers on them. This page lists what the subset holds, what it does not, and what `cute-rs` adds. Status at commit `5eaa6cbda`, 2026-10-22. +cute-rs is a **preview**. It compiles Rust CuTe kernels ahead of time through the pinned CUTLASS 4.7 compiler for two targets, `sm_100a` and `sm_120a`, and covers a validated subset of what CuTeDSL 4.7 offers on them. This page lists what the subset holds, what it does not, and what cute-rs adds. Status at commit `5eaa6cbda`, 2026-10-22. ```{figure} images/support-matrix.svg :align: center diff --git a/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md b/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md index 61793a456e..cec0a77fba 100644 --- a/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md +++ b/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md @@ -35,10 +35,12 @@ This generates a ready-to-run project: ```text my_first_kernel/ +├── .cargo/ +│ └── config.toml # sets RUSTC_BOOTSTRAP = "1" for cargo check / rust-analyzer ├── Cargo.toml # dependencies on cuda-device, cuda-host, cuda-core ├── README.md # doctor + run instructions ├── .gitignore # target/ and generated device artifacts -├── rust-toolchain.toml # pins the required nightly toolchain +├── rust-toolchain.toml # pins the required Rust toolchain (1.99.1) └── src/ └── main.rs # kernel + host code in one file ``` diff --git a/cuda-oxide/cuda-oxide-book/getting-started/installation.md b/cuda-oxide/cuda-oxide-book/getting-started/installation.md index f8114dcc48..7557eb8e13 100644 --- a/cuda-oxide/cuda-oxide-book/getting-started/installation.md +++ b/cuda-oxide/cuda-oxide-book/getting-started/installation.md @@ -20,7 +20,7 @@ without a driver and fails on that call with `CUDA_ERROR_NOT_INITIALIZED`; the error message names the library files the loader tried. A driver whose CUDA major version is older than the toolkit's (a 12.x driver with a 13.x build) fails the same way with a "driver too old" message. -| **Rust** | Nightly (pinned) | Pinned in `cuda-oxide/rust-toolchain.toml` | +| **Rust** | 1.99.1 (pinned) | Pinned in `cuda-oxide/rust-toolchain.toml` | :::{note} cuda-oxide currently targets **Linux only**. Windows is not supported. @@ -32,7 +32,7 @@ cuda-oxide currently targets **Linux only**. Windows is not supported. The repository includes a cuda-oxide devcontainer in `cuda-oxide/.devcontainer/`. Using it is the quickest way to get a reproducible development environment with -CUDA Toolkit 13.0, LLVM 21, Clang 21, and the pinned Rust nightly already +CUDA Toolkit 13.0, LLVM 21, Clang 21, and the pinned Rust toolchain already installed. The host does not need the CUDA Toolkit installed. It does need: @@ -76,7 +76,7 @@ Rust setup sections below. ## Nix / flake.nix The repository also ships `cuda-oxide/flake.nix`, providing a reproducible dev shell -(CUDA 13, LLVM 22, Clang, pinned Rust nightly). Requires +(CUDA 13, Clang/libclang, and the pinned Rust toolchain with LLVM 23 `llvm-tools`). Requires [Nix](https://nixos.org/download/) with flakes enabled, an NVIDIA driver on the host, and Linux (x86\_64 or aarch64). diff --git a/cuda-oxide/cuda-oxide-book/index.md b/cuda-oxide/cuda-oxide-book/index.md index 92579e7d96..3922b14645 100644 --- a/cuda-oxide/cuda-oxide-book/index.md +++ b/cuda-oxide/cuda-oxide-book/index.md @@ -114,7 +114,7 @@ Schedule across stream pools. Await results with `.await`. ::: :::{grid-item-card} 🧩 Typed CuTe Kernels (preview) -Encode layouts, swizzles, TMA pipelines, and tensor-core MMA in Rust's type system with [`cute-rs`](cute/overview.md), compiled through CUTLASS 4.7 for `sm_100a` and `sm_120a`. See the [support matrix](cute/support-matrix.md). +Encode layouts, swizzles, TMA pipelines, and tensor-core MMA in Rust's type system with [cute-rs](cute/overview.md), compiled through CUTLASS 4.7 for `sm_100a` and `sm_120a`. See the [support matrix](cute/support-matrix.md). ::: :::: diff --git a/cuda-oxide/cute-rs-preview/examples/blockscale_gemm_ctm/README.md b/cuda-oxide/cute-rs-preview/examples/blockscale_gemm_ctm/README.md index 2ac27c218c..e1899793a2 100644 --- a/cuda-oxide/cute-rs-preview/examples/blockscale_gemm_ctm/README.md +++ b/cuda-oxide/cute-rs-preview/examples/blockscale_gemm_ctm/README.md @@ -1,4 +1,4 @@ -# MXFP4 GEMM with Rust CTM (`ctm-rs`) +# MXFP4 GEMM with Rust CTM (ctm-rs) This cuda-oxide example multiplies packed E2M1 matrices with UE8M0 scales per 32 K values, FP32 accumulation and FP16 output. The fixed problem is diff --git a/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/README.md b/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/README.md index b9f48f9300..6dd6c33fce 100644 --- a/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/README.md +++ b/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/README.md @@ -1,6 +1,6 @@ -# `qwen3_8b_cute` — `Qwen/Qwen3-8B` in `cute-rs` (BF16 & NVFP4 on RTX 5090) +# qwen3_8b_cute: Qwen/Qwen3-8B in cute-rs (BF16 & NVFP4 on RTX 5090) -End-to-end 36-layer **`Qwen/Qwen3-8B`** transformer decoder (`H = 4096`, `I = 12288`, `Q = 32`, `KV = 8`, `D = 128`, `RoPE θ = 1e6`, per-head `QK-Norm`) written in Rust with [`cute-rs`](../../cute-rs), supporting both **BF16 Tensor-Core MMA** and **true NVFP4 (`Float4E2M1FN` + `Float8E4M3FN` scale, `SF_VEC = 16`)** execution on Blackwell (`sm_120a`). +End-to-end 36-layer **Qwen/Qwen3-8B** transformer decoder (`H = 4096`, `I = 12288`, `Q = 32`, `KV = 8`, `D = 128`, `RoPE θ = 1e6`, per-head `QK-Norm`) written in Rust with [cute-rs](../../cute-rs), supporting both **BF16 Tensor-Core MMA** and **true NVFP4 (`Float4E2M1FN` + `Float8E4M3FN` scale, `SF_VEC = 16`)** execution on Blackwell (`sm_120a`).

Qwen/Qwen3-8B dual-precision architecture map @@ -10,7 +10,7 @@ End-to-end 36-layer **`Qwen/Qwen3-8B`** transformer decoder (`H = 4096`, `I = 12 --- -## RTX 5090 (`sm_120a`) Live Benchmark — Full 36 Layers (`cute-rs` vs. `CuTeDSL`) +## RTX 5090 (sm_120a) Live Benchmark: Full 36 Layers (cute-rs vs. CuTeDSL)

RTX 5090 BF16 vs NVFP4 benchmark comparison @@ -18,11 +18,11 @@ End-to-end 36-layer **`Qwen/Qwen3-8B`** transformer decoder (`H = 4096`, `I = 12 | Implementation & Precision Mode | 36-Layer Weight Footprint | Compression | Prefill (`32` tok) | Decode (`B = 1`) | | --- | ---: | ---: | ---: | ---: | -| **`cute-rs` NVFP4 (`Float4E2M1FN` + `Float8E4M3FN`, `SF_VEC = 16`)** | **`3,726.59 MiB` (`3.64 GiB`)** | **`3.56×`** | **`50.57 ms` (`632.8 tok/s`)** | **`8.77 ms/tok` (`114.0 tok/s`)** | +| **cute-rs NVFP4 (`Float4E2M1FN` + `Float8E4M3FN`, `SF_VEC = 16`)** | **`3,726.59 MiB` (`3.64 GiB`)** | **`3.56×`** | **`50.57 ms` (`632.8 tok/s`)** | **`8.77 ms/tok` (`114.0 tok/s`)** | | **CuTeDSL (`cutlass.cute`) NVFP4 Reference** | `3,726.59 MiB` (`3.64 GiB`) | `3.56×` | `53.08 ms` (`602.9 tok/s`) | `13.20 ms/tok` (`75.8 tok/s`) | -| **`cute-rs` BF16 (`m16n8k16` Tensor Core)** | `13,248.59 MiB` (`12.94 GiB`) | `1.00×` | `95.89 ms` (`333.7 tok/s`) | `14.32 ms/tok` (`69.8 tok/s`) | +| **cute-rs BF16 (`m16n8k16` Tensor Core)** | `13,248.59 MiB` (`12.94 GiB`) | `1.00×` | `95.89 ms` (`333.7 tok/s`) | `14.32 ms/tok` (`69.8 tok/s`) | -### Why `cute-rs` NVFP4 Wins (At a Glance) +### Why cute-rs NVFP4 Wins (At a Glance) ```text 1. 3.56× Less Memory Traffic per Token (NVFP4 vs. BF16) @@ -35,7 +35,7 @@ End-to-end 36-layer **`Qwen/Qwen3-8B`** transformer decoder (`H = 4096`, `I = 12 ``` - **NVFP4 vs. BF16 (`1.89×` prefill, `1.63×` decode):** Reading `3.64 GiB` instead of `12.94 GiB` per token cuts memory traffic by `3.56×`, and `Nvfp4Tile64::load()` lowers directly to Blackwell hardware `nvgpu.cvt_fpext`. -- **`cute-rs` (1 warp/output) vs. `CuTeDSL` (1 thread/output):** Assigning a 32-lane warp to each output row coalesces global `K`-tile loads and saturates all 170 SMs on the RTX 5090, while stream-ordered async execution over pre-allocated `Nvfp4Scratch` buffers eliminates per-layer host synchronization (**`632.8` vs. `602.9 tok/s` prefill**, **`114.0` vs. `75.8 tok/s` decode**). +- **cute-rs (1 warp/output) vs. CuTeDSL (1 thread/output):** Assigning a 32-lane warp to each output row coalesces global `K`-tile loads and saturates all 170 SMs on the RTX 5090, while stream-ordered async execution over pre-allocated `Nvfp4Scratch` buffers eliminates per-layer host synchronization (**`632.8` vs. `602.9 tok/s` prefill**, **`114.0` vs. `75.8 tok/s` decode**). --- diff --git a/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/svg/03-rtx5090-benchmark-comparison.svg b/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/svg/03-rtx5090-benchmark-comparison.svg index e3dd01d7cd..e5d1786717 100644 --- a/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/svg/03-rtx5090-benchmark-comparison.svg +++ b/cuda-oxide/cute-rs-preview/models/qwen3_8b_cute/svg/03-rtx5090-benchmark-comparison.svg @@ -26,43 +26,43 @@ - RTX 5090 (SM120A) LIVE BENCHMARK — FULL 36-LAYER QWEN/QWEN3-8B (CUTE-RS VS. CUTEDSL) + RTX 5090 (SM120A) LIVE BENCHMARK – FULL 36-LAYER QWEN/QWEN3-8B (CUTE-RS VS. CUTEDSL) - 1. 36-Layer Weight Footprint (Lower is Better — 3.56× Compression) + 1. 36-Layer Weight Footprint (Lower is Better – 3.56× Compression) cute-rs BF16 (m16n8k16) - - 13,248.6 MiB (12.94 GiB) + + 13,248.6 MiB (12.94 GiB) cute-rs / CuTeDSL NVFP4 - - 3,726.6 MiB (3.64 GiB — saves 9.30 GiB VRAM) + + 3,726.6 MiB (3.64 GiB – saves 9.30 GiB VRAM) - 2. Prefill Throughput (32-Token Prompt, 36 Layers — Higher is Better) + 2. Prefill Throughput (32-Token Prompt, 36 Layers – Higher is Better) cute-rs NVFP4 (E2M1+E4M3) - - 632.8 tok/s (50.6 ms — 1.89× vs BF16) + + 632.8 tok/s (50.6 ms – 1.89× vs BF16) CuTeDSL NVFP4 Reference - - 602.9 tok/s (53.1 ms) + + 602.9 tok/s (53.1 ms) cute-rs BF16 (m16n8k16) - - 333.7 tok/s (95.9 ms) + + 333.7 tok/s (95.9 ms) - 3. Autoregressive Decode Throughput (B=1, 36 Layers + KV Cache — Higher is Better) + 3. Autoregressive Decode Throughput (B=1, 36 Layers + KV Cache – Higher is Better) cute-rs NVFP4 (E2M1+E4M3) - - 114.0 tok/s (8.77 ms/tok — 1.63× vs BF16) + + 114.0 tok/s (8.77 ms/tok – 1.63× vs BF16) CuTeDSL NVFP4 Reference - - 75.8 tok/s (13.20 ms/tok) + + 75.8 tok/s (13.20 ms/tok) cute-rs BF16 (m16n8k16) - - 69.8 tok/s (14.32 ms/tok) + + 69.8 tok/s (14.32 ms/tok) diff --git a/cuda-oxide/flake.lock b/cuda-oxide/flake.lock index d938292504..7f363161e8 100644 --- a/cuda-oxide/flake.lock +++ b/cuda-oxide/flake.lock @@ -64,11 +64,11 @@ ] }, "locked": { - "lastModified": 1791707771, - "narHash": "sha256-w3m8HtAVCfGTHw+gcCTAy7jFHSdGbokm8NngbeFrkek=", + "lastModified": 1792653800, + "narHash": "sha256-E8zFJilWlORqzIUJizyhZ/0WKBOizuM6s1nRSgPRomo=", "owner": "oxalica", "repo": "rust-overlay", - "rev": "c1cfa7a2c809d364ac1a7e9b3ec418afa3f19793", + "rev": "5f267b7f323a4d6df0b9c0e3eaa043cbe85386f2", "type": "github" }, "original": { diff --git a/cuda-oxide/flake.nix b/cuda-oxide/flake.nix index 9b0c548c50..17eaa2e37c 100644 --- a/cuda-oxide/flake.nix +++ b/cuda-oxide/flake.nix @@ -55,7 +55,7 @@ # LLVM llvmPkgs = pkgs.llvmPackages_22; - # Nightly Rust + # Pinned Rust toolchain (from rust-toolchain.toml) rustToolchain = pkgs.rust-bin.fromRustupToolchainFile ./rust-toolchain.toml; # cuda @@ -111,6 +111,7 @@ CUDA_HOME = cudaSymlinked; LIBCLANG_PATH = "${llvmPkgs.libclang.lib}/lib"; + RUSTC_BOOTSTRAP = "1"; }; cargoOxideDeps = craneLib.buildDepsOnly ( @@ -202,6 +203,7 @@ shellHook = '' export CUDA_HOME="${cudaSymlinked}" export LIBCLANG_PATH="${llvmPkgs.libclang.lib}/lib" + export RUSTC_BOOTSTRAP="1" # GPU driver setup borrowed from https://github.com/NVlabs/cutile-rs # NixOS provides /run/opengl-driver/lib @@ -209,8 +211,12 @@ export LD_LIBRARY_PATH="/run/opengl-driver/lib:$LD_LIBRARY_PATH" else # Non-NixOS Linux (Ubuntu, Arch, etc. running Nix) - # Symlink only the NVIDIA driver to avoid host glibc pollution - _nv_drv_dir=$(mktemp -d /tmp/nix-nvidia-driver.XXXXXX) + # Symlink only the NVIDIA driver into a deterministic per-user + # directory to avoid host glibc pollution and avoid leaking a new + # /tmp directory when `nix develop --command` execs a child. + _nv_drv_dir="''${XDG_RUNTIME_DIR:-/tmp}/cuda-oxide-nvidia-driver-$(id -u)" + mkdir -p "$_nv_drv_dir" + rm -f "$_nv_drv_dir"/libcuda.so* "$_nv_drv_dir"/libnvidia-ptxjitcompiler.so* "$_nv_drv_dir"/libnvidia-gpucomp.so* for d in /usr/lib/x86_64-linux-gnu \ /lib/x86_64-linux-gnu \ @@ -228,11 +234,8 @@ if [ -n "$(ls -A "$_nv_drv_dir" 2>/dev/null)" ]; then export LD_LIBRARY_PATH="$_nv_drv_dir:$LD_LIBRARY_PATH" - # cleanup when the user exits the nix shell - trap 'rm -rf "$_nv_drv_dir"' EXIT else - # Clean up immediately if no NVIDIA drivers were found - rm -rf "$_nv_drv_dir" + rmdir "$_nv_drv_dir" 2>/dev/null || true fi fi diff --git a/cuda-oxide/scripts/check-toolchain-parity.sh b/cuda-oxide/scripts/check-toolchain-parity.sh index 099c7df9f2..9f75a81437 100755 --- a/cuda-oxide/scripts/check-toolchain-parity.sh +++ b/cuda-oxide/scripts/check-toolchain-parity.sh @@ -26,10 +26,18 @@ # symptom lands later and elsewhere: a missing `rustfmt` surfaces as # `cargo oxide fmt` failing, not as a bad toolchain file. # -# 4. The dated commands and prose across the book and the READMEs: -# `rustup toolchain install nightly-...`, `cargo +nightly-... install`, -# and sentences naming the pin. Readers run those commands outside a -# checkout, where no rust-toolchain.toml can correct a stale date. +# 4. The explicit commands and prose across the book and the READMEs: +# `rustup toolchain install `, `cargo + install`, and +# `rustup component add ... --toolchain `. Readers run those +# commands outside a checkout, where no rust-toolchain.toml can correct a +# stale version. +# +# 5. The `rust-overlay` input in `cuda-oxide/flake.lock`. `flake.nix` +# resolves `./rust-toolchain.toml` through `oxalica/rust-overlay`, so when +# bumping to a newer release or nightly, update `rust-overlay` in +# `cuda-oxide/flake.lock` (`nix flake update rust-overlay` and verify with +# `nix eval .#devShells.x86_64-linux.default.drvPath`) so `nix develop` +# has the new manifest. # # This is exactly what happened: #727 added `rustfmt` to both real files and # left both book quotes at five components. @@ -52,6 +60,7 @@ cd "$(dirname "$0")/.." ROOT_PIN=rust-toolchain.toml SCAFFOLD=crates/cargo-oxide/src/commands/scaffold.rs DEVCONTAINER=.devcontainer/devcontainer.json +FLAKE_LOCK=flake.lock if ! command -v python3 >/dev/null 2>&1; then echo "error: python3 is required to verify the toolchain pin" >&2 @@ -71,17 +80,20 @@ fi test -s "${ROOT_PIN}" test -s "${SCAFFOLD}" test -s "${DEVCONTAINER}" +test -s "${FLAKE_LOCK}" GIT_ROOT="$(git rev-parse --show-toplevel)" -python3 - "${ROOT_PIN}" "${SCAFFOLD}" "${DEVCONTAINER}" "${GIT_ROOT}" <<'PY' +python3 - "${ROOT_PIN}" "${SCAFFOLD}" "${DEVCONTAINER}" "${FLAKE_LOCK}" "${GIT_ROOT}" <<'PY' import glob import json import os import re import subprocess import sys +import urllib.error +import urllib.request -root_path, scaffold_path, devcontainer_path, git_root = sys.argv[1:5] +root_path, scaffold_path, devcontainer_path, flake_lock_path, git_root = sys.argv[1:6] CHANNEL = re.compile(r'^\s*channel\s*=\s*"([^"]+)"', re.M) # Match the whole bracketed span and pull the quoted names out of it so @@ -296,6 +308,54 @@ if not pinned_refs: "this block of the guard) or the pattern is stale (fix it)" ) +# Verify `cuda-oxide/flake.lock` locks an `oxalica/rust-overlay` revision that +# actually contains the manifest for `root_channel`. +flake_lock = json.loads(read(flake_lock_path)) +rust_overlay_locked = ( + flake_lock.get("nodes", {}).get("rust-overlay", {}).get("locked", {}) +) +overlay_owner = rust_overlay_locked.get("owner") +overlay_repo = rust_overlay_locked.get("repo") +overlay_rev = rust_overlay_locked.get("rev") +if not overlay_owner or not overlay_repo or not overlay_rev: + sys.exit( + f"parse self-test failed: expected nodes.rust-overlay.locked in " + f"{flake_lock_path}, found {rust_overlay_locked!r}" + ) +if root_channel.startswith("nightly-"): + _, year, month, day = root_channel.split("-") + manifest_rel = f"manifests/nightly/{year}/{month}-{day}.nix" +else: + manifest_rel = f"manifests/stable/{root_channel}.nix" +manifest_url = ( + f"https://raw.githubusercontent.com/{overlay_owner}/{overlay_repo}/" + f"{overlay_rev}/{manifest_rel}" +) +try: + req = urllib.request.Request(manifest_url, method="HEAD") + with urllib.request.urlopen(req, timeout=10) as resp: + if resp.status != 200: + failures.append( + f"{flake_lock_path} locks {overlay_owner}/{overlay_repo} at " + f"{overlay_rev[:10]}, where {manifest_rel} returned HTTP " + f"{resp.status} (run `nix flake update rust-overlay` in cuda-oxide/)" + ) +except urllib.error.HTTPError as err: + if err.code == 404: + failures.append( + f"{flake_lock_path} locks {overlay_owner}/{overlay_repo} at " + f"{overlay_rev[:10]}, which does not contain {manifest_rel} for " + f"{root_channel!r} (HTTP 404; run `nix flake update rust-overlay` " + "in cuda-oxide/)" + ) + else: + sys.exit(f"error checking {manifest_url}: HTTP {err.code}") +except urllib.error.URLError as err: + print( + f"note: skipping online rust-overlay manifest check ({err.reason})", + file=sys.stderr, + ) + if failures: print("error: the toolchain pin is not consistent across its copies", file=sys.stderr) for failure in failures: