[hipblaslt] Fix Float4/Float6/BFloat6 elementBytes assertion and totalAllocatedBytes (AIHPBLAS-3570) - #7672
Merged
bnemanich merged 3 commits intoMay 23, 2026
Conversation
…lAllocatedBytes PR #6499 (gfx950 mxfp4) changed BaseTypeInfo::ElementSize from sizeof(T) to sizeof(T) / Packing, so info.elementSize is now the per-element segment size in bytes (0.5 for Float4, 0.75 for Float6/BFloat6). Two sites in TensorDescriptor.hpp were not updated to match: elementBytes() still hard-asserted m_dataType was none of {Float4,Float6,BFloat6} and carried a stale comment claiming elementSize was the unsegmented container size, and totalAllocatedBytes() divided by info.packing a second time on top of multiplyElementSize, returning N/4 bytes for N Float4 elements instead of the correct N/2. Drop the obsolete assertion and the obsolete /packing, replace the misleading comment with one that reflects post-PR-6499 semantics, and keep totalAllocatedBytes safe by asserting the element count is a whole number of packed containers. Add a TensorDescriptor_test.cpp covering Float (control), Half, Float8, Float4, Float6, and BFloat6 in tensilelite-tests, including an explicit Debug-mode regression check that elementBytes() does not abort on the packed MX types and that totalAllocatedBytes() is not double-divided. Non-packed types are unchanged because their packing == 1. Co-authored-by: Cursor <cursoragent@cursor.com>
Match the style used by MXScalePadding_test.cpp and MXDataGen_test.cpp. Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
bnemanich
changed the title
talumbau
approved these changes
May 21, 2026
talumbau
left a comment
talumbau
left a comment
Collaborator
There was a problem hiding this comment.
thanks for the quick fix!
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (77.83%) 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 #7672 +/- ##
===========================================
- Coverage 61.76% 61.73% -0.03%
===========================================
Files 2082 2083 +1
Lines 356042 356515 +473
Branches 53628 53752 +124
===========================================
+ Hits 219883 220067 +184
- Misses 117417 117697 +280
- Partials 18742 18751 +9
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
bnemanich
enabled auto-merge (squash)
May 22, 2026 12:13
bnemanich
disabled auto-merge
May 23, 2026 14:43
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
3 participants
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.
Motivation
PR #6499 (gfx950 mxfp4) changed BaseTypeInfo::ElementSize from sizeof(T)
to sizeof(T) / Packing, so info.elementSize is now the per-element
segment size in bytes (0.5 for Float4, 0.75 for Float6/BFloat6). Two
sites in TensorDescriptor.hpp were not updated to match: elementBytes()
still hard-asserted m_dataType was none of {Float4,Float6,BFloat6} and
carried a stale comment claiming elementSize was the unsegmented
container size, and totalAllocatedBytes() divided by info.packing a
second time on top of multiplyElementSize, returning N/4 bytes for N
Float4 elements instead of the correct N/2.
Drop the obsolete assertion and the obsolete /packing, replace the
misleading comment with one that reflects post-PR-6499 semantics, and
keep totalAllocatedBytes safe by asserting the element count is a whole
number of packed containers. Add a TensorDescriptor_test.cpp covering
Float (control), Half, Float8, Float4, Float6, and BFloat6 in
tensilelite-tests, including an explicit Debug-mode regression check
that elementBytes() does not abort on the packed MX types and that
totalAllocatedBytes() is not double-divided. Non-packed types are
Addresses AIHPBLAS-3570.
Test Plan
Tested on debug build to make sure everything passed.
Submission Checklist