Repository navigation
cuda-oxide: move SIMT toolchain pin to stable 1.99.1 - #1609
Open
nihalpasham wants to merge 104 commits into
Open
nihalpasham wants to merge 104 commits into
nihalpasham wants to merge 104 commits into
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
Signed-off-by: nihalpasham <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
…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 <nihalp@nvidia.com>
…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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
…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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
…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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
…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 <nihalp@nvidia.com>
…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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
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 <nihalp@nvidia.com>
Pin `cuda-oxide/rust-toolchain.toml` (and matching scaffold, devcontainer, and book/README copies) to official stable Rust `1.99.1` with `RUSTC_BOOTSTRAP=1` configured via `.cargo/config.toml`, `Justfile`, and `cargo-oxide`. `1.99.1` ships LLVM 23 (`23.1.1`) with an NVPTX backend and TableGen intrinsic definitions identical to `nightly-2026-10-11`, plus the `rustc-dev` (`librustc_driver.so` and `librustc_public.rlib`) component, removing daily nightly drift while preserving the full `rustc_public` and `rustc-codegen-cuda` pipeline. Signed-off-by: nihalpasham <nihalp@nvidia.com>
nihalpasham
requested review from
binarybana,
cryos,
elibol,
kkraus14,
roivanov and
rparolin
as code owners
October 22, 2026 16:48
This branch has not been deployed
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
1 participant
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.
Summary
Move
cuda-oxidefrom the pinnednightly-2026-10-11toolchain to official stable Rust1.99.1:1.99.1across all SIMT toolchain surfaces:cuda-oxide/rust-toolchain.toml,cargo oxide new(scaffold.rs),.devcontainer/devcontainer.json,intrinsics/upstream.lock, and the book/README install commands.RUSTC_BOOTSTRAP=1automatically: Set incuda-oxide/.cargo/config.toml,cuda-oxide/Justfile,cargo-oxide(backend.rs,codegen_env.rs,fmt.rs,host_cargo.rs), and the example CI loops so-Zcompiler flags,#![feature(rustc_private)],#![feature(f16)], andcargo-features = ["profile-rustflags"]work transparently on stable1.99.1.1.99.1compiler APIs:cuda-host: Usecore::intrinsics::type_id::<T>()intype_id_u128.mir-importer: Querytcx.fn_abi_of_instanceinsidefacts.rs(reference_param_validity) for typedArgAttributes(NoAlias,NoFree,pointee_size), and matchValueAbi::Scalarintypes.rs.rustc-codegen-cuda: Match1.99.1'sCodegenBackend(init,target_config,CompiledModule),ShimKind::FnPtrAddr, andlayout.peel_transparent_wrappers.cuda-oxide-book: Add the toolchain stabilization,facts.rsupstreaming workflow, and standalone bundle sections with four SVG diagrams undercuda-oxide-book/compiler/images/.Validation
just check&just book:fmt-check,clippy(workspace + backend + all 258 example/model manifests),test,test-cuda,check-guards,check-intrinsics(all 1,042 routes / 2,427 target probes onLLVM version 23.1.1-rust-1.99.1-stable),doc-check, and Sphinx HTML book build (-W) all pass.scripts/smoketest.shon RTX 5090): 240 / 240 pass.cute-rs-previewexamples and models: 16 / 16 pass (all 12sm_120aworkloads execute and pass on RTX 5090; all 4sm_100aworkloads compile to validated direct-ABI cubins via CUTLASS 4.7).Checklist
git commit -s)