Fix naming bug that would incorrectly discard kernels as duplicates#8745
Open
AlexBrownAMD wants to merge 3 commits into
Open
Fix naming bug that would incorrectly discard kernels as duplicates#8745AlexBrownAMD wants to merge 3 commits into
AlexBrownAMD wants to merge 3 commits into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
Contributor
|
Is this the right approach? |
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 #8745 +/- ##
========================================
Coverage 71.50% 71.50%
========================================
Files 2612 2612
Lines 407895 407895
Branches 60982 60982
========================================
Hits 291632 291632
Misses 94879 94879
Partials 21384 21384
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Contributor
Author
There are 2 xcc mapping algorithms: regular xcc mapping and the new chunking algorithm. For now I think both are needed since we have many kernels tuned using both. |
aliry95amd
approved these changes
Jun 24, 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
4 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
Fixes a bug that could cause kernels to hang or produce invalid results.
Technical Details
WGMXCC value was ignored in the kernel name (flagged as an internal arg). But, WGMXCC=-1 generates different assembly code from WGMXCC set to a fixed value. One version generates code to handle chunking, the other is regular xcc mapping.
If WGMXCC is left out of the name, kernels with different values are discarded as duplicates, even though they contain different assembly code. This can lead to errors like kernel hangs or validation errors if a kernel with regular xcc mapping is called with arguments meant for the chunking algorithm.
This change includes WGMXCC in kernel name as either WGMXCCn1 for -1 (chunk) or WGMXCC1 for any other fixed value since they would produce the same kernel code.
Test Plan
Tested locally and it fixes the kernel hang that was discovered. Running CI to verify other cases.