jinchen62 · Oct 31, 2026 · Oct 29, 2026 · Oct 29, 2026 · Oct 31, 2026 · Oct 31, 2026
@@ -183,6 +183,9 @@ struct Arguments
     int32_t fast_check_repeat;
     int32_t fast_check_inject;
     int8_t  requires_streamk;
+    // Self-test of the checks of D across solutions: the position in the solution list of a
+    // solution that is never launched, which the allclose check must report (-1 = none)
+    int32_t skip_launch_solution;
     // integer_exact value pattern: '' (standard), ternary or sparse_k (hipblaslt_init.hpp)
     char    integer_exact_pattern[16];
     // skip, rather than fail, when the library offers no solution (size-threshold sweeps)
@@ -328,6 +331,7 @@ struct Arguments
     OPER(fast_check_repeat) SEP      \
     OPER(fast_check_inject) SEP      \
     OPER(requires_streamk) SEP       \
+    OPER(skip_launch_solution) SEP   \
     OPER(integer_exact_pattern) SEP  \
     OPER(allow_no_solution) SEP      \
     OPER(timing) SEP                 \

@@ -1236,6 +1236,35 @@ inline void report_failed_solution(int64_t                failures_before,
                        << solution_description(handle, algo, sol, count, true) << std::endl;
 }
 
+// The skip_launch_solution self-test, after check(): the solution that was never launched must
+// fail the checks of D, and every other solution must pass them. Returns whether sol is the
+// solution that was never launched.
+inline bool check_skip_launch_self_test(const Arguments&       arg,
+                                        int64_t                failures_before,
+                                        hipblasLtHandle_t      handle,
+                                        hipblasLtMatmulAlgo_t& algo,
+                                        size_t                 sol,
+                                        size_t                 count)
+{
+    if(arg.skip_launch_solution < 0)
+        return false;
+    const bool skipped  = int64_t(sol) == arg.skip_launch_solution;
+    const bool reported = hipblaslt_bench_options::verification_failures() > failures_before;
+    if(skipped != reported)
+    {
+        const std::string error
+            = "skip_launch_solution: " + solution_description(handle, algo, sol, count, true)
+              + (skipped ? " was never launched, but D passed its checks"
+                         : " failed the checks of D");
+#ifdef GOOGLE_TEST
+        ADD_FAILURE() << error;
+#else
+        hipblaslt_cerr << error << std::endl;
+#endif
+    }
+    return skipped;
+}
+
 void check(hipStream_t                   stream,
            const Arguments&              arg,
            const uint32_t&               gemm_count,
@@ -2485,6 +2514,31 @@ void testing_matmul_with_bias(const Arguments& arg,
         }
     }
 
+    // The self-test counts the allclose failures that check() reports per solution: unit_check
+    // and norm_check would fail the case outright on the NaN in the skipped solution's D. The
+    // skipped solution follows one that wrote D.
+    if(arg.skip_launch_solution != -1)
+    {
+        std::string why;
+        if(arg.skip_launch_solution < 1)
+            why = "skip_launch_solution must be -1 (none) or a position after the first solution";
+        else if(!arg.allclose_check || arg.unit_check || arg.norm_check || arg.fast_check
+                || arg.c_equal_d)
+            why = "skip_launch_solution requires allclose_check, without unit_check, norm_check, "
+                  "fast_check or c_equal_d";
+        else if(arg.use_ext || arg.grouped_gemm || batchMode != HIPBLASLT_BATCH_MODE_STRIDED)
+            why = "skip_launch_solution supports only the C API with strided batches";
+        if(!why.empty())
+        {
+#ifdef GOOGLE_TEST
+            FAIL() << why;
+#else
+            hipblaslt_cerr << why << std::endl;
+            return;
+#endif
+        }
+    }
+
     IntegerExactPattern      iePattern = IntegerExactPattern::standard;
     IntegerExactPatternScope iePatternScope;
     if(arg.integer_exact_pattern[0])
@@ -5407,6 +5461,20 @@ void testing_matmul_with_bias(const Arguments& arg,
         }
     }
 
+    if(arg.skip_launch_solution >= int64_t(heuristicResult.size()))
+    {
+        const std::string why = "skip_launch_solution " + std::to_string(arg.skip_launch_solution)
+                                + " needs at least " + std::to_string(arg.skip_launch_solution + 1)
+                                + " solutions, and this case has "
+                                + std::to_string(heuristicResult.size());
+#ifdef GOOGLE_TEST
+        GTEST_SKIP() << why;
+#else
+        hipblaslt_cout << why << std::endl;
+        return;
+#endif
+    }
+
     // A placed workspace replaces the normal one, which placement runs never use: placement
     // requires fast_check, and fast_check refuses timing runs.
     const bool placeWorkspace = !strcmp(arg.placement, "workspace");
@@ -5907,6 +5975,7 @@ void testing_matmul_with_bias(const Arguments& arg,
             CHECK_HIP_ERROR(err);
         }
     };
+    bool skippedSolutionChecked = false;
     if(!arg.timing)
     {
         // fast_check launches and checks each solution fast_check_repeat times, solution by
@@ -6094,7 +6163,7 @@ void testing_matmul_with_bias(const Arguments& arg,
                                                           stream),
                                           HIPBLAS_STATUS_SUCCESS);
                 }
-                else
+                else if(int64_t(sol) != arg.skip_launch_solution)
                 {
                     CHECK_HIP_ERROR(hipStreamSynchronize(stream));
                     EXPECT_HIPBLAS_STATUS(hipblasLtMatmul(handle,
@@ -6241,6 +6310,12 @@ void testing_matmul_with_bias(const Arguments& arg,
                                        heuristicResult[sol].algo,
                                        sol,
                                        heuristicResult.size());
+                skippedSolutionChecked |= check_skip_launch_self_test(arg,
+                                                                      failures_before,
+                                                                      handle,
+                                                                      heuristicResult[sol].algo,
+                                                                      sol,
+                                                                      heuristicResult.size());
             }
             if(arg.fast_check)
             {
@@ -6860,6 +6935,12 @@ void testing_matmul_with_bias(const Arguments& arg,
                                        heuristicResult[sol].algo,
                                        sol,
                                        heuristicResult.size());
+                skippedSolutionChecked |= check_skip_launch_self_test(arg,
+                                                                      failures_before,
+                                                                      handle,
+                                                                      heuristicResult[sol].algo,
+                                                                      sol,
+                                                                      heuristicResult.size());
                 first_launch_failed
                     = hipblaslt_bench_options::verification_failures() > failures_before;
             }
@@ -7077,6 +7158,8 @@ void testing_matmul_with_bias(const Arguments& arg,
             {
                 // One launch of candidate sol on rotation block i % block_count, cold or timed.
                 auto launchMatmul = [&](size_t sol, int64_t i) {
+                    if(int64_t(sol) == arg.skip_launch_solution)
+                        return HIPBLAS_STATUS_SUCCESS;
                     int  b          = static_cast<int>(i % block_count);
                     auto ptr_matmul = matmul[b][0];
                     return hipblasLtMatmul(
@@ -7311,6 +7394,17 @@ void testing_matmul_with_bias(const Arguments& arg,
         }
     }
 
+    if(arg.skip_launch_solution >= 0 && !skippedSolutionChecked)
+    {
+        const std::string error = "skip_launch_solution: solution "
+                                  + std::to_string(arg.skip_launch_solution) + " was never checked";
+#ifdef GOOGLE_TEST
+        ADD_FAILURE() << error;
+#else
+        hipblaslt_cerr << error << std::endl;
+#endif
+    }
+
     for(auto it : ptrs)
     {
         CHECK_HIP_ERROR(hipFree(it));

@@ -135,6 +135,7 @@ void Arguments::init()
     fast_check_repeat = 1;
     fast_check_inject = -1;
     requires_streamk  = 0;
+    skip_launch_solution = -1;
     integer_exact_pattern[0] = 0;
     allow_no_solution        = 0;
     timing         = 0;

@@ -655,6 +655,7 @@ Arguments:
   - fast_check_repeat: c_int32
   - fast_check_inject: c_int32
   - requires_streamk: c_int8
+  - skip_launch_solution: c_int32
   - integer_exact_pattern: c_char*16
   - allow_no_solution: c_int8
   - timing: c_int8
@@ -773,6 +774,7 @@ Defaults:
   fast_check_repeat: 1
   fast_check_inject: -1
   requires_streamk: 0
+  skip_launch_solution: -1
   integer_exact_pattern: ''
   allow_no_solution: 0
   timing: 0

@@ -3358,6 +3358,39 @@ Tests:
   # shape on gfx1250-strict (hardware run on PR #13504).
   gpu_arch: '1201'
 
+###############################################################################
+# Checks of D across solutions (ROCM-32357).
+#
+# A case that checks several solutions launches them all into the same D, in
+# the hipblaslt-test loop and in the hipblaslt-bench candidate search (timing
+# 1). D is filled with the fast_check NaN sentinel before each solution (all-ones
+# bytes for D types the sentinel does not cover), so a solution that never
+# stores D fails instead of passing on the previous solution's result.
+###############################################################################
+
+# Self-test: solution 1 is never launched. The case passes only if the allclose
+# check reports solution 1 and no other solution.
+# Runtime on gfx942: 4 cases, under 1 s.
+- name: matmul_check_reports_unlaunched_solution
+  category: pre_checkin
+  function:
+    matmul:
+      - *single_precision
+      - *hpa_bf16_precision
+  matrix_size:
+    - { M: 131, N: 67, K: 200 }
+  transA: N
+  transB: T
+  alpha: 2.0
+  beta: -2.0
+  requested_solution_num: 4
+  skip_launch_solution: 1
+  unit_check: 0
+  allclose_check: 1
+  timing: [0, 1]
+  iters: 0
+  cold_iters: 0
+
 ###############################################################################
 # fast_check: verification without a host reference GEMM.
 #

@@ -509,6 +509,8 @@ namespace
 
                 if(arg.c_equal_d)
                     name << "_C_EQUAL_D";
+                if(arg.timing)
+                    name << "_TIMING";
                 if(arg.placement[0])
                     name << "_PLACE_" << arg.placement;
                 if(arg.integer_exact_pattern[0])