Repository navigation
fix(hipblaslt): refuse the unsupported epilogue aux scale - #13522
Merged
tony-davis merged 3 commits intoOct 14, 2026
Merged
Conversation
No kernel applies an aux scale, so a non-null HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER or GemmInputs::setScaleAux was silently ignored and E came out unscaled. Refuse it with HIPBLAS_STATUS_NOT_SUPPORTED and a log message instead; null stays allowed. Remove the quarantined matmul_fast_check_aux_scale_e_ROCM32138 case, whose premise is gone; AlgoErrors.smoke_AuxScaleIsRefused is the regression test. ROCM-32138 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❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #13522 +/- ##
===========================================
- Coverage 74.64% 71.99% -2.65%
===========================================
Files 2881 2879 -2
Lines 480122 478672 -1450
Branches 71209 70933 -276
===========================================
- Hits 358358 344611 -13747
- Misses 97497 110041 +12544
+ Partials 24267 24020 -247
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
newling
approved these changes
Oct 13, 2026
newling
left a comment
newling
left a comment
Contributor
There was a problem hiding this comment.
LGTM, just one nit about error messages
Resolve known_bugs.yaml and matmul_gtest.yaml: develop already removed the ROCM-32062, ROCM-32139 and ROCM-32277 entries, and this PR removes the ROCM-32138 entry and its quarantined case. Co-authored-by: Cursor <cursoragent@cursor.com>
Include the sizes, ops, types and the pointer in the error, and name the call site (Gemm, GroupedGemm, MatmulDesc) so each message is unique. Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
enabled auto-merge (squash)
October 14, 2026 15:22
tony-davis
deleted the
users/todavis_amdeng/ROCM-32138-refuse-aux-scale
branch
October 14, 2026 23:32
Contributor
TheRock Submodule Bump ActivityNewest first
|
Merged
2 tasks done
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-32138
Targets
develop. It removes the quarantine that the fast_check stack (#13021, AIHPBLAS-5018, now merged) added for this defect.Motivation
HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER(andGemmInputs::setScaleAuxin the extension API) let a caller set a scale for the epilogue's auxiliary output, E. No kernel applies it. The library stored the pointer, passed it down to TensileLite, and returned success, so E came out unscaled and the caller had no way to know. Implementing the scale would mean new kernel arguments and codegen across every epilogue that writes E. I think refusing it is the right fix until someone needs it.Technical Details
hipblasLtMatmulDescSetAttributereturnsHIPBLAS_STATUS_NOT_SUPPORTED, with a log message, for a non-null aux scale pointer. Setting it to null still succeeds. The check is inhipblaslt.cpp, because rocblaslt has no status that maps toNOT_SUPPORTED; the file already returns it directly elsewhere.GemmInputs::setScaleAuxreturnsvoid, soGemm::setProblemandGroupedGemm::setProblemrefuse a non-null aux scale instead, the same way. The shortersetProblemoverloads forward to these two, and the deviceUserArgumentsstruct has no aux scale field, so these are the only ways in.hipblaslt.handhipblaslt-ext.hppnow say the aux scale is not supported and what setting it returns.matmul_fast_check_aux_scale_e_ROCM32138comes out ofmatmul_gtest.yamlandknown_bugs.yaml. It set the aux scale, which now fails at the set call, so its premise is gone. It was the only YAML case withscaleE: 1, and no bench flag sets it. The client's scaleE plumbing intesting_matmul.hppstays, unused.AlgoErrors.smoke_AuxScaleIsRefused. It checks that the C API refuses a non-null pointer and accepts null, and that both extensionsetProblemcalls refuse one.aux_matmul_set_get_attrexpected success for a non-null pointer; it now expectsNOT_SUPPORTED, and success for null.Any caller that sets this attribute, even to a pointer to 1.0, now gets an error where it used to get silently unscaled E. That's on purpose. A follow-up could implement the scale if a user needs it.
Test Plan
On gfx942 and gfx950:
once with this branch's library, and once with the same test binary against a
developlibrary without this change.Test Result
ROCm 10.2.0 (TheRock 10.2.0a20260928), built from this branch for gfx942 and gfx950:
pre_checkinaux case. gfx950 (MI355X): the same filter runs 209 tests, and all pass.developlibrary at 24bdd07), on both: 2 tests fail. The smoke test's C-API set,Gemm::setProblemandGroupedGemm::setProblemeach returnHIPBLAS_STATUS_SUCCESSwhereHIPBLAS_STATUS_NOT_SUPPORTEDis expected, andpre_checkin_aux_matmul_set_get_attr_f16_rfails at the same set call.Submission Checklist
Risk level
Low (2/5): host-side checks on one attribute and one extension input. Callers that set the aux scale now get an error, but they never got a scaled E.