Skip to content

feat(intrinsics): add ordered sparse Ampere floating MMA - #1139

Merged
nihalpasham merged 6 commits into
NVIDIA:mainfrom
almightychang:sparse-mma-float
Aug 30, 2026
Merged

nihalpasham merged 6 commits into
NVIDIA:mainfrom
almightychang:sparse-mma-float

Conversation

@almightychang

@almightychang almightychang commented Aug 22, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Adds first-class generated APIs for eight PTX 8.5 mma.sp::ordered_metadata Ampere floating-point register-MMA forms (f16, bf16, and tf32; sm_80+). On an RTX 5080 (sm_120), a memory-fed f16 register-MMA inner loop measured a median kernel time of 0.4039 ms for plain mma.sp against 0.2084 ms for the ordered form. Two independent runs put the ratio at 1.93x and 1.94x. This is an inner-loop microbenchmark, not an end-to-end GEMM claim. This relates to tracking issue #274 but does not close it, and closes no issue.

Changes

  • Admit the eight reviewed forms through a schema-63 compact overlay and generate their Rust APIs, ABI declarations, dialect/import/lowering routes, target matching, probes, compile coverage, and reference documentation.
  • Append ABI IDs i1017..i1024 after the 1,016-entry main ledger.
  • Expose the ordered metadata operand and selector domains, including TF32's 1:2 metadata domain, while retaining the warp-uniformity and valid-metadata safety contract.
  • Add sparse_mma_ordered_float, a normal-path example with nonzero lane-distributed inputs and an independent host GEMM oracle covering all eight variants, all 32 lanes, and four logical accumulators per lane.

Testing

  • just check passes (the local mirror of CI: fmt, clippy, tests, guards, docs)
  • cargo oxide run <example> passes, or scripts/smoketest.sh -o '^<example>$'
  • New example added (if applicable)

Reproduced on head 0c2810ba:

  • CUDA_TOOLKIT_PATH=/home/sam/Workspaces/dinov3-candle/.pixi/envs/default just check — passed.
  • CUDA_TOOLKIT_PATH=/home/sam/Workspaces/dinov3-candle/.pixi/envs/default cargo oxide run sparse_mma_ordered_float --arch sm_120 — all eight variants, 32 lanes, and four logical accumulators per lane matched the host oracle on an RTX 5080 with driver 580.173.02.
  • CUDA_TOOLKIT_PATH=/home/sam/Workspaces/dinov3-candle/.pixi/envs/default scripts/smoketest.sh --no-color --only '^sparse_mma_ordered_float$' --keep-logs — 1/1 passed.
  • compute-sanitizer --tool memcheck crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/target/release/sparse_mma_ordered_float — 0 errors.
  • cargo run -p cuda-intrinsics-gen -- check-abi-history --base-ref 1f6a0fdd — preserves 1,016 base entries and has 1,024 total entries.
  • bash scripts/check-book-catalog-stamp.sh — passed (a74497e1).

LLVM 21 end-to-end lowering and CUDA 13.1.1 ptxas/libNVVM remain unverified locally.

Checklist

  • All commits signed off (git commit -s)
  • SPDX headers on new source files

@nihalpasham nihalpasham added enhancement New feature or request intrinsics Device intrinsics and libdevice math mappings codegen Device code-generation pipeline (Rust MIR to IR to PTX) IR-lowering Lowering between dialects (dialect-mir to LLVM dialect) cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) device-apis User-facing device-side and kernel-authoring APIs (cuda-device, cuda-macros) tests Adds or improves test coverage examples Example kernels and host harnesses under crates/rustc-codegen-cuda/examples documentation Improvements or additions to documentation safety Memory safety, soundness, or undefined behavior soundness-audit Correctness case under active maintainer audit; not cleared for merge yet labels Aug 23, 2026
@nihalpasham

nihalpasham commented Aug 23, 2026 •

Copy link
Copy Markdown
Collaborator

@nihalpasham nihalpasham added needs-implementation Depends Blocked on another PR/issue or an upstream dependency and removed safety Memory safety, soundness, or undefined behavior labels Aug 23, 2026

@nihalpasham nihalpasham left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

  • Benefit required: name the tensor-memory workload and state whether ordered collection improves performance, correctness, or essential API coverage. The zero-filled example does not establish the value of adding this feature.
  • Metadata blocker: TF32 ordered variants are valid only for immediate kinds 0x4 and 0xe, but the generated API currently accepts six values.
  • Test blocker: the example's normal-path smoketest classification is not established, its zero-filled A/B inputs cannot expose many operand or lane-selection errors, and it checks only lanes 0 and 3 of an 8-lane result.
  • Dependency blocker: the generated ABI tail depends on #1138, which is not merge-ready, so this branch must be regenerated after that dependency is repaired.
  • LLVM 21/22 lowering, libNVVM constraints, and the PTX 8.5 / sm_80 floor are otherwise coherent in local compile/assembly checks.
  • Complete this with an independent nonzero GPU oracle over all lanes and a sanitizer run. This is related to the #274 tracking epic, but it does not close it.

@almightychang

almightychang commented Aug 24, 2026 •

Copy link
Copy Markdown
Contributor Author

@nihalpasham
Concrete benefit: on an RTX 5080 (sm_120, CUDA 12.9.86), ordered metadata improved a memory-fed sparse-MMA inner loop.

The benchmark used 256 warps, 1,000 dependency-chained m16n8k16.row.col.f32.f16.f16.f32 MMAs per thread, 10 warmups, and 30 CUDA-event samples per form. Each iteration loaded fresh A, B, and metadata operands from global memory:

  • mma.sp::ordered_metadata: 0.209440 ms median
  • plain mma.sp: 0.403968 ms median
  • plain / ordered: 1.929× (median)

Plain mma.sp calls a metadata-normalization helper before each HMMA; the ordered form emits the HMMA directly. This is an inner-loop result for 2:4-pruned f16/bf16 GEMM, not an end-to-end GEMM claim.

@nihalpasham nihalpasham left a comment •

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The four earlier blockers are addressed on paper, and the contracts check out against the PTX ISA (metadata domains, selector ranges, PTX 8.5 / sm_80 floor, operand order in all eight templates). But the GPU oracle still can't catch the errors it exists to catch, so the test blocker stands.

The problem: every fragment comes from val(x, y) = 1 + ((x + y) & 1), which only depends on parity. Walking the lane mappings:

aa(l, n): parity(r + q) = (l/4) & 1       // the 8*_ and 2*_ terms are even
bb(l, n): parity(k + c) = (l/4) & 1       // same value
at/bt:    parity = ((l/4) + (l & 3)) & 1  // same as each other

So A and B register arrays are bitwise identical in every lane, in all eight variants (C too for f16.f16 k16). Swap a and b in any template and the hardware sees the same bits: the oracle passes. Nothing else catches it either: a/b register counts are equal in all variants, and the probes exercise the LLVM route, not these templates.

Needed:

  • oracle inputs where A, B, and C are positionally distinct and asymmetric, e.g. matrix_id * 1000 + row * K + col over a small exact-in-f16 range
  • at least one nonzero selector per selector domain, and a second metadata code (e.g. 0x8), executed at runtime; today a lowering that hardwires selector 0 passes
  • the sub-sm_80 message in the harness form skipping: ... (the current skip sm_XY makes pre-Ampere hosts report FAIL)

Sequencing: this hardcodes first_abi_id = "i1017" and derives IDs positionally, while #1093 removes that pattern crate-wide and #1138 also claims the ledger tail. #1093 lands first; this then adopts blank-ID + bind_generated_abi_ids, and whichever tail PR lands second regenerates.

After the rework: fresh adversarial pass on the new head, then our local validation (5090 run, smoketest, ABI history, LLVM 21 spot-check since the recorded evidence is LLVM 22 only).

almightychang and others added 3 commits August 30, 2026 15:36
…tf32, sm_80+)

Add the eight PTX 8.5 mma.sp::ordered_metadata floating forms through the
catalog/generator pipeline, symmetric to the ordered integer sparse family:

- m16n8k8  f32.tf32          - m16n8k16 f32.tf32
- m16n8k16 f32.bf16          - m16n8k32 f32.bf16
- m16n8k16 f16.f16/f32.f16   - m16n8k32 f16.f16/f32.f16

Contracts follow the PTX ISA exactly: operand order d,a,b,c,metadata,selector;
selector as an "n" immediate with domain 0-3 for the 2-register carriers
(k8, k16-f16/bf16) and 0-1 for the 4-register carriers (k32, k16-tf32);
metadata codes 0x4/0x8/0x9/0xc/0xd/0xe (tf32: 0x4/0xe only, with its own
reference rule text). The f16/bf16 forms carry <2 x half> multiplicands in
LLVM; tf32 carries packed i32.

Ported from PR NVIDIA#1139 onto the split resolve/model/render module layout and
the ledger-bound blank-ABI-id convention (the shard carries no first_abi_id
or product_count; identity binds from intrinsics/abi-v1.toml by catalog id).
Evidence identity fields are rebound to the pinned LLVM 23.1.0 backend,
matching how all recorded evidence was migrated in the LLVM 23 bump.

Signed-off-by: Sam, Joochul Chang <almightychang@icloud.com>
Signed-off-by: nihalpasham <nihalp@nvidia.com>
…and multi-code

The ported oracle was parity-degenerate: val(x,y) = 1 + ((x+y) & 1) made
the A and B fragments bitwise identical in every lane across all eight
variants, metadata was the single code 0x4 repeated, and the selector was
always 0. An a<->b template swap, a hardwired selector, or a metadata
misroute all still passed. Rework:

- A, B, C are positionally distinct and asymmetric (A = 1+(3r+5j)%7,
  B = 1+(2k+3c)%8, C = (5r+7c)%32), small enough that f16/bf16/tf32
  accumulation stays exact and the host reference is pure integer math
- every valid ordered metadata code executes (0x4/0x8/0x9/0xc/0xd/0xe
  cycling by row+chunk; tf32 alternates 0x4/0xe with an aperiodic index
  so no wrong (lane, nibble) -> (row, chunk) reading aliases to the truth)
- each selector domain runs a nonzero selector (2/1/3 in the 0-3 domain,
  1 in the 0-1 domain) and non-selected lanes carry a valid decoy pattern,
  so selector misrouting changes the result
- pair-selector shapes (k32 f16/bf16, k16 tf32) split the K dimension
  across the thread pair while keeping the [row g | row g+8] register
  halves; verified against sm_120 hardware, where a row-per-thread split
  fails the oracle
- the sub-sm_80 skip message now uses the harness "skipping: ..." form

Mutation check on an RTX 5090 (cc 12.0): swapping the a/b arguments of one
variant fails 127 of its 128 checked accumulators; reverted, all 8 variants
pass and compute-sanitizer memcheck reports 0 errors. The evidence runtime
stages now record this re-measured run.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
Append the eight ordered Ampere float sparse-MMA entries to
intrinsics/abi-v1.toml (catalog ids, operation keys, and raw Rust
signatures derived from the closed recipe; blank shard ids bind to them
by catalog id) and regenerate every catalog-stamped output:

- intrinsics/catalog.json, generated-reference.md, and the 8 new
  selector-swept probe fixtures (plus the stamp header in every probe)
- generated sources across cuda-device, dialect-nvvm, mir-importer,
  mir-lower, and the collector targets
- the book's catalog stamp and 1016 -> 1024 record count

Gates run on this tree: cargo run -p cuda-intrinsics-gen -- generate /
check / check-abi-history --base-ref origin/main (1016 base entries
preserved, 1024 total); probe --all --skip-terminal --per-target (2373
target configurations across all 1024 routes, pinned LLVM 23.1.0); cargo
test -p cuda-intrinsics-gen -p dialect-nvvm -p mir-importer -p mir-lower;
fmt + clippy -D warnings for the workspace and the new example;
scripts/sync-example-locks.sh --check.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
…UDA 13.3 toolchain

Signed-off-by: nihalpasham <nihalp@nvidia.com>
@nihalpasham

Copy link
Copy Markdown
Collaborator

Full re-review, and this time we finished it ourselves; the structural rework was all fallout from our own landings (#1093 ledger-bound ABI ids, #1190 module split, LLVM 23), and the two of us round-tripping it again felt like a waste of your 1.93x. Your commit is the base of the new head, authorship and sign-off intact.

What held up under two independent checks: all eight instruction contracts (operand order d,a,b,c,metadata,selector; selector immediate domains 0-3 vs 0-1 per carrier width; tf32 metadata 0x4/0xe). Those are byte-for-byte your templates.

What we changed

  1. The oracle. The old inputs were parity-symmetric, which made the test blind to the exact bugs it existed to catch:
old:  val(x,y) = 1 + ((x+y) & 1)     ->  A and B fragments bitwise
      META = 0x4444_4444, sel = 0        IDENTICAL in every lane;
                                         an a/b swap in any template passes

new:  A = 1+(3r+5j)%7,  B = 1+(2k+3c)%8,  C = (5r+7c)%32
      all six metadata codes cycle (tf32: both), nonzero selector
      per selector domain, decoy metadata in non-selected lanes

Proof it has teeth now: we mutated one template's a<->b on the RTX 5090 and 127/128 of that variant's accumulators failed; reverted, re-passed. The old oracle passed that same mutation.

  1. The structural redo on main: overlay shard rewritten ledger-bound (no first_abi_id/product_count), generator edits ported into the split resolve//model//render/ layout, ABI ids i1017-i1024, skip message in the harness skipping: form.

  2. Evidence re-attested on the local CUDA 13.3 toolchain. Two reasons: repo precedent (every evidence file on main uses local tools), and it wasn't optional; the per-target probe hard-fails on a recorded tool path that doesn't exist here (recorded evidence tool does not exist: /home/sam/...). Everything re-measured: llc PTX for all 8 probes, local ptxas assembly, the libNVVM route end to end (nvJitLink cubin, SASS shows exactly 8 HMMA.SP, one per variant), and runtime on the 5090 through BOTH routes, memcheck clean.

Gates on the final head: generate/check round-trip, ledger history vs main (1016 preserved, 1024 total), per-target probe of all 8 incl. terminal assembly, probe --all gate (2373 configs / 1024 routes), 310 generator tests, fmt/clippy, per-example gates, lockfile check, smoketest PASS.

Ordering note: #1101 takes i1017 and lands first; this regenerates after (one command, ours). Dropping Depends/needs-implementation; CI approval + landing next.

@nihalpasham nihalpasham added the perf Performance of generated code or of the compiler itself label Aug 30, 2026
@nihalpasham nihalpasham removed Depends Blocked on another PR/issue or an upstream dependency needs-implementation labels Aug 30, 2026
The evidence re-attestation renamed the libNVVM evidence profile, which
is recorded in intrinsics/catalog.json, so the catalog hash moved again
after the page was stamped for the i1017-i1024 regeneration:

  stamp    a6962ae6 -> 91c449f0
  counts   unchanged (the eight float forms resolve onto the shared
           sparse-MMA structural op; entries were already 1024)

Signed-off-by: nihalpasham <nihalp@nvidia.com>
…il to i1018-i1025

NVIDIA#1101 landed mbarrier_arrive_no_complete as i1017, which this branch
had claimed for the eight ordered sparse float forms. Resolution:

  - abi-v1.toml: main's i1017 kept, our eight entries appended as
    i1018-i1025 (same relative order)
  - the two count asserts move to the merged totals (1025 entries,
    65 shards); the pinned-binding tests move to 1018..=1025
  - every generated file resolved by regeneration, not by hand
  - book dialect page re-stamped (20bc41c2, 1025 entries; op counts
    unchanged, the float forms share the sparse-MMA structural op)

Re-validated: generate/check round-trip, ledger history vs main
(preserves 1017 base entries, 1025 total), probe gate (2374 configs /
1025 routes), per-target terminal probes of all eight plus the
mbarrier form, 310 generator tests, fmt, all nine status-guard
scripts, lockfile check, and a fresh RTX 5090 run of
sparse_mma_ordered_float (all 8 variants match the host GEMM).

Signed-off-by: nihalpasham <nihalp@nvidia.com>
@nihalpasham
nihalpasham merged commit 8b41ac1 into NVIDIA:main Aug 30, 2026
52 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

codegen Device code-generation pipeline (Rust MIR to IR to PTX) cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) device-apis User-facing device-side and kernel-authoring APIs (cuda-device, cuda-macros) documentation Improvements or additions to documentation enhancement New feature or request examples Example kernels and host harnesses under crates/rustc-codegen-cuda/examples intrinsics Device intrinsics and libdevice math mappings IR-lowering Lowering between dialects (dialect-mir to LLVM dialect) perf Performance of generated code or of the compiler itself soundness-audit Correctness case under active maintainer audit; not cleared for merge yet tests Adds or improves test coverage

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants