WIP: feat(tensilelite): Use StinkyTofu with Subtile.#9303
Draft
memmett wants to merge 8 commits into
Draft
Conversation
…riter Move rocIsaPass and StinkyTofu conversion logic from kernelBody() into reusable helpers with no behavior change for classic or subtile kernels. Co-authored-by: Cursor <cursoragent@cursor.com>
…ubtile Set _StinkySubtile on gfx1250 subtile solutions and run rocIsaPass plus a minimal StinkyTofu pipeline in kernelBodySubtile. Add unit tests. Co-authored-by: Cursor <cursoragent@cursor.com>
Emit PRELOOP/MAINLOOP/NGLL as loopBody and NLL as noLoadLoopBody so gfx1250 waitcnt region detection works without StinkyTofu converter changes. Co-authored-by: Cursor <cursoragent@cursor.com>
…itcnt Attach LDS buffer tokens to GR/LR loads and barriers so gfx1250 subtile kernels can pass MemTokenConsistencyCheck in the waitcnt-only pipeline. Co-authored-by: Cursor <cursoragent@cursor.com>
Expose loopBody as a top-level kernel-body module so ST can detect loopWithPrefetch, enable loop-carried token deps, and skip hand-rolled vmcnt adjustment when the waitcnt-only pipeline owns synchronization. Co-authored-by: Cursor <cursoragent@cursor.com>
Expose the tail loop as top-level noLoadLoopBody modules so Gfx1250 waitcnt insertion covers TAILLOOP and replaces hand-rolled tail waits. Co-authored-by: Cursor <cursoragent@cursor.com>
…n passes Centralize waitcnt-only StinkyTofu overrides and add asm checks that ST does not double-insert cluster barriers or wait-alu while OptLevel 0 skips DAG rescheduling. Co-authored-by: Cursor <cursoragent@cursor.com>
…is on Let StinkyTofu own GR/LR waitcnt insertion for gfx1250 subtile kernels by no-oping WaitGROp/WaitLROp emission while keeping scheduler deps. Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (76.92%) 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 #9303 +/- ##
===========================================
+ Coverage 64.99% 65.00% +0.01%
===========================================
Files 2685 2685
Lines 422183 422070 -113
Branches 62726 62719 -7
===========================================
- Hits 274365 274339 -26
+ Misses 127113 127026 -87
Partials 20705 20705
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
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
2 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.
No description provided.