fix(tensilelite): Fix output-store 32-bit address overflow#8801
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (77.89%) 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 #8801 +/- ##
===========================================
- Coverage 64.58% 64.58% -0.00%
===========================================
Files 2675 2658 -17
Lines 420244 415455 -4789
Branches 62446 61689 -757
===========================================
- Hits 271410 268304 -3106
+ Misses 128348 126984 -1364
+ Partials 20486 20167 -319
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
✅ 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. |
jimbochi
changed the title
jimbochi
changed the title
Contributor
There was a problem hiding this comment.
Pull request overview
This PR fixes a TensileLite conversion-kernel bug where large output tensors could trigger 32-bit address/offset truncation in buffer_store, causing silent writes to the wrong address. It does this by widening the output byte-offset computation and plumbing a 64-bit offset path through the buffer_store helpers.
Changes:
- Added
splitBufferOffset()plusuint64_t voffsetoverloads forbuffer_storespecializations (1/2/4/8/16B) to support >4GB byte offsets via base-pointer adjustment. - Updated
KernelWriterConversion.pyto computebyteOffsetDas a 64-bit value and pass it intobuffer_store. - Added unit tests validating the codegen/header patterns, and updated
known_bugs.yamlcommentary/entries for overflow tracking.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 2 comments.
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Source/memory_gfx.h | Adds splitBufferOffset() and buffer_store uint64 overloads to safely handle large byte offsets. |
| projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py | Switches conversion-kernel store offset computation to byteOffsetD and passes it to buffer_store. |
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py | Adds unit tests to guard the new 64-bit offset path and helper/overload presence. |
| projects/hipblaslt/clients/tests/data/known_bugs.yaml | Removes obsolete entries and adds an overflow-coverage tracking section/template. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
jichangjichang
approved these changes
Jul 9, 2026
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.
Motivation
Conversion kernels compute the output byte-offset as
idxD * sizeof(destType)in32-bit arithmetic. When the output tensor's leading dimension is large enough that
an element's byte-offset exceeds 2^31, the 32-bit multiplication wraps and the
buffer_storeintrinsic writes to a wrong address, producing silent data corruption.This PR widens the output-store offset path to 64-bit, fixing the address overflow
for large-output matmul cases.
JIRA ID : ROCM-26455
Technical Details
memory_gfx.h— Added asplitBufferOffset()helper that splits a 64-bitvoffsetinto a low-32-bit portion (passed as the hardwarevoffset) and foldsthe high-32 bits into the buffer base pointer. Each
buffer_storespecialisation(1B / 2B / 4B / 8B / 16B) gains a new constructor overload accepting
uint64_t voffsetthat calls this helper before issuing the intrinsic.KernelWriterConversion.py— The codegen now pre-computesuint64_t byteOffsetD = idxD * sizeof(destType)and passes it tobuffer_storeinstead of the inline
idxD * sizeof(...)expression, ensuring the multiplicationis performed in 64-bit.
known_bugs.yaml— Removed the obsolete gfx950 MXFP4 subtile scaleAlphaVecentries (ROCM-27082, resolved). Added a size_t / 32-bit overflow tracking section
(AIHPBLAS-3806) with a template and the ROCM-26455 bf16 output-overflow reproducer
entry (currently commented out, ready to activate when the matching gtest lands).
Test Plan
test_KernelWriterConversion_64bit_offset.py
Test Result
python3 -m pytest Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py -v
=====================test session starts=====================
platform linux -- Python 3.12.3, pytest-8.4.2, pluggy-1.6.0 -- /opt/venv/bin/python3
cachedir: .pytest_cache
rootdir: /src/Issues/projects/hipblaslt/tensilelite
configfile: pytest.ini
plugins: hydra-core-1.3.2
collected 8 items
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestKernelWriterConversion64BitOffset::test_byteOffsetD_declared_as_uint64 PASSED [ 12%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestKernelWriterConversion64BitOffset::test_buffer_store_uses_byteOffsetD PASSED [ 25%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestKernelWriterConversion64BitOffset::test_no_inline_idxD_multiply_in_buffer_store PASSED [ 37%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestMemoryGfxSplitBufferOffset::test_splitBufferOffset_exists PASSED [ 50%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestMemoryGfxSplitBufferOffset::test_splitBufferOffset_signature PASSED [ 62%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestMemoryGfxSplitBufferOffset::test_splitBufferOffset_masks_high_bits PASSED [ 75%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestMemoryGfxSplitBufferOffset::test_buffer_store_uint64_overloads_exist PASSED [ 87%]
Tensile/Tests/unit/test_KernelWriterConversion_64bit_offset.py::TestMemoryGfxSplitBufferOffset::test_overloads_call_splitBufferOffset PASSED [100%]
===================== 8 passed in 0.06s=====================
Submission Checklist