Repository navigation
fix(hipblaslt): clamp workspace limits above 4 GiB instead of rejecting - #13410
Merged
tony-davis merged 1 commit intoOct 10, 2026
Merged
Conversation
#12856 made any workspace limit or size above UINT32_MAX return HIPBLAS_STATUS_INVALID_VALUE. Callers pass SIZE_MAX to mean "no limit", so driftcheck now fails at hipblasLtMatmulPreferenceSetAttribute before any GEMM runs. A limit is a ceiling, so clamp it to UINT32_MAX at the same four entry points: the MAX_WORKSPACE_BYTES preference, hipblasLtMatmul, and the extension API's algoGetHeuristic and initialize. Kernels still never see more than a 32-bit range, and the int truncation fix from #12856 stays. The clamp is logged at the hints level. The reject test becomes two clamp tests. One checks the preference and needs no solution library. The other compares heuristic results at 4 GiB and SIZE_MAX against UINT32_MAX, and runs hipblasLtMatmul and extension initialize with an oversized workspace size. JIRA ID : ROCM-32535 Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
changed the title
tony-davis
marked this pull request as ready for review
October 9, 2026 19:21
✅ 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 9, 2026
newling
left a comment
newling
left a comment
Contributor
There was a problem hiding this comment.
LGTM. Clamping makes sense.
Alex-Vasile
approved these changes
Oct 9, 2026
Alex-Vasile
left a comment
Alex-Vasile
left a comment
Contributor
There was a problem hiding this comment.
Sounds reasonable. Code looks good.
tony-davis
enabled auto-merge (squash)
October 9, 2026 21:54
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status 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 #13410 +/- ##
========================================
Coverage 71.86% 71.86%
========================================
Files 2877 2877
Lines 475878 475860 -18
Branches 70414 70410 -4
========================================
- Hits 341980 341973 -7
+ Misses 110042 110035 -7
+ Partials 23856 23852 -4
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
tony-davis
deleted the
users/tony-davis/ROCM-32535-clamp-workspace-limit
branch
October 10, 2026 01:58
Contributor
TheRock Submodule Bump ActivityNewest first
|
pghysels
added a commit
that referenced
this pull request
Oct 13, 2026
tony-davis
added a commit
that referenced
this pull request
Oct 14, 2026
… 26.13 pin) Refreshes projects/hipblaslt, shared/stinkytofu, shared/origami, shared/mxdatagenerator and projects/hipsparselt to their contents at develop 01f6f9c, the rocm-libraries commit that TheRock's 26.13 BKC branch (release/bkc/therock-10.2-20261013) pins. Everything else stays at release/bkc/therock-10.2-20260922 (201f54b), the rocm-libraries in BKC 26.12. Adds #13410 (ROCM-32535, workspace clamp) on top of the previous patch commit (#12170, #13011, #13274). hipBLASLt in this build is now identical to BKC 26.13's. Co-authored-by: Cursor <cursoragent@cursor.com>
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-32535
Motivation
#12856 fixed a real bug where the C++ heuristic truncated the workspace limit to
int. It also made any workspace limit or size aboveUINT32_MAXreturnHIPBLAS_STATUS_INVALID_VALUE. The thing is, callers commonly passSIZE_MAXto mean "no limit." driftcheck 1.0.6 does exactly that, so it now fails athipblasLtMatmulPreferenceSetAttributebefore any GEMM runs.A limit is a ceiling, so I think a value above what the kernels can address should be clamped, not rejected.
Technical Details
validateWorkspaceSize()becomesclampWorkspaceSize()inrocblaslt_mat_utils.hpp. It returnsmin(value, UINT32_MAX)and logs at the hints level when it clamps (the rocblaslt logger has no warning level).HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTESpreference,hipblasLtMatmul, and the extension API'salgoGetHeuristicandinitialize. Kernels still never see more than a 32-bit range.UINT32_MAX. Reading it back after settingSIZE_MAXgivesUINT32_MAX.inttosize_tfix ingetBestSolutionsfrom fix(hipblaslt): check workspace limits and name complex types #12856 stays.Test Plan
smoke_RejectsWorkspaceAboveUint32Maxlocked in the behavior this PR fixes, so two tests inheuristic_workspace_gtest.cppreplace it:smoke_PreferenceClampsLimitAboveUint32Max: at 4 GiB andUINT64_MAX, setting the preference succeeds and reads backUINT32_MAX. It needs no solution library, so it can't skip.smoke_ClampsWorkspaceAboveUint32Max: at 4 GiB andUINT64_MAX, the C and C++ heuristics return the same results as atUINT32_MAX.hipblasLtMatmulsucceeds with an oversized workspace size (the driftcheck pattern), and extensioninitializesucceeds with an oversizedsetMaxWorkspaceBytes.Restoring the reject, or dropping the clamp at any one entry point, fails a specific assertion.
smoke_CppQueryKeepsLimitsFrom2GiBToUint32Maxis unchanged and still guards theintfix. Both new tests use thesmoke_prefix, so they run in the existing CI smoke lane.Adjacent tests:
AlgoWorkspace.*(the #13184 undersized-workspace test on the same path) and the full*smoke*:*aux_test*set.hipblaslt-test --gtest_filter='HeuristicWorkspaceLimit.*:DataTypeNames.*:AlgoWorkspace.*', with this PR's library and again with adeveloplibraryhipblaslt-test --gtest_filter='*smoke*:*aux_test*'on this PR and ondevelopTest Result
All runs on one MI300X (gfx942), ROCm 10.2.0a20260928, Release build for gfx942,
developate56d6561886.developlibrary fails both new tests; the setter returns 3 (INVALID_VALUE) at 4 GiB andSIZE_MAX.*smoke*:*aux_test*: 3236/3236 pass with this PR, 3235/3235 ondevelop. The only difference in the passing lists is the reject test swapped for the two clamp tests.Device coverage: host-only, with no change to kernels, kernel selection, or defaults for limits at or below
UINT32_MAX. So I think standard PR CI on any gfx lane is enough, and no multi-arch sweep is needed.Submission Checklist
Risk level
Low (2/5). It's host-only and narrow. It turns an error back into the ceiling behavior callers expect, and kernels still see at most 32 bits. One visible difference from before #12856: reading back the preference after setting a huge value returns
UINT32_MAXinstead of the value set.Related: #12856 (AIHPBLAS-4984), the regression source. #13184 (ROCM-32226) touches the same path.