From 6f9518a86fb8c7f94ce6918f3cbf6a4ea5464fb0 Mon Sep 17 00:00:00 2001 From: nihalp Date: Sun, 11 Oct 2026 16:28:46 +0530 Subject: [PATCH 1/2] fix(cuda-oxide): honor the sparse FP8 metadata contract Signed-off-by: nihalp --- .../cuda-device/src/generated/sparse_mma.rs | 8 +- .../src/render/families/mma.rs | 6 ++ .../src/render/tests/mma.rs | 44 ++++++-- .../cuda-intrinsics/src/generated/abi_v1.rs | 8 +- .../examples/sparse_mma_fp8/src/main.rs | 101 +++++++----------- cuda-oxide/intrinsics/generated-reference.md | 8 +- 6 files changed, 92 insertions(+), 83 deletions(-) diff --git a/cuda-oxide/crates/cuda-device/src/generated/sparse_mma.rs b/cuda-oxide/crates/cuda-device/src/generated/sparse_mma.rs index a1e8bd0f63..886487ef99 100644 --- a/cuda-oxide/crates/cuda-device/src/generated/sparse_mma.rs +++ b/cuda-oxide/crates/cuda-device/src/generated/sparse_mma.rs @@ -416,7 +416,7 @@ pub unsafe fn mma_sp_m16n8k32_s32_u8_satfinite( /// # Safety /// All 32 lanes must execute the same instruction with the same qualifiers, and no lane may have exited. /// `c`, `a`, and `b` must contain this lane's 4-, 4-, and 4-register fragments; `metadata` must contain its sparse metadata. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// `selector` must be the compile-time constant `0`. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). /// This register-only operation is not a memory fence. @@ -440,7 +440,7 @@ pub unsafe fn mma_sp_m16n8k64_f32_e4m3_e4m3_f32( /// # Safety /// All 32 lanes must execute the same instruction with the same qualifiers, and no lane may have exited. /// `c`, `a`, and `b` must contain this lane's 4-, 4-, and 4-register fragments; `metadata` must contain its sparse metadata. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// `selector` must be the compile-time constant `0`. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). /// This register-only operation is not a memory fence. @@ -464,7 +464,7 @@ pub unsafe fn mma_sp_m16n8k64_f32_e4m3_e5m2_f32( /// # Safety /// All 32 lanes must execute the same instruction with the same qualifiers, and no lane may have exited. /// `c`, `a`, and `b` must contain this lane's 4-, 4-, and 4-register fragments; `metadata` must contain its sparse metadata. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// `selector` must be the compile-time constant `0`. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). /// This register-only operation is not a memory fence. @@ -488,7 +488,7 @@ pub unsafe fn mma_sp_m16n8k64_f32_e5m2_e4m3_f32( /// # Safety /// All 32 lanes must execute the same instruction with the same qualifiers, and no lane may have exited. /// `c`, `a`, and `b` must contain this lane's 4-, 4-, and 4-register fragments; `metadata` must contain its sparse metadata. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// `selector` must be the compile-time constant `0`. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). /// This register-only operation is not a memory fence. diff --git a/cuda-oxide/crates/cuda-intrinsics-gen/src/render/families/mma.rs b/cuda-oxide/crates/cuda-intrinsics-gen/src/render/families/mma.rs index f2de275c66..87ddd80577 100644 --- a/cuda-oxide/crates/cuda-intrinsics-gen/src/render/families/mma.rs +++ b/cuda-oxide/crates/cuda-intrinsics-gen/src/render/families/mma.rs @@ -366,6 +366,12 @@ pub(in crate::render) const SPARSE_MMA_ORDERED_TF32_METADATA_RULE: &str = pub(in crate::render) fn sparse_mma_metadata_rule(mma: &SparseMma) -> &'static str { match (mma.metadata, mma.a_element) { + // PTX 8.4 requires these six ordered index pairs for plain FP8 mma.sp + // too; this does not change the instruction to mma.sp::ordered_metadata. + // https://docs.nvidia.com/cuda/archive/12.4.1/parallel-thread-execution/index.html#warp-level-sparse-matrix-storage + (SparseMmaMetadata::Standard, SparseMmaElement::E4m3 | SparseMmaElement::E5m2) => { + SPARSE_MMA_ORDERED_METADATA_RULE + } (SparseMmaMetadata::Standard, _) => SPARSE_MMA_STANDARD_METADATA_RULE, (SparseMmaMetadata::Ordered, SparseMmaElement::Tf32) => { SPARSE_MMA_ORDERED_TF32_METADATA_RULE diff --git a/cuda-oxide/crates/cuda-intrinsics-gen/src/render/tests/mma.rs b/cuda-oxide/crates/cuda-intrinsics-gen/src/render/tests/mma.rs index 67123ff4e6..c9fb824f4f 100644 --- a/cuda-oxide/crates/cuda-intrinsics-gen/src/render/tests/mma.rs +++ b/cuda-oxide/crates/cuda-intrinsics-gen/src/render/tests/mma.rs @@ -791,6 +791,38 @@ fn register_mma_kind_is_explicit_and_legacy_ir_stays_compatible() { assert!(targets.contains("kind_matches && accumulator_matches")); } +#[test] +fn sparse_fp8_metadata_contract_requires_ordered_indices_without_changing_integer_rules() { + let repo_root = Path::new(env!("CARGO_MANIFEST_DIR")).join("../.."); + let catalog = crate::resolve::resolve(&repo_root).unwrap(); + for id in [ + "mma_sp_m16n8k64_f32_e4m3_e4m3_f32", + "mma_sp_m16n8k64_f32_e4m3_e5m2_f32", + "mma_sp_m16n8k64_f32_e5m2_e4m3_f32", + "mma_sp_m16n8k64_f32_e5m2_e5m2_f32", + ] { + let record = sparse_mmas(&catalog) + .find(|record| record.id == id) + .unwrap(); + assert_eq!( + sparse_mma_metadata_rule(record.sparse_mma.as_ref().unwrap()), + "Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior.", + "{id}" + ); + assert!(sparse_mma_template(record).starts_with("mma.sp.sync.aligned.")); + } + for id in ["mma_sp_m16n8k64_s32_s8", "mma_sp_m16n8k64_s32_s4"] { + let record = sparse_mmas(&catalog) + .find(|record| record.id == id) + .unwrap(); + assert_eq!( + sparse_mma_metadata_rule(record.sparse_mma.as_ref().unwrap()), + "Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior.", + "{id}" + ); + } +} + #[test] fn sparse_mma_rendering_enforces_selector_and_keeps_family_distinct() { let repo_root = Path::new(env!("CARGO_MANIFEST_DIR")).join("../.."); @@ -1072,8 +1104,8 @@ fn sparse_mma_rendering_enforces_selector_and_keeps_family_distinct() { assert!(targets.contains("GeneratedSparseMmaOverflow::NotApplicable")); assert!(targets.contains("GeneratedHardwareAlternative::ExactArchitecture(120)")); - assert_eq!(raw.matches(SPARSE_MMA_STANDARD_METADATA_RULE).count(), 36); - assert_eq!(raw.matches(SPARSE_MMA_ORDERED_METADATA_RULE).count(), 88); + assert_eq!(raw.matches(SPARSE_MMA_STANDARD_METADATA_RULE).count(), 32); + assert_eq!(raw.matches(SPARSE_MMA_ORDERED_METADATA_RULE).count(), 92); assert_eq!( raw.matches(SPARSE_MMA_ORDERED_TF32_METADATA_RULE).count(), 2 @@ -1082,13 +1114,13 @@ fn sparse_mma_rendering_enforces_selector_and_keeps_family_distinct() { compatibility .matches(SPARSE_MMA_STANDARD_METADATA_RULE) .count(), - 36 + 32 ); assert_eq!( compatibility .matches(SPARSE_MMA_ORDERED_METADATA_RULE) .count(), - 88 + 92 ); assert_eq!( compatibility @@ -1146,11 +1178,11 @@ fn sparse_mma_rendering_enforces_selector_and_keeps_family_distinct() { assert!(reference.contains("LLVM source record uses A, B, C, metadata, selector order")); assert_eq!( reference.matches(SPARSE_MMA_STANDARD_METADATA_RULE).count(), - 36 + 32 ); assert_eq!( reference.matches(SPARSE_MMA_ORDERED_METADATA_RULE).count(), - 88 + 92 ); assert_eq!( reference diff --git a/cuda-oxide/crates/cuda-intrinsics/src/generated/abi_v1.rs b/cuda-oxide/crates/cuda-intrinsics/src/generated/abi_v1.rs index ce0289c0dd..0d77f9e1d4 100644 --- a/cuda-oxide/crates/cuda-intrinsics/src/generated/abi_v1.rs +++ b/cuda-oxide/crates/cuda-intrinsics/src/generated/abi_v1.rs @@ -9248,7 +9248,7 @@ pub unsafe fn i0169( /// `_arg0`, `_arg1`, and `_arg2` hold this lane's 4-register C, 4-register A, and 4-register B fragments; `_arg3` holds its sparse metadata. /// `_arg4` must be the compile-time constant `0`. /// The operation is register-only and is not a memory fence. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). #[must_use] #[inline(never)] @@ -9274,7 +9274,7 @@ pub unsafe fn i1026( /// `_arg0`, `_arg1`, and `_arg2` hold this lane's 4-register C, 4-register A, and 4-register B fragments; `_arg3` holds its sparse metadata. /// `_arg4` must be the compile-time constant `0`. /// The operation is register-only and is not a memory fence. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). #[must_use] #[inline(never)] @@ -9300,7 +9300,7 @@ pub unsafe fn i1027( /// `_arg0`, `_arg1`, and `_arg2` hold this lane's 4-register C, 4-register A, and 4-register B fragments; `_arg3` holds its sparse metadata. /// `_arg4` must be the compile-time constant `0`. /// The operation is register-only and is not a memory fence. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). #[must_use] #[inline(never)] @@ -9326,7 +9326,7 @@ pub unsafe fn i1028( /// `_arg0`, `_arg1`, and `_arg2` hold this lane's 4-register C, 4-register A, and 4-register B fragments; `_arg3` holds its sparse metadata. /// `_arg4` must be the compile-time constant `0`. /// The operation is register-only and is not a memory fence. -/// Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. +/// Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. /// See the [PTX sparse MMA fragment layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-mma-sp). #[must_use] #[inline(never)] diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/src/main.rs index 7e0f14c7f6..45c9b4ba13 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/src/main.rs @@ -3,8 +3,8 @@ //! Checks all four plain SM89 sparse FP8 MMA forms against an exact host GEMM. //! -//! Each form covers all twelve standard-metadata codes and zero/nonzero C. -//! Values vary across K bands so swapped fragment registers change the result. +//! Each form covers all six valid FP8 metadata codes and zero/nonzero C. +//! Values vary across K bands so swapped B fragment registers change the result. //! See the PTX ISA's "Matrix Fragments for sparse mma.m16n8k64" for the layout. //! `--probe` additionally compares candidate layouts on the current GPU. @@ -15,7 +15,7 @@ use cuda_device::{DisjointSlice, cuda_module, kernel, thread, wmma}; const M: usize = 16; const N: usize = 8; const K: usize = 64; -const METADATA_CODES: usize = 12; +const METADATA_CODES: usize = 6; const VARIANTS: usize = 2 * 4 * METADATA_CODES; /// Compressed K: two kept elements per 4-wide group. const KC: usize = K / 2; @@ -59,9 +59,19 @@ fn f8_bits(v: u32, e4m3: bool) -> u32 { } } -/// Standard metadata accepts two distinct 2-bit indices, in either order. +/// Plain FP8 `mma.sp` requires two positions in increasing order, even without +/// `::ordered_metadata`: `0x4` encodes (0, 1), but reversed `0x1` is undefined. +/// See PTX 8.4's FP8-specific rule under "Sparse matrix storage": +/// . fn nibble_of(i: u32) -> u32 { - probe_nibble(i) + match i { + 0 => 0x4, + 1 => 0x8, + 2 => 0x9, + 3 => 0xc, + 4 => 0xd, + _ => 0xe, + } } #[cuda_module] @@ -132,7 +142,7 @@ mod kernels { } /// Variant `v` occupies `out[v*128 + lane*4 .. +4]`, with - /// `v = nonzero_accumulator*48 + form*12 + metadata_code`. + /// `v = (nonzero_accumulator*4 + form)*METADATA_CODES + metadata_code`. #[kernel] pub fn oracle(mut out: DisjointSlice) { let l = thread::threadIdx_x(); @@ -211,14 +221,9 @@ mod kernels { let mut mi = 0u32; while mi < META_OPTS { - let nib = probe_nibble(mi % 12); - // `sw` exchanges the two 2-bit fields of the nibble. - let word = if mi / 12 == 0 { - nib - } else { - ((nib & 3) << 2) | (nib >> 2) - }; - let meta = word * 0x1111_1111; + // Vary fragment layouts, but never reverse the metadata's + // two indices: those FP8 inputs have undefined behavior. + let meta = nibble_of(mi) * 0x1111_1111; let d = unsafe { wmma::mma_sp_m16n8k64_f32_e4m3_e4m3_f32(zero, a, b, meta, 0) }; let p = combo * 128 + (l as usize) * 4; let mut j = 0usize; @@ -292,7 +297,7 @@ fn run_oracle(ctx: &std::sync::Arc) { } } println!( - " {name:12} {:>7} all 12 standard metadata codes: {} mismatches", + " {name:12} {:>7} all {METADATA_CODES} FP8 metadata codes: {} mismatches", if acc == 1 { "C!=0," } else { "C=0," }, form_bad ); @@ -301,7 +306,7 @@ fn run_oracle(ctx: &std::sync::Arc) { } assert_eq!(bad, 0, "sparse FP8 MMA accumulator mismatches"); println!( - "SUCCESS: all 4 sparse FP8 m16n8k64 forms x 12 metadata codes x C=0/C!=0; \ + "SUCCESS: all 4 sparse FP8 m16n8k64 forms x {METADATA_CODES} metadata codes x C=0/C!=0; \ all 32 lanes and 4 logical accumulators/lane match host GEMM exactly" ); } @@ -316,12 +321,12 @@ fn run_oracle(ctx: &std::sync::Arc) { // A order `byte_ord(ao, i)` -- byte position i -> column offset // B band `b_band(bb, n)` -- which dense K rows a register holds // B order `byte_ord(bo, i)` -- byte position i -> K offset -// code the nibble and which of its two 2-bit fields names the even column +// code one of the six valid FP8 metadata nibbles // // Every candidate is filled into the registers and executed once; the host -// compares the result with an exact GEMM whose answer does not depend on the -// candidate at all. A wrong axis changes the sum, so only the true wiring -// reproduces it. +// compares the result with a GEMM using that candidate's layout and metadata. +// Exact matches are diagnostic candidates: repeated input patterns can leave +// more than one layout indistinguishable. // --------------------------------------------------------------------------- const A_ROW_OPTS: u32 = 4; @@ -329,11 +334,9 @@ const A_BAND_OPTS: u32 = 6; const A_ORD_OPTS: u32 = 2; const B_BAND_OPTS: u32 = 6; const B_ORD_OPTS: u32 = 2; -const META_SWAP_OPTS: u32 = 2; - const A_LAYOUTS: u32 = A_ROW_OPTS * A_BAND_OPTS * A_ORD_OPTS; const B_LAYOUTS: u32 = B_BAND_OPTS * B_ORD_OPTS; -const META_OPTS: u32 = 12 * META_SWAP_OPTS; +const META_OPTS: u32 = METADATA_CODES as u32; const COMBOS: u32 = A_LAYOUTS * B_LAYOUTS * META_OPTS; /// Separate K-band patterns for the optional layout sweep. @@ -392,35 +395,6 @@ fn byte_ord(o: u32, i: u32) -> u32 { if o == 0 { i } else { 3 - i } } -/// The twelve legal nibbles: two 2-bit fields that differ. -fn probe_nibble(idx: u32) -> u32 { - if idx == 0 { - 0x1 - } else if idx == 1 { - 0x2 - } else if idx == 2 { - 0x3 - } else if idx == 3 { - 0x4 - } else if idx == 4 { - 0x6 - } else if idx == 5 { - 0x7 - } else if idx == 6 { - 0x8 - } else if idx == 7 { - 0x9 - } else if idx == 8 { - 0xb - } else if idx == 9 { - 0xc - } else if idx == 10 { - 0xd - } else { - 0xe - } -} - #[derive(Clone, Copy)] struct Candidate { rp: u32, @@ -429,7 +403,6 @@ struct Candidate { bb: u32, bo: u32, nib: u32, - sw: u32, } impl Candidate { @@ -471,11 +444,7 @@ impl Candidate { fn kmap(&self, c: u32) -> u32 { let low = self.nib & 3; let high = self.nib >> 2; - let field = if c.is_multiple_of(2) == (self.sw == 0) { - low - } else { - high - }; + let field = if c.is_multiple_of(2) { low } else { high }; 4 * (c / 2) + field } @@ -546,8 +515,7 @@ fn run_probe(ctx: &std::sync::Arc) { ao: ai % A_ORD_OPTS, bb: bi / B_ORD_OPTS, bo: bi % B_ORD_OPTS, - nib: probe_nibble(mi % 12), - sw: mi / 12, + nib: nibble_of(mi), }; let idx = total; total += 1; @@ -567,8 +535,8 @@ fn run_probe(ctx: &std::sync::Arc) { for c in &exact { let (ap, bp) = c.coverage(); println!( - " exact rp={} ab={} ao={} bb={} bo={} nib={:#x} sw={} (A {ap}/512, B {bp}/512 placed)", - c.rp, c.ab, c.ao, c.bb, c.bo, c.nib, c.sw + " exact rp={} ab={} ao={} bb={} bo={} nib={:#x} (A {ap}/512, B {bp}/512 placed)", + c.rp, c.ab, c.ao, c.bb, c.bo, c.nib ); } } @@ -607,15 +575,18 @@ mod tests { } #[test] - fn metadata_covers_every_distinct_pair() { + fn metadata_covers_every_valid_fp8_pair() { let mut seen = [false; 16]; for index in 0..METADATA_CODES as u32 { let code = nibble_of(index) as usize; - assert_ne!(code & 3, code >> 2); assert!(!seen[code]); seen[code] = true; } - assert_eq!(seen.iter().filter(|&&present| present).count(), 12); + // FP8 requires the first stored value's position to precede the second. + // This checks both that reversed pairs are absent and no valid pair is lost. + for (code, present) in seen.into_iter().enumerate() { + assert_eq!(present, (code & 3) < (code >> 2), "metadata {code:#x}"); + } } #[test] diff --git a/cuda-oxide/intrinsics/generated-reference.md b/cuda-oxide/intrinsics/generated-reference.md index 71c68648ee..a90564c21d 100644 --- a/cuda-oxide/intrinsics/generated-reference.md +++ b/cuda-oxide/intrinsics/generated-reference.md @@ -5349,10 +5349,10 @@ This table is generated from the resolved catalog. Evidence stages below disting - `mma_sp_m16n8k32_s32_u8_s8` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k32.row.col.s32.u8.s8.s32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0` or `1`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Integer overflow wraps. Runtime validation is not executed on a GPU. - `mma_sp_m16n8k32_s32_u8_s8_satfinite` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k32.row.col.satfinite.s32.u8.s8.s32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0` or `1`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Integer overflow uses finite saturation. Runtime validation is not executed on a GPU. - `mma_sp_m16n8k32_s32_u8_satfinite` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k32.row.col.satfinite.s32.u8.u8.s32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0` or `1`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Integer overflow uses finite saturation. Runtime validation is not executed on a GPU. -- `mma_sp_m16n8k64_f32_e4m3_e4m3_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e4m3.e4m3.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. -- `mma_sp_m16n8k64_f32_e4m3_e5m2_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e4m3.e5m2.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. -- `mma_sp_m16n8k64_f32_e5m2_e4m3_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e5m2.e4m3.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. -- `mma_sp_m16n8k64_f32_e5m2_e5m2_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e5m2.e5m2.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. +- `mma_sp_m16n8k64_f32_e4m3_e4m3_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e4m3.e4m3.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. +- `mma_sp_m16n8k64_f32_e4m3_e5m2_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e4m3.e5m2.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. +- `mma_sp_m16n8k64_f32_e5m2_e4m3_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e5m2.e4m3.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. +- `mma_sp_m16n8k64_f32_e5m2_e5m2_f32` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.f32.e5m2.e5m2.f32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must be `0x4`, `0x8`, `0x9`, `0xc`, `0xd`, or `0xe`; any other value is undefined behavior. Overflow mode is not applicable. Runtime validation is not executed on a GPU. - `mma_sp_m16n8k64_s32_s4` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.s32.s4.s4.s32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0` or `1`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Integer overflow wraps. Runtime validation is not executed on a GPU. - `mma_sp_m16n8k64_s32_s4_satfinite` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.satfinite.s32.s4.s4.s32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0` or `1`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Integer overflow uses finite saturation. Runtime validation is not executed on a GPU. - `mma_sp_m16n8k64_s32_s4_u4` takes fragments in C, A, B, metadata, selector order and lowers to one convergent, register-only `mma.sp.sync.aligned.m16n8k64.row.col.s32.s4.u4.s32` instruction. Its LLVM source record uses A, B, C, metadata, selector order. The selector must be the compile-time constant `0` or `1`. Every non-exited warp lane must execute the same instruction and qualifiers. Every 4-bit metadata group must encode two distinct 2-bit indices; `0x0`, `0x5`, `0xa`, and `0xf` are undefined behavior. Integer overflow wraps. Runtime validation is not executed on a GPU. From f46ce3f764ecc4af8491881a0425421e1570fe91 Mon Sep 17 00:00:00 2001 From: nihalp Date: Sun, 11 Oct 2026 18:53:21 +0530 Subject: [PATCH 2/2] Use default Cargo features in rust-analyzer Signed-off-by: nihalp --- .vscode/settings.json | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.vscode/settings.json b/.vscode/settings.json index f04c586f3e..a06fba831c 100644 --- a/.vscode/settings.json +++ b/.vscode/settings.json @@ -2,7 +2,7 @@ // Project-level editor settings, tracked in git. Keep personal preferences // (fonts, colours, keybindings) in your VS Code User settings instead. "rust-analyzer.rustc.source": "discover", - "rust-analyzer.cargo.features": "all", + "rust-analyzer.cargo.features": [], "rust-analyzer.procMacro.enable": true, "rust-analyzer.procMacro.attributes.enable": true, // "rust-analyzer.cargo.extraEnv": {