Test/sample improvements + block-scale & SDPA fixes (9.18–9.24 fuzzer mining)

include/cudnn_frontend_Tensor.h

@@ -447,7 +447,7 @@ class TensorBuilder_v8 {
                 set_error_and_throw_exception(
                     &m_tensor,
                     status,
-                    "CUDNN_BACKEND_TENSOR_DESCRIPTOR: SetAttribute CUDNN_ATTR_TENSOR_BYTE_ALIGNMENT Failed");
+                    "CUDNN_BACKEND_TENSOR_DESCRIPTOR: SetAttribute CUDNN_ATTR_TENSOR_IS_VIRTUAL Failed");
                 return std::move(m_tensor);
             }
         }

samples/cpp/sdpa/mxfp8_fwd.cpp

@@ -126,12 +126,11 @@ TEST_CASE("sdpa_mxfp8_fprop", "[graph][sdpa][mxfp8][forward]") {
                                      .set_data_type(fe::DataType_t::FP8_E8M0)
                                      .set_reordering_type(fe::TensorReordering_t::F8_128x4));
 
-    // Block scale tensor for V (FP8_E8M0 with F8_128x4 reordering)
-    // V scales the s (sequence) dimension since BMM2 (S @ V) contracts on s_kv
-    // The contracting dimension (s_scale) must be contiguous, so use COL_MAJOR-like strides
-    auto SF_V_dims = std::vector<int64_t>({b, h, s_scale_padded, d_padded});
-    auto SF_V_strides =
-        std::vector<int64_t>({h * s_scale_padded * d_padded, s_scale_padded * d_padded, 1, s_scale_padded});
+    // Block scale tensor for V (FP8_E8M0, F8_128x4). V is block-scaled along s; declared
+    // d-contiguous (stride[3]==1) uniformly with SF_Q/SF_K -- the kernel reads SF via the
+    // F8_128x4 swizzle, so the declared inner stride is not load-bearing.
+    auto SF_V_dims    = std::vector<int64_t>({b, h, s_scale_padded, d_padded});
+    auto SF_V_strides = std::vector<int64_t>({h * s_scale_padded * d_padded, s_scale_padded * d_padded, d_padded, 1});
 
     auto SF_V = mha_graph.tensor(fe::graph::Tensor_attributes()
                                      .set_name("SF_V")

test/python/sdpa/random_config.py

@@ -9,8 +9,14 @@
 # fmt: off
 
 def generate_test_seeds(*, num_tests, rng_seed):
+    # Env overrides for focused/expanded sweeps without editing each suite:
+    #   MHAS_NUM_TESTS   - cap (or expand) the number of configs (e.g. small subset for sanitizer runs)
+    #   MHAS_SEED_OFFSET - shift the geometry RNG seed to explore a fresh set of configs
+    import os
+    rng_seed = rng_seed + int(os.environ.get("MHAS_SEED_OFFSET", "0"))
+    n = int(os.environ.get("MHAS_NUM_TESTS", "0")) or num_tests
     rng = random.Random(rng_seed)
-    return [(i+1, num_tests, rng.randint(65536, 2147483647)) for i in range(num_tests)]
+    return [(i+1, n, rng.randint(65536, 2147483647)) for i in range(n)]
 
 
 def get_strides_from_indices(shape, indices=[0, 1, 2, 3], gaps=[0, 0, 0, 0], rng_geom=None):
@@ -498,6 +504,7 @@ def __init__(
             s_kv_min = min(s_kv_min, s_kv_max)
         self.s_q_gen = RandomIntValue(min=s_q_min, max=s_q_max)
         self.s_kv_gen = RandomIntValue(min=s_kv_min, max=s_kv_max)
+        self._s_q_max = s_q_max  # post-cap upper bound, used to clamp the s_q>s_kv branch
         self.distribution = RandomChoice(s_q_distribution)
 
     def __call__(self, rng):
@@ -510,6 +517,16 @@ def __call__(self, rng):
             s_q = 1
         elif distribution == "s_q=s_kv":
             s_q = s_kv
+        elif distribution == "s_q>s_kv":
+            # Query longer than key/value (cross-attention, chunked prefill).
+            # The default "s_q=random" branch below structurally caps s_q<=s_kv, so
+            # this regime is otherwise never exercised. Caught class: NVBug 5829882
+            # (SDPA hang when S_Q > S_KV).
+            # Clamp to s_q_max so we don't overshoot the suite bound (esp. the SM80
+            # OOM cap) and trigger spurious torch CUDA OOM instead of testing cuDNN.
+            if s_kv < self._s_q_max:
+                s_q = min(s_kv + rng.randint(1, max(1, s_kv)), self._s_q_max)
+            # else: s_kv already at the cap -> no room for s_q>s_kv within bounds; leave s_q.
         else:
             s_q = self.s_q_gen(rng)
 

test/python/test_low_precision_matmul.py

@@ -277,12 +277,11 @@ def create_matmul_dequantize_graph(cudnn_handle, A, B, A_descale, B_descale, BLO
 
     with cudnn.graph(cudnn_handle) as (g, _):
 
-        batch_size, M, N, K = (
-            A_descale.shape[0],
-            A_descale.shape[1],
-            B_descale.shape[1],
-            A_descale.shape[2],
-        )
+        # Derive dims from the data tensors (float4_e2m1fn_x2 packs two fp4 per byte).
+        batch_size = A.shape[0]
+        M = A.shape[1]
+        K = A.shape[2] * 2
+        N = B.shape[2] * 2
 
         A_cudnn_tensor = g.tensor(
             name="tensor_a",
@@ -298,18 +297,23 @@ def create_matmul_dequantize_graph(cudnn_handle, A, B, A_descale, B_descale, BLO
             data_type=convert_to_cudnn_type(B.dtype),
         )
 
+        # F8_128x4 descale: block dim (K) reduced by BLOCK_SIZE & rounded to 4, contiguous; other dim to 128.
+        k_scale = ((K + BLOCK_SIZE - 1) // BLOCK_SIZE + 3) // 4 * 4
+        m_pad = (M + 127) // 128 * 128
+        n_pad = (N + 127) // 128 * 128
+
         A_descale_tensor = g.tensor(
             name="block_descale_a",
-            dim=A_descale.shape,
-            stride=(M * K, K, 1),
+            dim=(batch_size, m_pad, k_scale),
+            stride=(m_pad * k_scale, k_scale, 1),
             data_type=convert_to_cudnn_type(A_descale.dtype),
             reordering_type=cudnn.tensor_reordering.F8_128x4,
         )
 
         B_descale_tensor = g.tensor(
             name="block_descale_b",
-            dim=B_descale.shape,
-            stride=(N * K, 1, K),
+            dim=(batch_size, k_scale, n_pad),
+            stride=(k_scale * n_pad, 1, k_scale),
             data_type=convert_to_cudnn_type(B_descale.dtype),
             reordering_type=cudnn.tensor_reordering.F8_128x4,
         )
@@ -413,8 +417,12 @@ def test_low_precision_fp4_matmul(cudnn_handle):
     A = _bfloat16_to_float4_e2m1fn_x2(A_ref)
     B = _bfloat16_to_float4_e2m1fn_x2(B_ref)
 
-    A_descale = torch.full((batch_size, M, K), 1.0, dtype=torch.float8_e4m3fn, device="cuda")
-    B_descale = torch.full((batch_size, K, N), 1.0, device="cuda", dtype=torch.float8_e4m3fn)
+    # Canonical F8_128x4 block-reduced descale shapes (K reduced by BLOCK_SIZE & rounded to 4; other dim to 128).
+    k_scale = ((K + BLOCK_SIZE - 1) // BLOCK_SIZE + 3) // 4 * 4
+    m_pad = (M + 127) // 128 * 128
+    n_pad = (N + 127) // 128 * 128
+    A_descale = torch.full((batch_size, m_pad, k_scale), 1.0, dtype=torch.float8_e4m3fn, device="cuda")
+    B_descale = torch.full((batch_size, k_scale, n_pad), 1.0, device="cuda", dtype=torch.float8_e4m3fn)
 
     g, uids = create_matmul_dequantize_graph(cudnn_handle, A, B, A_descale, B_descale, BLOCK_SIZE)
 

test/python/test_matmul_fuzzer.py

@@ -326,6 +326,42 @@ def generate(self) -> MatmulConfig:
         N = self.random_dim()
         K = self.random_dim()
 
+        # Degenerate / GEMV shapes. random_dim() squares a sqrt-range int then rounds up
+        # to a multiple of 8, so M/N/K are always >= 8 and never 1 — the matrix-vector
+        # (M=1 / N=1) and tiny-K corners are structurally unreachable. Force them ~12% of
+        # the time. Caught class: NVBug 5866266 (force_jit core dump on a degenerate
+        # matmul) and 5990039 (no engine for IDENTITY/degenerate matmul).
+        #
+        # OPT-IN (default off): enabling this surfaced a real FORT-native matmul illegal
+        # memory access on K=1 with an int8 operand + large M (SM90, dev 9.30 + rel 9.24;
+        # root-caused to get_kernel_access_size() unclamped — bug filed separately). An IMA
+        # corrupts the CUDA context and kills the test process (it cannot be caught as a
+        # skip), so this axis is gated behind MATMUL_FUZZ_DEGENERATE=1 until that backend
+        # bug is fixed, after which the default can be flipped on.
+        if os.environ.get("MATMUL_FUZZ_DEGENERATE", "0") == "1" and self.rng.random() < 0.12:
+            kind = self.rng.choice(['m1', 'n1', 'tiny_k', 'm1_tiny_k', 'k1'])
+            if kind == 'm1':
+                M = 1
+            elif kind == 'n1':
+                N = 1
+            elif kind == 'tiny_k':
+                K = self.rng.choice([1, 2, 4, 8, 16])
+            elif kind == 'm1_tiny_k':
+                M = 1; K = self.rng.choice([1, 8, 16])
+            else:  # k1
+                K = 1
+
+        # Unaligned leading dims (K/N not a multiple of 4). random_dim() rounds up to mult-of-8,
+        # so the bits_per_access<32 corner is normally unreachable -- that's where the FORT-native
+        # engine takes the LDG+STS smem-staging path, which over-runs for a widening cast (int8/fp8
+        # -> fp16/fp32): silent wrong-results when K is unaligned, IMA when N is. Opt-in (IMA aborts).
+        if os.environ.get("MATMUL_FUZZ_UNALIGNED", "0") == "1" and self.rng.random() < 0.15:
+            unaligned = self.rng.choice([1, 2, 3, 5, 6, 7, 9, 13, 17, 30])  # bpa<32 for 8-bit load
+            if self.rng.random() < 0.5:
+                K = unaligned
+            else:
+                N = unaligned
+
         # Data types - ensure compatible combinations
         if self.rng.random() < 0.8:
             # 80% of tests use same dtype for A and B (more stable)
@@ -492,8 +528,19 @@ def compute_reference(config: MatmulConfig, A: torch.Tensor, B: torch.Tensor, bi
 
 
 def run_cudnn_matmul(config: MatmulConfig, A: torch.Tensor, B: torch.Tensor, C: torch.Tensor,
-                     bias: Optional[torch.Tensor], cudnn_handle) -> Tuple[bool, str]:
-    """Run matmul using cuDNN and return success status and message."""
+                     bias: Optional[torch.Tensor], cudnn_handle,
+                     det_reruns: Optional[int] = None) -> Tuple[bool, str]:
+    """Run matmul using cuDNN and return success status and message.
+
+    If det_reruns > 0, re-execute the *same built plan* that many extra times into a
+    freshly garbage-poisoned output (and re-poisoned workspace) and assert the output
+    hash is identical run-to-run. This catches nondeterministic kernels (lost-update /
+    smem-pipeline races, uninitialized split-K workspace) without needing a golden ref
+    and without re-running heuristics. Caught class: NVBug 5897577 (sm120 matmul buffer
+    race), 6140341 (matmul-fuzzer numeric regression), 6217343-class.
+    """
+    if det_reruns is None:
+        det_reruns = int(os.environ.get("MATMUL_DET_RERUNS", "2"))
     try:
         stream = torch.cuda.current_stream().cuda_stream
         cudnn.set_stream(handle=cudnn_handle, stream=stream)
@@ -544,18 +591,31 @@ def run_cudnn_matmul(config: MatmulConfig, A: torch.Tensor, B: torch.Tensor, C:
         # Build and execute
         graph.validate()
         graph.build_operation_graph()
+        if det_reruns > 0:
+            # The determinism rerun-assert below must only judge plans cuDNN declares
+            # deterministic: a legitimate FP-atomic split-K plan carries the
+            # NONDETERMINISTIC numeric note and is *expected* to vary run-to-run, so
+            # deselect those here to avoid a spurious determinism failure.
+            try:
+                graph.deselect_numeric_notes([cudnn.numerical_note.NONDETERMINISTIC])
+            except Exception:
+                pass
         graph.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
         graph.check_support()
         graph.build_plans(cudnn.build_plan_policy.HEURISTICS_CHOICE)
 
         # Allocate workspace and fill with garbage to catch uninitialized memory bugs
         workspace_size = graph.get_workspace_size()
         workspace = torch.empty(workspace_size, device='cuda', dtype=torch.uint8)
-        if workspace_size > 0:
-            # Fill with random garbage + some NaN patterns to test proper workspace init
-            workspace.random_(0, 256)
-            nan_mask = torch.rand(workspace_size, device='cuda') < 0.1
-            workspace[nan_mask] = 0xFF
+
+        def poison_workspace():
+            if workspace_size > 0:
+                # Fill with random garbage + some NaN patterns to test proper workspace init
+                workspace.random_(0, 256)
+                nan_mask = torch.rand(workspace_size, device='cuda') < 0.1
+                workspace[nan_mask] = 0xFF
+
+        poison_workspace()
 
         # Build variant pack
         variant_pack = {A_tensor: A, B_tensor: B, result: C}
@@ -566,6 +626,20 @@ def run_cudnn_matmul(config: MatmulConfig, A: torch.Tensor, B: torch.Tensor, C:
         graph.execute(variant_pack, workspace, handle=cudnn_handle)
         torch.cuda.synchronize()
 
+        # Determinism check: re-execute the SAME plan into a re-poisoned output and
+        # workspace; the output hash must be bit-identical run-to-run.
+        if det_reruns > 0:
+            h0 = print_tensor_stats(C, tag=None)
+            for _ in range(det_reruns):
+                fill_with_garbage(C)            # re-poison output: kernel must fully overwrite
+                poison_workspace()              # re-poison workspace: catches uninit split-K
+                graph.execute(variant_pack, workspace, handle=cudnn_handle)
+                torch.cuda.synchronize()
+                hi = print_tensor_stats(C, tag=None)
+                if hi != h0:
+                    return False, (f"NONDETERMINISTIC output: hash 0x{h0 >> 32:08X} != "
+                                   f"0x{hi >> 32:08X} across reruns (same plan, re-poisoned)")
+
         return True, "success"
 
     except cudnn.cudnnGraphNotSupportedError as e:
@@ -711,8 +785,8 @@ def get_test_params(request):
 
 
 # Fixed test list for default runs
-DEFAULT_NUM_TESTS = 2048
-DEFAULT_SEED = 42
+DEFAULT_NUM_TESTS = int(os.environ.get("MATMUL_NUM_TESTS", "2048"))
+DEFAULT_SEED = int(os.environ.get("MATMUL_FUZZ_SEED", "42"))
 TEST_PARAMS = tlist_with_configs(num_tests=DEFAULT_NUM_TESTS, rng_seed=DEFAULT_SEED)