Repository navigation
feat(intrinsics): add ordered sparse Ampere floating MMA - #1139
Conversation
|
nihalpasham
left a comment
There was a problem hiding this comment.
- 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
0x4and0xe, 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_80floor 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.
a04d9e4 to
0c2810b
Compare
|
@nihalpasham The benchmark used 256 warps, 1,000 dependency-chained
Plain |
There was a problem hiding this comment.
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 + colover 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 currentskip sm_XYmakes 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).
…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>
0c2810b to
61c0b21
Compare
|
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
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.
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. |
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>
Summary
Adds first-class generated APIs for eight PTX 8.5
mma.sp::ordered_metadataAmpere 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 plainmma.spagainst 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
i1017..i1024after the 1,016-entrymainledger.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 checkpasses (the local mirror of CI: fmt, clippy, tests, guards, docs)cargo oxide run <example>passes, orscripts/smoketest.sh -o '^<example>$'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
git commit -s)