Repository navigation
Add the cute-rs preview: CuTe kernels in Rust for sm_100a and sm_120a - #1598
Merged
Merged
Conversation
…gnostic rendering nightly-2026-08-28 prints the mismatched 'found' type in tiled_copy_role_mismatch as `SharedTensor<_, f16, (), Nkl>` where the snapshot (recorded under the previous nightly) had `SharedTensor<..., f16, (), ...>`. Same error, same span; re-blessed with TRYBUILD=overwrite.
Same change as on main, with the three standalone cuteir example workspaces listed so rust-analyzer loads them on this branch.
Express the persistent 256x352 FP16 GEMM with typed CuTe layouts, two-CTA paired MMA and collector reuse, cluster TMA, tensor memory, and pipeline operations through the CUTLASS translation backend. Preserve vectors through epilogue conversion and shared-memory stores. Group CUTLASS compiler regressions under src/tests/cutlass, add GPU validation cases, and document the example with concise diagrams and single-launch commands. Keep the local Python reference, comparison harness, benchmark results, and performance notes out of version control. Validation: pinned CUTLASS compiler regressions, exporter/frontend checks, six GPU correctness cases, memcheck, synccheck, and single-launch Python oracle comparisons for both performance shapes passed.
Use Apache-2.0 headers for the Rust kernel and host runner and align the example package license. This change only updates license comments and metadata. Signed-off-by: nihalpasham <[email protected]>
Carry grid-constant parameter size and alignment through CUTLASS MLIR export and cubin ABI validation. Pass the SM100 GEMM descriptors directly with the launch and preserve assertion behavior after rebasing onto main. Run the local CuTe DSL comparison by default, with standalone execution available through --rust-only. Keep the local Python reference assets untracked and document the setup and timing differences. Update CuTe examples for the shared host crates. Simplify documentation, fix test compilation errors and Clippy warnings, and include the new crates in the book inventory and CI test matrix. Validation: 5,636 workspace tests, 34 doctests, strict root/backend Clippy, all-feature analyzer checks, formatting, rustdoc and book warning checks, and repository consistency guards passed. Signed-off-by: nihalpasham <[email protected]>
Add top-8 and exhaustive-grid cuBLASLt tuning with matched repeated CUDA-event timing and per-shape results. Preserve B200 measurements, validation evidence, and reproduction commands. Update GEMM benchmark documentation and include the CuTe GEMM README changes.
Merge current main while retaining the experimental CuTe translation work. Move its crates and examples under cuda-oxide, update workspace membership, host dependencies, editor projects, validation scripts, and documentation. Keep stable host features opt-in in the editor. Resolve examples from the selected backend component, and cover both example trees in formatting, cleanup, linting, and lockfile checks. Adapt CUTLASS export to validated target selection and current Pliron symbol attributes, preserving the existing compiler checks and adding regression coverage for the combined repository layout. Signed-off-by: nihalpasham <[email protected]>
Add an unsafe trapping SharedRegion constructor with the same null, alignment, capacity, and storage/lifetime contract as from_raw_parts. Use it in the SM120 blockscale example, which already trapped on errors, and test validation and pointer preservation for its three shared formats. Remove direct region-helper calls from the generated cubin without changing kernel scheduling or raw TMA semantics. Record all Nsight samples and comparison limits: median kernel time falls from 12.960 to 11.648 us across three rounds, with CuTe-RS at 11.264 us and CuTeDSL at 10.112 us. Validation: 123 cuda-device tests passed; all 4,194,304 outputs bit-exact at 170, 1, and 17 CTAs with replay; memcheck and synccheck report zero errors. The final cubin matches the GPU-validated cubin byte for byte. Signed-off-by: nihalpasham <[email protected]>
Signed-off-by: nihalpasham <[email protected]>
Implement the matching SM120a MXFP4 GEMM in CUDA C++ with typed instruction wrappers, physical-layout checks, and reproducible builds. Add a serial Nsight kernel-active comparison runner and record fresh correctness, sanitizer, and timing evidence against Rust after the main rebase. Signed-off-by: nihalpasham <[email protected]>
Move the Rust primitives GEMM into cuteir/examples as a standalone sibling of the CuTe-RS example. Vendor the measured CuTeDSL 4.6.2 source and benchmark harness under the shared reference directory, and move the C++ instruction-primitives counterpart there. Update build routing, dependency paths, IDE discovery, documentation, and license checks while preserving kernel sources and historical benchmark evidence. Signed-off-by: nihalpasham <[email protected]>
Keep cuda-device as the raw instruction layer and move typed operands, collective TMEM views, barrier adapters, and optional host tensor-map encoding into ctm-rs. Migrate examples, compile fixtures, safety tests, and CI. Forward instruction APIs explicitly to preserve compiler intrinsic identity. Verify the production API and both GEMM targets; SM120 blockscale cubin is byte-identical and all correctness, memcheck, and synccheck runs pass on GPU 1. Default/host tests, strict Clippy/rustdoc, formatting, and license checks pass. SM100 execution remains untested on this SM120 host. Signed-off-by: nihalpasham <[email protected]>
Use the CTM programming-model name for the package, executable, example directory, build checks, and documentation links. Preserve historical benchmark names and hashes. Verified the SM120 production build, byte-identical cubin, exact GPU/replay correctness for all 4,194,304 outputs, 12 comparison-runner tests, dependency inventory, and SPDX headers. Signed-off-by: nihalpasham <[email protected]>
Rename the package and executable to fp16_gemm_256x352_ctm, update relative dependencies and compiler checks, and place it beside the separate CuTe implementation. Document that the CuTe example currently uses cute-rs and cuda-device without ctm-rs. Verified production SM100 compilation, identical GPU instruction and constant bytes, CLI help, formatting, dependency inventory, and SPDX headers. SM100 execution remains untested on this SM120 host. Signed-off-by: nihalpasham <[email protected]>
Port the SM120 Rust CTM schedule to the actual CUTLASS experimental primitives API and verify identical packed fixtures and all FP16 output bits. Keep the custom CUDA C++ reference labeled separately. Record serial Nsight kernel-active results: Rust CTM 11.648 us, CUTLASS primitives 10.752 us. Both pass resident, single-CTA, phase-wrap, memcheck and synccheck validation. Document compiler 4.7/4.8 and wait-hint differences; 24 CPU checks and SPDX guard pass. Signed-off-by: nihalpasham <[email protected]>
Place the reference directly in blockscale_gemm_ctm/reference and update links and reproduction commands. Preserve kernel sources and historical measurement records byte for byte. All 24 CPU checks pass from the new location. Signed-off-by: nihalpasham <[email protected]>
Hoist private enum scratch allocations to function entry blocks so LLVM can eliminate temporary storage, and preserve inline(always) on pointer helpers. Replace FP16 kernel macros with functions and compute MMA descriptors outside elected-lane branches. Move the CuTe example's low-level calls to ctm-rs while retaining its layouts and tensor operations. Document the stack-allocation transformation and B200 validation. Paired three-run medians are 778.976 us for Rust CTM and 779.456 us for CUTLASS CTM; the controlled-clock profile retains a 2.44 percent gap. Validation: 165 affected-package tests; SM100 primitives/FP16 and SM120 blockscale compile probes; six GPU numerical cases, fractional CPU oracle, memcheck and synccheck. CuTe rebuild produces an identical cubin. Formatting and SPDX checks pass. Signed-off-by: nihalpasham <[email protected]>
Rename TensorMut<'a, T, Tile1D<TILE>, T>::store_linear(index, value) to store_at(index, value), the verb the Contiguous1D row and StridedTensorMut already use for one scalar write at an absolute index, so the stack has one name for the DSL's scalar __setitem__. The body still calls the tensor_store_element_abs boundary, whose rustdoc now links the new name. No alias or forwarder remains. The device compiler recognizes this method by name, so its constants move in the same change: the classify.rs arm that maps the Tile1D read-write receiver to CuteFn::TensorStoreElementAbs now reads "store_at", and the identity test fixture carries the new path. The collector's TensorMut method table keyed its stubs on the ADT name alone, which after the rename would also have swept the ordinary TensorMut<Contiguous1D>::store_at into the skipped-body set; the owner is now layout-aware, with a TensorMutTile1D variant that lists store and store_at while TensorMut keeps from_disjoint_slice, zipped_divide and slice, and unit tests pin both sides. The hook path, the IR operation and its verifier are unchanged, so the exported MLIR is unchanged. The four examples with a ragged tail, the two book snippets and the stale comment in the codegen prep fixture move with it. The figure that prints the old name is redrawn in the documentation batch. Signed-off-by: nihalpasham <[email protected]>
Rename cute_rs::copy::vector::assume_div::<D>(x) to assume::<D>(x), CuTeDSL's cute.assume(x, divby=D) with the divisor as a const generic. The crate root exports assume in place of assume_div, the tile alias module follows for free, and the rustdoc names the DSL call. No alias or forwarder remains. The device compiler recognizes this hook by its defining path, so its constants move in the same change: the registry's canonical path is now cute_rs::copy::vector::assume with the same CuteFn::AssumeDiv handler and const schema, the hook-path tests in the registry and the collector carry the new path, and the emitter and divisor decoder docs and diagnostics print the new name. The IR operation stays cute.assume_div, so CuteAssumeDivOp, emit_assume_div and the exported MLIR are unchanged; only the dialect-cute op doc that shows the Rust call moves. The two internal callers, the block-scaled load boundaries and the tile load and store boundaries, move with it. No example, book page or snapshot names the old function. Signed-off-by: nihalpasham <[email protected]>
…riples Add cute_rs::arch::block_in_cluster_idx() -> (u32, u32, u32), CuTeDSL's cute.arch.block_in_cluster_idx() over the ctm_rs::cluster cluster_ctaidX, cluster_ctaidY and cluster_ctaidZ readers, and cute_rs::arch::block_in_cluster_dim() -> (u32, u32, u32), cute.arch.block_in_cluster_dim() over cluster_nctaidX, cluster_nctaidY and cluster_nctaidZ. Both are inline-always forwarders in the style of thread_idx and block_dim; ctm-rs keeps its PTX names and block_idx_in_cluster, the linear rank, stays as it is. The rustdoc on block_in_cluster_dim states the two meanings of cluster_dim: the compiler's #[cluster_launch] attribute and intrinsic family mean CTAs per cluster, the value this triple reads, while CuTeDSL's cute.arch.cluster_dim() means the number of clusters in the grid. Nothing in the compiler is renamed. The cute.arch.cluster_dim() and cute.arch.cluster_idx() triples over %nclusterid and %clusterid are not added: cuda-device holds those per-axis readers only as generated pub(crate) items, and no hand-written module or ctm-rs forwarder exposes them. No compiler constant names these functions and no kernel calls them, so the exported MLIR is unchanged. Signed-off-by: nihalpasham <[email protected]>
…t diagram The numeric rename left two rustdoc examples in cute-rs/src/tensor/mod.rs writing the retired E2M1 type in type position; they now read Float4E2M1FN. The work-tile type diagram in dialect-cute labelled the arrow to the four results with tile_idx, the Rust accessor that returns only (m, n, l); the arrow now names the dialect operation that yields all four. Documentation only; no code or generated output changes. Signed-off-by: nihalpasham <[email protected]>
Add Tensor<'a, T, Contiguous1D, T>::local_tile::<TILE>(tile) and the TensorMut twin, each equal to zipped_divide::<TILE>().slice(tile). The name is CuTeDSL's cute.local_tile(tensor, tiler, coord); the Rust form covers a one-dimensional tiler and a single tile coordinate, and the rustdoc says so. Both only change the view. The device compiler recognizes the direct kernel call: CuteFn gains LocalTileRead and LocalTileWrite, classify.rs maps a Contiguous1D receiver with a Tile1D destination of the same access to them, and emit_local_tile inserts cute.tensor_zipped_divide immediately followed by cute.tensor_slice through the same builders as emit_zipped_divide and emit_slice, finishing on the slice result. A kernel that adopts local_tile therefore lowers to the two operations the chain produces, in the same order, and the tensor verifier's one-divide-one-slice rule holds. The collector's Tensor and TensorMut method tables list local_tile so the wrapper body is left to the importer, with positive assertions, and the identity test rejects the fn-item path. The collector test that asserted a store method the Contiguous1D impl does not have now pins the routing that exists: from_disjoint_slice, zipped_divide and local_tile stay stubs on TensorMut, store stays a stub on the Tile1D owner, and store_at on the Contiguous1D owner does not. A host test compares base() and is_full() of local_tile against the two-step chain for a full and a ragged tile on both views. The support matrix cell no longer lists local_tile as absent; its figure is redrawn in the documentation batch. No example changes, so the exported MLIR of the examples is unchanged. Signed-off-by: nihalpasham <[email protected]>
Add cute_rs::copy::tiled::ThrCopy<Atom, ThreadLayout, ValueLayout, TileLayout>, CuTeDSL's ThrCopy, holding one thread's index, and TiledCopy::get_slice(thread_idx) -> ThrCopy, CuTeDSL's TiledCopy.get_slice(thr_idx). ThrCopy::copy(source, tile_coord, destination) takes over the body of TiledCopy::copy, which is removed: it calls copy_g2s with the plan's type parameters and the slice's thread index, so the chain lowers to the one cooperative-copy hook the importer registers, in the same order as before. Both methods are inline-always and ThrCopy lives in copy::tiled, so the importer body-translates them as it did the removed method. The rustdoc says what the Rust form covers: one two-dimensional global tile and one shared destination tile, with the source and destination partitions derived from the plan's layouts inside the compiler operation. The crate root exports ThrCopy beside TiledCopy. The module diagram and docs follow, the size test asserts the slice holds only a u32, and a test checks that get_slice records the index. The two compile-fail probes call get_slice(0).copy(..); their snapshots change only in the call line and caret column, and the expected-type notes are the same. The importer's identity test asserts the ThrCopy::copy path as a source-abstraction path. No example or model calls TiledCopy, so the exported MLIR of the examples is unchanged. Signed-off-by: nihalpasham <[email protected]>
cute_rs::arch gains a sealed MemOrder marker trait, Acquire as its only implementor, and an unsafe fence_sync_restrict::<Order>() forwarder. It carries the CuTeDSL name fence_sync_restrict(order), with the order as a type where the DSL takes a MemOrder string enum, and lowers through ctm_rs::fence::fence_sync_restrict_acquire to cuda_device::barrier::fence_acquire_shared_cluster_cluster, the PTX fence.acquire.sync_restrict::shared::cluster.cluster instruction. The release order has no catalog intrinsic yet and is not offered. The attention pipeline keeps calling the ctm-rs primitive through its existing ctm-rs module imports, so its exported MLIR is unchanged. Signed-off-by: nihalpasham <[email protected]>
RegisterTile<T, N> gains the four operator impls beside its Add impl, over the same element bounds and the same in-place lane loop. They carry the CuTeDSL TensorSSA operator names (-, *, / and unary - on register vectors) and are body-translated by the importer like Add, so no compiler path changes: the importer pins only the type path cute_rs::tensor::RegisterTile. Host tests cover lane order, integer truncation and empty tiles; the support matrix row lists the operators. No fill constructor is added because from_array([value; N]) is the Rust form of cute.full, and no to::<U>() because cute-rs has no element conversion trait and map keeps conversions explicit. Signed-off-by: nihalpasham <[email protected]>
cute_layout::ComposedLayout gains get_swizzle_portion(&self) -> Swizzle and get_nonswizzle_portion(&self) -> &Layout, the CuTeDSL names get_swizzle_portion and get_nonswizzle_portion from core.py. They return what inner() and outer() return today: the inner swizzle applied last and the outer layout applied first in inner(offset + outer(coordinate)). A layout wrapped by from_layout answers with Swizzle::IDENTITY, the (0, 4, 3) swizzle the DSL returns for a plain layout. inner() and outer() stay, since the compiler reads them. A unit test checks both accessors against their existing counterparts. Signed-off-by: nihalpasham <[email protected]>
…odule cute_rs::arch gains cluster_idx() -> (u32, u32, u32), the CuTeDSL cute.arch.cluster_idx reading %clusterid.x/y/z, and cluster_dim() -> (u32, u32, u32), the CuTeDSL cute.arch.cluster_dim reading %nclusterid.x/y/z, the number of clusters in the grid. They lower through new ctm_rs::cluster forwarders (cluster_idx_x/y/z and cluster_grid_dim_x/y/z in its coordinate list) to six hand-written cuda_device::cluster readers beside cluster_idx() and num_clusters(), each wrapping one generated leaf (__cluster_idxX, __cluster_grid_dimX and their Y and Z forms), so the device compiler sees the same llvm.nvvm.read.ptx.sreg.clusterid and nclusterid reads as before. The generated file, the intrinsic catalog and the ABI ledger are untouched. The cluster_dim rustdoc states the clash with the compiler's own cluster_dim vocabulary, which means CTAs per cluster, and the two block_in_cluster rustdocs no longer say the triples are missing. The cluster book page shows the per-axis readers in its position example. Signed-off-by: nihalpasham <[email protected]>
DenseBlockScaleKMajor::offset(row, group) and Sm1xxBlockScaleKMajor::offset(batch, row, group) become crd2idx, the CuTeDSL cute.crd2idx name every other coordinate-to-index map in cute-rs already carries. The bodies are unchanged, so the dense row-major byte index and the 512-byte Blackwell scale-block order lower as before; no importer or collector pin names either method. The block-scaled tensor loaders, the in-crate layout tests and the blockscale_gemm_cute and nvfp4_gemv_cute hosts move to the new name. Signed-off-by: nihalpasham <[email protected]>
…erences The elementwise example's f32 kernel now takes its read and write views through local_tile, while the f16 kernel keeps the zipped_divide and slice chain, so the compiler route behind local_tile is exercised by a committed kernel and both forms stay side by side in one file. The exported MLIR of the local_tile form is structurally identical to the one-statement chain. Three rustdoc notes record where the Rust forms differ from CuTeDSL: the RegisterTile `/` keeps the element type and truncates integer lanes toward zero where the DSL divides integer vectors as f32 and floors with `//`; unary minus is a true negation where the DSL computes 0 - x, so a +0.0 lane differs in sign; and the arch module states that fence_sync_restrict is a primitives-layer name hosted beside the cute.arch forwarders. The MemOrder doc now says what is missing for the release order (a catalog entry bound in cuda-device), and module_paths.rs pins the ThrCopy root export. Signed-off-by: nihalpasham <[email protected]>
The code labels in nine book figures and one qwen3 model figure still printed names retired by the naming cleanup: TmaDesc, make_tma_desc_2d, TmaLoadPipeline, Sm100TmaStorePipeline, Mxfp4TiledMma, accumulate_k64, the scheduler's current_tile().coordinates() and advance(), the PipelineState slot accessor, cluster::block_rank() with a warp::id() that never existed, store_linear, and Bf16 as a type word. They now read TensorMap, create_tensor_map_tiled, PipelineTmaAsync, sm100::PipelineTmaStore, Mxf4TiledMma, gemm, get_current_work().tile_idx() and advance_to_next_work(), index, cluster_ctarank() / warp_id(), store_at and BFloat16. The support matrix figure's tiling cell lists 1D local_tile, the additions figure lists the RegisterTile operators, and the matrix status line names the commit the figure describes. Each label was replaced as a text node with an occurrence check, measured against its box, and the figures were rendered in light and dark mode. Format words that name data encodings (E2M1, E4M3, UE8M0) stay as they are. Signed-off-by: nihalpasham <[email protected]>
…ng paths cute_rs re-exported eight modules under a second path: sm100_attention (attention::sm100), cooperative (copy::cooperative), tiled_copy (copy::tiled), tma (copy::tma), tile (copy::vector), bf16_mma (mma::bf16), block_scaled_mma (mma::block_scaled) and partition (tensor::partition). The one-name rule retires them. Root item exports and the parent-module re-exports in attention and epilogue stay. Callers moved to the defining paths: two doctests in attention/sm100, five integration tests, 25 trybuild probes, the four examples blockscale_gemm_cute, fp16_gemm_256x352_cute, sm100_fmha_prefill_cute and sm100_mla_decode_fp16_cute, and the layouts-and-tiling book page. Eight compile-fail snapshots were regenerated; only path text changed. tests/module_paths.rs keeps the root and parent-module identity checks and drops the alias ones. The README import table is replaced by a short note on root and parent-module re-exports. The registry and collector negative tests keep the alias forms as paths that must not be recognized; only their comments changed. No function body changed, so generated code is unchanged. Signed-off-by: nihalpasham <[email protected]>
…sh the cute docs The comparison page gains concept-map rows for the tensor view, local_tile, register fragments, TMA encoding and copies, pipeline state, the producer step, the persistent scheduler, the tiled copy and register vectors, and a new section that lists every deliberate deviation from CuTeDSL 4.7 with its reason: producer_expect_tx kept outside producer_acquire, has_work and the one-dimensional get_grid_shape, PipelineState::index as usize beside RingCursor::index as u32, TensorMap and create_tensor_map_tiled as the one descriptor and encoder names with TmaCopy::load and store kept apart, cluster_dim following the DSL while the compiler's cluster_dim vocabulary means CTAs per cluster, NamedBarrier::sync alone, fence_sync_restrict::<Acquire> in cute_rs::arch, and the integer division and signed-zero differences on RegisterTile. It also states that TiledCopy, get_slice and ThrCopy exist but no kernel can reach them yet, and that kernels call copy_g2s directly. layouts-and-tiling.md names local_tile as the one-call form of zipped_divide and slice, and the overview lists it. The support matrix status line moves to 5eaa6cb (2026-10-22) with the figure, its tiling cell states the TiledCopy gap, and its primitives row names the barrier and fence families. The ctm-rs README coverage row now lists the named CTA barrier forms and the fence family that its module table already had. No code changed. Signed-off-by: nihalpasham <[email protected]>
The support matrix cell uses the same wording as its figure, the two em dashes on the comparison page are gone, and the elementwise README says that the f32 kernel reads and writes through local_tile while the f16 kernel keeps the zipped_divide and slice chain. Signed-off-by: nihalpasham <[email protected]>
cute_rs::epilogue re-exported the sm120::__compiler module under a second path, which the one-path rule and the crate documentation no longer allow. The re-export and the test that pinned the second path are gone; the negative tests in the importer registry and the codegen collector keep asserting that the old path is never recognized. The compiler hooks keep their defining path, cute_rs::epilogue::sm120::__compiler, and nothing called the retired one. Signed-off-by: nihalpasham <[email protected]>
2 tasks done
1 task done
nihalpasham
added a commit
that referenced
this pull request
Oct 23, 2026
#1615) ## Summary Update `cuda-bindings`, `cuda-core`, and `cuda-core-derive` from `0.4.0` to `0.5.0` in the 17 `Cargo.lock` files introduced by #1598 (`primitives_api_compile` plus the 14 `cute-rs-preview` examples and 2 models) to align with the `v0.5.0` bump in #1614 (`0143af5cc`). ```text Before: 17 cute-rs-preview / fixture Cargo.lock files -> cuda-core 0.4.0 (fails --locked on main) After: 17 cute-rs-preview / fixture Cargo.lock files -> cuda-core 0.5.0 ``` ## Validation - `bash cuda-oxide/scripts/sync-example-locks.sh --check` - `bash cuda-oxide/scripts/check-dependency-licenses.sh` - `bash cuda-oxide/scripts/check-example-license-policy.sh` ## Checklist - [x] All commits signed off (`git commit -s`) Signed-off-by: nihalpasham <[email protected]>
1 task done
nihalpasham
added a commit
that referenced
this pull request
Oct 23, 2026
… clean up docs (#1616) ## Summary Follow-up cleanup after moving `cuda-oxide` to stable Rust `1.99.1` (#1609) and landing `cute-rs-preview` (#1598): ```text 1. cuda-oxide/flake.lock + flake.nix -> bump rust-overlay for Stable 1.99.1 + set RUSTC_BOOTSTRAP=1 2. cargo oxide new (scaffold.rs) -> write .cargo/config.toml with RUSTC_BOOTSTRAP = "1" 3. check-toolchain-parity.sh -> verify flake.lock's rust-overlay rev has the pinned manifest 4. .github/workflows/unit-tests.yml -> drop unused apt llvm-21 install (sysroot llvm-tools provides LLVM 23) 5. CONTRIBUTING / AGENTS / README / book -> update remaining nightly wording + clean up book backticks & SVG ``` ### Changes - **Nix flake (`cuda-oxide/flake.lock`, `cuda-oxide/flake.nix`):** - Bump `oxalica/rust-overlay` from `c1cfa7a2c8` (`2026-10-11`, two days before `1.99.1` was published) to `5f267b7f32` (`2026-10-22`) so `nix develop` finds `Stable 1.99.1`. - Set `RUSTC_BOOTSTRAP = "1"` in `flake.nix` (`cargoOxideCommonArgs` and `devShells.default.shellHook`). - Use a deterministic per-user directory (`${XDG_RUNTIME_DIR:-/tmp}/cuda-oxide-nvidia-driver-$(id -u)`) for non-NixOS driver symlinks so `nix develop --command` does not leak `/tmp/nix-nvidia-driver.XXXXXX` directories. - **Project scaffold (`crates/cargo-oxide/src/commands/scaffold.rs`):** - Write `.cargo/config.toml` (`[env]\nRUSTC_BOOTSTRAP = "1"\n`) in `cargo oxide new` so plain `cargo check`, `cargo build`, and `rust-analyzer` work out of the box in standalone projects on stable `1.99.1`. - **Toolchain parity guard (`scripts/check-toolchain-parity.sh`):** - Verify via an HTTP `HEAD` check that the `oxalica/rust-overlay` revision locked in `cuda-oxide/flake.lock` actually contains `manifests/stable/<channel>.nix` (or `manifests/nightly/...`), and document the `nix flake update rust-overlay` step in `check-toolchain-parity.sh` and `compiler/rustc-public.md`. - **CI workflow (`.github/workflows/unit-tests.yml`, `.github/workflows/status-guard.yml`):** - Remove the redundant `Install LLVM 21` step on `cuda-oxide-codegen` (`LlvmToolchain::resolve` already selects the Rust sysroot's LLVM 23 `llc`/`opt` from `llvm-tools`). - Update the `status-guard.yml` comments to match the current single SIMT stable pin and `flake.lock` check. - **Contributor, agent, and book docs (`CONTRIBUTING.md`, `AGENTS.md`, `cuda-oxide/README.md`, `cuda-oxide-book/`):** - Replace remaining "nightly toolchain" wording in `CONTRIBUTING.md`, `AGENTS.md` (`CLAUDE.md`), `cuda-oxide/README.md`, `getting-started/installation.md`, and `getting-started/hello-gpu.md`. - Remove backticks from `# ` sidebar page titles, `## ` / `### ` section headings, and prose crate/tool names across the `cute/` chapter and `compiler/rustc-public.md`. - Scale bar widths in `models/qwen3_8b_cute/svg/03-rtx5090-benchmark-comparison.svg` so the right-hand throughput labels fit inside the panel with >=45px margin. ## Validation - `cargo test -p cargo-oxide scaffold` - `bash cuda-oxide/scripts/check-toolchain-parity.sh` - `just fmt-check` - `just check-guards` - `just book` (Sphinx HTML build with `-W`) ## Checklist - [x] All commits signed off (`git commit -s`) Signed-off-by: nihalpasham <[email protected]>
juliankang4
added a commit
to juliankang4/cuda-rust
that referenced
this pull request
Oct 25, 2026
Backported from NVIDIA/cuda-rust 927d462 (NVIDIA#1598) by nihalpasham <[email protected]>. Reuse the existing normalization helper under its current name and preserve its type and provenance checks for ordinary Rust calls. Build both non-inlined shared-memory helper variants in the workflow on main and bloomery, and document their call-boundary contract. Signed-off-by: nihalpasham <[email protected]> Signed-off-by: juliankang4 <[email protected]>
3 of 5 tasks
This was referenced Oct 26, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What this adds
A preview of CuTe kernels in Rust, under
cuda-oxide/cute-rs-preview. Layouts, swizzles, operand roles, TMA pipelines and tensor-core MMA are Rust types; the compiler decodes them, verifies the whole module, and hands CUTLASS MLIR to the pinned CUTLASS 4.7 compiler. Two targets,sm_100aandsm_120a; everything else is rejected before compilation.cute-rs(tensors, tiles, copies, pipelines, tiled MMA),ctm-rs(typed instruction primitives: TMA, mbarrier, tcgen05, TMEM, mma.sync),cute-layout(layout algebra and compile-time proofs) anddialect-cute(the Pliron dialect and verifier).mir-importerrecognisescute-rscalls by definition path and emitsdialect-cute;cuda-oxide-mlir-exportmaps MIR, NVVM and CuTe operations to CUTLASS MLIR through the publishedpliron-mlir-interopprinter;cuda-oxide-codegendrives the pinnedlibCutlassCompiler.soand embeds the validated cubin.cargo oxide toolchain install cutlassinstalls the compiler and checks its digests.sm_120a, 4 onsm_100a) and two end-to-end models onsm_120a: Qwen3-TTS 1.7B and Qwen3-8B in BF16 and NVFP4.m16n8k64MXF4 register MMA and three relaxed cluster barrier operations.The preview covers a validated subset of CuTeDSL 4.7 on the two targets, not feature parity: one of eleven tcgen05 MMA kinds, three of seven warp-level kinds, 32x32b TMEM only, fixed epilogues, a static scheduler. The support matrix page lists every row with what the preview has.
What we fixed
cute-rs-previewand the word preview appears where readers look: the folder README, the book chapter and index card, the four crate descriptions.sm_121.Implementation
cuda-oxide-mlir-export/src/profile.rsholds the target gate:CutlassFullCuteMlir22::newacceptssm_100aorsm_120a;tcgen05operations are refused onsm_120aand MXF4 register MMA offsm_120a. The primitives-only profilecutlass-primitives-mlir22has the same gate.dialect-cute/src/verify/runs before any MLIR is built: tensor provenance, pipeline lifecycle, cursor provenance, scheduler grids, MMA placement, epilogue phases, and a ghost-escape audit so compiler-only types never reach runtime storage.mir-importer/src/translator/cute/is a closed world: 88 recognised entry points over 60 operation kinds, 15 audited source-abstraction modules; any othercute_rs::call is an error.mainrenumbered the four new ABI ledger entries toi1039toi1042behind main'si1030toi1038, kept main's lockfiles and let cargo add the new crates, and regenerated every catalog-derived file;check,probe --allandcheck-abi-history --base-ref mainpass.cute_rs::definition paths are unchanged.Verification
Checked at
56384d465:just fmt-check,just clippy,just test,just test-cuda,just check-guards,just doc-checkandjust check-intrinsics origin/mainpass on the merged tree.cargo test --workspace --all-targets --no-fail-fastover 115 test binaries,cargo test -p oxide-artifacts --features object, and the nestedrustc-codegen-cudaworkspace pass. One pre-existing race,ptx::tests::debug_builds_fail_when_opt_drops_the_debug_graph, failed once under parallel load (a just-written temp file reported busy) and passed three times alone.mainwith 1,042 entries.scripts/sync-example-locks.sh --checkafter aligning the 13 cute example lockfiles with the standard examples.scripts/check-primitives-api-compile.pyfor theapi,fp16andblockscalefixtures, and the 13 CUTLASS consumer tests with the pinned 4.7.0 library.elementwise_cute,rmsnorm_cute,swiglu_cute,ffn_cute,rope_cache_cute,blockscale_gemm_cute,blockscale_gemm_ctm,nvfp4_gemv_cute,fmha_prefill_sm120a_cuteandfmha_decode_sm120a_cuterun and pass their reference comparisons; the FMHA pair also times against CuTeDSL 4.7.0 (prefill 18.0 us against 23.1 us, decode 4.4 us against 12.0 us). The foursm_100aexamples compile.vecadd,mma_mxf8f6f4andclusterrun on the native path.scripts/smoketest.sh --compile-only, which publishes to the shared backend cache; the examples-compile lane covers it here.Caveats
sm_100f,sm_103a,sm_110a,sm_120f,sm_121a, Hopper and Ampere are rejected at the profile gate.CUDA_OXIDE_DEVICE_BACKEND=cutlass-mlir; the native backend has no CuTe lowering and currently stops with a verification error rather than a dedicated message.cute.dot; the NVFP4 tensor-core atom exists in the dialect and exporter but has no producer yet.cute-rsanddialect-cutestay for now.