Repository navigation
fix(hipblaslt): lay out E with lde and stride_e in the CPU reference - #13215
Merged
tony-davis merged 2 commits intoOct 13, 2026
Merged
Conversation
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (42.73%) 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 #13215 +/- ##
===========================================
- Coverage 71.99% 71.86% -0.12%
===========================================
Files 2879 2876 -3
Lines 478104 475289 -2815
Branches 70852 70231 -621
===========================================
- Hits 344165 341563 -2602
+ Misses 109972 109948 -24
+ Partials 23967 23778 -189
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
tony-davis
changed the title
tony-davis
force-pushed
the
users/todavis_amdeng/AIHPBLAS-5017-side-operands
branch
2 times, most recently
from
October 7, 2026 20:03
dccd0c0 to
a007b01
Compare
tony-davis
force-pushed
the
users/todavis_amdeng/ROCM-32139-cpu-ref-e-lde
branch
from
October 7, 2026 20:24
392af76 to
13305ea
Compare
tony-davis
force-pushed
the
users/todavis_amdeng/AIHPBLAS-5017-side-operands
branch
from
October 7, 2026 21:05
a007b01 to
519e7a1
Compare
tony-davis
force-pushed
the
users/todavis_amdeng/ROCM-32139-cpu-ref-e-lde
branch
from
October 7, 2026 21:24
13305ea to
3030497
Compare
tony-davis
force-pushed
the
users/todavis_amdeng/AIHPBLAS-5017-side-operands
branch
2 times, most recently
from
October 8, 2026 17:03
1e40a94 to
aa4b65c
Compare
tony-davis
force-pushed
the
users/todavis_amdeng/ROCM-32139-cpu-ref-e-lde
branch
from
October 8, 2026 18:08
3030497 to
3e6b887
Compare
An error occurred while trying to automatically change base from
users/todavis_amdeng/AIHPBLAS-5017-side-operands
to
users/todavis_amdeng/AIHPBLAS-5016-exactness-large-k
October 9, 2026 03:32
… reference The client CPU reference's epilogue (epilogue_func in testing_matmul.hpp) wrote the aux output E, and read it back for gradient epilogues, at D's position (j * ldd + i), and started each batch's E at D's batch offset (stride_d). The library lays E out with lde and stride_e, so whenever lde differs from ldd, unit_check failed a correct GPU result. Pass lde to epilogue_func and index E with it, and offset each batch's E by stride_e. matmul_fast_check_aux_padded_ldd_ROCM32139 comes out of known_bugs.yaml and becomes the regression test. It now also runs with two batches, which the stride_e change needs, and with a padded lde, which its fast_check-only copy covered while it was quarantined; that copy is removed. Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
force-pushed
the
users/todavis_amdeng/ROCM-32139-cpu-ref-e-lde
branch
from
October 9, 2026 16:16
3e6b887 to
b56f244
Compare
tony-davis
changed the base branch from
users/todavis_amdeng/AIHPBLAS-5017-side-operands
to
develop
October 9, 2026 16:16
✅ 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. |
newling
approved these changes
Oct 12, 2026
newling
left a comment
newling
left a comment
Contributor
There was a problem hiding this comment.
Looks good to me. I left one suggestion about argument names that might be invalid, or better done as a NFC follow-up.
tony-davis
added a commit
that referenced
this pull request
Oct 13, 2026
Resolve the matmul_gtest.yaml conflict with #13215. Each PR removed its own fast_check-only copies from the shared block, so with both fixes in, the whole block comes out. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : ROCM-32139
Targets
develop. It removes the quarantine that the fast_check stack (#13015, AIHPBLAS-5017, now merged) added for these cases.Motivation
hipblaslt-testchecks the GPU's answer against a reference computed on the CPU. For epilogues with an auxiliary output, E (the pre-activation result some training workloads save for the backward pass), the reference writes its own copy of E, and gradient epilogues read E back.The library lays E out with its own leading dimension (
lde) and batch stride (stride_e). The reference used D's instead: it wrote E at D's position,j * ldd + i, and started each batch's E at D's batch offset. Wheneverldediffers fromldd,unit_checkcompared the GPU's correct E against a reference laid out on the wrong grid, and failed it. This affects the test client only; the library is not involved.Technical Details
epilogue_func(inclients/common/include/testing_matmul.hpp) takesldeand indexes E withj * lde + i, for both the write and the gradient read. D still usesldd.stride_einstead ofstride_d.matmul_fast_check_aux_padded_ldd_ROCM32139comes out ofknown_bugs.yamland becomes the regression test. It now also runs with two batches, which thestride_echange needs, and with a paddedlde(140, where M is 131 andlddis 151). The padded-ldesize was covered by a fast_check-only copy while the case was quarantined; that copy is removed.Test Plan
On gfx942 and gfx950, run the formerly quarantined case and every
pre_checkincase with an aux output:once with this branch's client, once with the client rebuilt without the
testing_matmul.hppchange, and once with only thestride_echange reverted.Test Result
ROCm 10.2.0 (TheRock 10.2.0a20260928), built from this branch:
lde131 and 140, one batch and two).stride_echange reverted: the 2 two-batch variants fail and the 2 one-batch variants pass, so each half of the fix has a test that fails without it.Submission Checklist
Risk level
Low (1/5): test-client only. The reference changes only when E's leading dimension or batch stride differs from D's, where it was wrong; library and device behavior are unchanged.