Repository navigation
fix(tensilelite): correct gfx1250 gate residual instruction size - #12235
Merged
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ 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. |
alex391a
force-pushed
the
users/awang105/fix-gfx1250-gate-inline-literal
branch
from
September 17, 2026 23:46
098c6f7 to
9d3e434
Compare
alex391a
force-pushed
the
users/awang105/fix-gfx1250-gate-inline-literal
branch
from
September 18, 2026 02:16
9d3e434 to
a40b130
Compare
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #12235 +/- ##
===========================================
- Coverage 70.80% 70.80% -0.00%
===========================================
Files 2845 2845
Lines 465630 465630
Branches 68575 68575
===========================================
- Hits 329660 329659 -1
Misses 112251 112251
- Partials 23719 23720 +1
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
vin-huang
approved these changes
Sep 18, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task
1 task done
tony-davis
pushed a commit
that referenced
this pull request
Oct 23, 2026
JIRA ID : AIHPBLAS-4764 ## Motivation `sparse/gfx1250/spmm_gate.yaml` doesn't run on the gfx1250 hardware lane. #12235 fixed its failure, and #12355 removed its xfail after it passed. Then #11777 marked the config `skip-gfx1250-strict` ("not supported yet"). ## Technical Details - Add `sparse/gfx1250-strict/spmm_gate.yaml`, a copy of the gfx1250 config following the pattern of the other 37 sparse strict copies: - `skip-gfx1250` instead of `skip-gfx1250-strict`. - The gfx1250 config is unchanged and keeps running where it runs today. ## Test Plan - This PR's `gfx1250 HW hipBLASLt + TensileLite Tests` run collects and passes `sparse/gfx1250-strict/spmm_gate.yaml`. ## Test Result - New strict test [passes](https://github.com/ROCm/rocm-libraries/actions/runs/39722965803/job/119640530321?pr=14175) on gfx1250 HW run. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. One new test config that runs only on gfx1250-strict. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude <noreply@anthropic.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 : AIHPSPAR-277
Motivation
The gate residual path selects
1.0when the gate pointer is null. It represented this value as the integer literal0x3f800000, which StinkyTofu counted as an additional 4-byte literal while LLVM encoded it as an inline floating-point constant. This caused the gfx1250 assembly-size check to report a 4-byte mismatch before client execution.Technical Details
Change the gate residual select operand from
hex(0x3f800000)to the floating-point literal1.0. The generated machine instruction remains unchanged, while StinkyTofu and LLVM now agree on its 4-byte size.Enable
CheckASMCodeSizeinspmm_gate.yamlso the gfx1250 FFM test checks the generated instruction size and detects future mismatches in this path.Test Plan
Test Result
Submission Checklist