fix(hipblaslt): Fix gfx1201 test for FP16 T-T GEMM + bias-fused epilogue#9204
Conversation
✅ 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. |
jinchen62
changed the title
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (76.84%) 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 #9204 +/- ##
========================================
Coverage 65.36% 65.36%
========================================
Files 2685 2685
Lines 422269 422269
Branches 62742 62742
========================================
Hits 276012 276012
Misses 125466 125466
Partials 20791 20791
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
newling
left a comment
newling
left a comment
There was a problem hiding this comment.
LGTM. My only suggestion is to trim the test down to a minimal case.
tony-davis
left a comment
tony-davis
left a comment
There was a problem hiding this comment.
Approving. Nice work, and thank you for the diligence here (tracking it down to the one BiasDataTypeList that was out of step with all its siblings is a clean root cause). I especially appreciate the before/after evidence on the regression test: returnedAlgoCount 0 -> 8 plus the numeric correctness numbers is exactly the kind of proof that makes a fix easy to trust.
One thing I want to name explicitly (and then accept): the new matmul_bias_fp16_gfx120x_tt test is gated to gpu_arch: '120[0-1]', and we don't currently have a gfx120x lane in this PR's CI, so it won't actually execute in shared CI. The validation we have is on real gfx1201 hardware, run locally. I'm accepting that for this merge (it's solid evidence, and it's the best we can do given the fleet today).
The thing is, this PR is a pretty good argument for why we need gfx120x GPU runners in CI. This exact bug would have been caught up front, and just as importantly, this regression test would keep running on every change instead of sitting dormant until someone happens to test on the hardware. I'm going to point at this one as more evidence for standing up gfx120x coverage.
Thanks again, this is good work.
4 participants
Technical Details
Root cause:
BiasDataTypeListinTensile/Logic/asm_full/gfx1201/GridBased/gfx1201_Cijk_Alik_Bjlk_HHS_BH_Bias_SHB_HA_S_SAB_SCD_SAV_UserArgs.yaml(the FP16 T-T logic file on gfx1201) was[0](Float only) instead of[0, 4](Float and Half). This is the only transpose combo on gfx1201 missing Half — its three sibling files (N-N, T-N, N-T) and the identical file on gfx1200 all correctly list[0, 4]. SinceBiasDataTypeListdrives which bias-dtype code pathsKernelWriterAssembly.pygenerates into the kernel, the compiled kernels for this one config genuinely had no Half-bias path.Fix: widen
BiasDataTypeListto[0, 4]to match every sibling combo, and rebuild so the kernel is regenerated with Half-bias support.Also added a regression test (
matmul_bias_fp16_gfx120x_ttinclients/tests/data/matmul_gtest.yaml, gated togpu_arch: '120[0-1]').Test Result
returnedAlgoCount == 0for FP16 T-T + BIAS/RELU_BIAS/GELU_BIAS.returnedAlgoCount == 8for all three, on real gfx1201 hardware.maxAbsErr=0.000977 maxRelErr=0.000479, 0/3072elements out of tolerance vs. CPU reference.
maxAbsErr=0.000974 maxRelErr=0.000472,0/3072 elements out of tolerance vs. CPU reference.
JIRA ID: ROCM-25598