Repository navigation
test(hipblaslt): self-test that the checks of D report a solution that never launches - #14647
Conversation
…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 <cursoragent@cursor.com>
…t 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 <cursoragent@cursor.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (68.40%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #14647 +/- ##
===========================================
- Coverage 77.35% 77.29% -0.07%
===========================================
Files 2894 2893 -1
Lines 487162 486074 -1088
Branches 73731 73490 -241
===========================================
- Hits 376823 375663 -1160
- Misses 85448 85637 +189
+ Partials 24891 24774 -117
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
tony-davis
left a comment
tony-davis
left a comment
There was a problem hiding this comment.
Hi Jinchen, this looks good to me, and thanks for the really clear write-up! I like that the self-test goes red with the fill disabled (your table shows all 4 cases failing), and I checked that all four matmul_check_reports_unlaunched_solution cases actually ran and passed (not skipped) in TheRock's gfx942 and gfx110X hipBLASLt lanes. I also walked the bench path. The fill happens in prepare, before the first cold launch that the check copies back, and it lands outside both the skip-slow warm-up timer and the timed loop, so I don't think it moves any timing numbers. A couple of non-blocking thoughts:
- I bet the nightly
comprehensivetier, which checks a lot more multi-solution cases thanstandard, may turn up a few newly red cases where a kernel skips its D store. That's the fix doing its job, but it's probably worth a heads-up to whoever watches nightly. - Tiny nit: "all ones" in the title reads like the value 1.0 (which a skipped solution could match when alpha=0, beta=1 and C is all ones). The code fills all-ones bytes, NaN in the IEEE types, so there's no real problem. "all-ones bytes" in the title would just match the description and keep the squash commit unambiguous.
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 <cursoragent@cursor.com>
jinchen62
changed the title
|
Thanks Tony. While merging develop I found #14743, which landed an equivalent fill of D on develop (
|
TheRock Submodule Bump ActivityNewest first
|
JIRA ID : ROCM-32357
Motivation
hipblaslt-benchandhipblaslt-testshare one test harness,testing_matmul.hpp. It can run several solutions (kernel configurations) for one matrix multiply and compare each result, the output matrix D, with a CPU reference. All solutions write into the same D buffer, so a solution that never stored D used to be checked against the previous solution's result and pass. On gfx942 before #12797,hipblaslt-bench --algo_method all -vpassed three hand-writtenCustomGSUs_*kernels this way. #14743 now fills D with a NaN sentinel before each solution, but no test fails if that fill stops working.Technical Details
skip_launch_solution, leaves the solution at that position in the list unlaunched (C interface with strided batches only), so that solution never writes D. After each solution's check, the harness fails the case if the allclose check reported a solution other than the unlaunched one or missed the unlaunched one, and also if the unlaunched solution was never checked.pre_checkin) casematmul_check_reports_unlaunched_solutionsets it for f32 and bf16 at 131×67×200, in thehipblaslt-testloop and in the bench's loop over candidate solutions (timing: 1). It uses allclose alone becauseunit_checkandnorm_checkfail a case outright when D holds NaN. Cases withtiming: 1now get a_TIMINGname suffix.Arguments, the clients' parameter struct, gains oneint32_tfield, added at the same position inhipblaslt_common.yaml, from whichhipblaslt_gtest.datais generated at build time.fillDWithSentinel, so this pull request now adds only the self-test.Test Plan
On MI300X (gfx942), with hipBLASLt built from this branch after the merge of develop and
HIPBLASLT_LIBLOGIC_PATHlimited to six logic files: fp8 × fp16 → f16, f16 → f32, grouped GEMM, aux output, and the f32 and bf16 N/T grid-based files cut to the 8 solutions nearest 131×67×200.matmul_check_reports_unlaunched_solutionwith develop's fill and with the fill disabled, plus the existingfast_check_injectself-test.-v, to check that the fill passes correct kernels: fp8 × fp16 at 4096×16×8192 with--verify_strict,--c_equal_d,--rotating 256with timed iterations,--api_method mixand--HMM; f16 → f32 at 4096×16×2048 with the C and C++ interfaces; f16 → f32 with pointer-array batches; grouped GEMM; aux output.Test Result
matmul_check_reports_unlaunched_solution, 4 cases--verify_strictexits 0), 39 f16 → f32 solutions with each interface, 37 with pointer-array batches, and the single grouped-GEMM and aux-output solutions. Thefast_check_injectself-test passes.Not addressed here
unit_checkandnorm_check, the harness still does not reset E (the aux output), the bias gradient or amaxD between solutions; onlyfast_checkfills them. ROCM-33792 tracks this.Submission Checklist