fix(cuda-oxide): use valid sparse FP8 metadata

.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": {

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.

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

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

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)]

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":
+/// <https://docs.nvidia.com/cuda/archive/12.4.1/parallel-thread-execution/index.html#warp-level-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<f32>) {
         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<CudaContext>) {
                 }
             }
             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<CudaContext>) {
     }
     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,24 +321,22 @@ fn run_oracle(ctx: &std::sync::Arc<CudaContext>) {
 //   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;
 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<CudaContext>) {
                     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<CudaContext>) {
     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]