From 354e24a3f326bbb4ab498ed349d4c56913e831c2 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Thu, 29 Oct 2026 10:58:15 +0000 Subject: [PATCH 1/3] fix(hipblaslt): poison D before checking each solution in the matmul client hipblaslt-bench -v launches every candidate solution into the same D and checks each one in turn, and so does the hipblaslt-test loop outside fast_check, but D was only restored between solutions under --c_equal_d. A solution that never stores D was therefore checked against the previous solution's result and reported that solution's error. On gfx942 before #12797, three CustomGSUs split-K kernels (indices 57 to 59) passed --algo_method all -v with the 1.8065e-05 norm_error of the generated solution checked before them, and failed with norm_error 1 when run alone. Fill D with all ones, NaN in the IEEE floating types, before each solution is launched whenever a result check is on and D is not C: in the candidate-search prepare hook and in the hipblaslt-test loop. fast_check keeps its own sentinel, and --c_equal_d keeps restoring C from the host. JIRA ID: ROCM-32357 Co-authored-by: Cursor --- .../clients/common/include/testing_matmul.hpp | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index f76a19f7b367..3dab7c7de966 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -5891,6 +5891,14 @@ void testing_matmul_with_bias(const Arguments& arg, } } } + // Every solution checked in this call writes the same D, so a solution that skips its store + // would be checked against the previous solution's result. Filling D with all ones first (NaN + // in the IEEE floating types) makes such a solution fail instead. + const bool checksResults = arg.unit_check || arg.norm_check || arg.allclose_check; + auto poisonD = [&]() { + for(auto& d : *dDp) + CHECK_HIP_ERROR(hipMemsetAsync(d.buf(), 0xFF, d.getNumBytes(), stream)); + }; if(!arg.timing) { // fast_check launches and checks each solution fast_check_repeat times, solution by @@ -6001,6 +6009,8 @@ void testing_matmul_with_bias(const Arguments& arg, CHECK_HIP_ERROR( fast_check_fill_sentinel_device((*dDp)[i].buf(), To, size_D[i], stream)); } + else if(checksResults && !arg.c_equal_d) + poisonD(); if(fcCheckE) for(int i = 0; i < gemm_count; i++) CHECK_HIP_ERROR( @@ -6551,7 +6561,7 @@ void testing_matmul_with_bias(const Arguments& arg, // Every path prepares a candidate with this first. auto prepareCandidate = [&](size_t) { - if((arg.unit_check || arg.norm_check || arg.allclose_check) && arg.c_equal_d) + if(checksResults && arg.c_equal_d) { if(batchMode == HIPBLASLT_BATCH_MODE_POINTER_ARRAY) //For General Batch GEMM { @@ -6568,6 +6578,8 @@ void testing_matmul_with_bias(const Arguments& arg, } } } + else if(checksResults) + poisonD(); perf_monitor = EfficiencyMonitor::create(); }; auto skipped = [&](size_t sol, double warm_time, double best_warm_time) { From 59d192a7a66799840d2050a2cff5fd92e9e91345 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Thu, 29 Oct 2026 11:34:41 +0000 Subject: [PATCH 2/3] test(hipblaslt): self-test that the checks of D report a solution that never launches Add skip_launch_solution, a test argument that leaves one solution in the list unlaunched, and a pre_checkin case that sets it for f32 and bf16 in the hipblaslt-test loop and in the hipblaslt-bench candidate search (timing 1). The case passes only if the allclose check reports that solution and no other, so it fails if a solution is again checked against the previous solution's D. It uses allclose alone because unit_check and norm_check fail a case outright on the NaN that the skipped solution leaves in D. Test names of cases that run the candidate search now end in _TIMING. JIRA ID: ROCM-32357 Co-authored-by: Cursor --- .../common/include/hipblaslt_arguments.hpp | 4 + .../clients/common/include/testing_matmul.hpp | 96 ++++++++++++++++++- .../common/src/hipblaslt_arguments.cpp | 1 + .../clients/tests/data/hipblaslt_common.yaml | 2 + .../clients/tests/data/matmul_gtest.yaml | 32 +++++++ .../clients/tests/src/matmul_gtest.cpp | 2 + 6 files changed, 136 insertions(+), 1 deletion(-) diff --git a/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp b/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp index 2f05e19fe925..fdadebe410e3 100644 --- a/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp +++ b/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp @@ -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 \ diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index 3dab7c7de966..978db4179674 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -1235,6 +1235,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, @@ -2484,6 +2513,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]) @@ -5406,6 +5460,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"); @@ -5899,6 +5967,7 @@ void testing_matmul_with_bias(const Arguments& arg, for(auto& d : *dDp) CHECK_HIP_ERROR(hipMemsetAsync(d.buf(), 0xFF, d.getNumBytes(), stream)); }; + bool skippedSolutionChecked = false; if(!arg.timing) { // fast_check launches and checks each solution fast_check_repeat times, solution by @@ -6092,7 +6161,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, @@ -6239,6 +6308,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) { @@ -6720,6 +6795,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()); } #define argument_param \ @@ -6921,6 +7002,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(i % block_count); auto ptr_matmul = matmul[b][0]; return hipblasLtMatmul( @@ -7147,6 +7230,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)); diff --git a/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp b/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp index aac4971724a4..6c23e1ea8501 100644 --- a/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp +++ b/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp @@ -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; diff --git a/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml b/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml index 43d74b8310b8..577c958ad7d3 100644 --- a/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml +++ b/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml @@ -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 diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index f4e4e370ff34..f63d3aa7b64e 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3339,6 +3339,38 @@ 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 all ones before each solution, 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. # diff --git a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp index 4ecac9adf27f..ad8ec55d6c83 100644 --- a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp @@ -462,6 +462,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]) From 749d4b9d6ed07354073d0ce9b472cda1c0e36ecb Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Sat, 31 Oct 2026 08:46:51 +0000 Subject: [PATCH 3/3] test(hipblaslt): describe develop's fill of D in the self-test section The fill of D that the self-test checks now comes from fillDWithSentinel: the fast_check NaN sentinel, with all-ones bytes for D types the sentinel does not cover. JIRA ID: ROCM-32357 Co-authored-by: Cursor --- projects/hipblaslt/clients/tests/data/matmul_gtest.yaml | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index aeb29f412003..54b676f9a5f5 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3363,7 +3363,8 @@ Tests: # # 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 all ones before each solution, so a solution that never +# 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. ###############################################################################