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 09e6aa432f0d..65b9cb9e8c54 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -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(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)); 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 e8cc36569366..54b676f9a5f5 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -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. # diff --git a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp index d5454d6095aa..4124d65bce8e 100644 --- a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp @@ -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])