fix(hipblaslt): drop stale KernArgsVersion 2 overrides from gfx950 logic - #12017
Open
jaopaulolc wants to merge 1 commit into
Open
fix(hipblaslt): drop stale KernArgsVersion 2 overrides from gfx950 logic#12017jaopaulolc wants to merge 1 commit into
jaopaulolc wants to merge 1 commit into
Conversation
467 solution entries across 9 gfx950 logic files still pin InternalSupportParams.KernArgsVersion to 2. The host then packs kernel arguments in v2 order (D, C, A, B, ...) while the generated kernel reads them in v3 order (A, B, strides, ...), so the kernel dereferences D as A and faults on its first global load. Remove the override so the affected entries inherit the default version 3. Only the KernArgsVersion token is removed; no solution entry is added, removed, or otherwise altered.
✅ 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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #12017 +/- ##
===========================================
- Coverage 70.23% 70.23% -0.00%
===========================================
Files 2810 2810
Lines 462535 462535
Branches 68105 68105
===========================================
- Hits 324828 324821 -7
- Misses 114187 114191 +4
- Partials 23520 23523 +3
*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
1 participant
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 : N/A
Motivation
467 solution entries across 9 gfx950 logic files still pin
InternalSupportParams.KernArgsVersionto 2, while the rest of the gfx950 logichas moved to the default version 3. When such an entry is selected, the host packs
kernel arguments in v2 order (
D, C, A, B, ...) but the generated kernel reads themin v3 order (
A, B, strides, ...); the kernel dereferencesDasAand faults onits first global load.
This is a pre-existing defect on
developthat causesHSA_STATUS_ERROR_MEMORY_FAULTon gfx950, and it blocks other gfx950 work by making any exhaustive solution run abort.
It also causes the current gfx950
precheckinfailure ondevelop: test_/matmul_test.matmul/pre_checkin_matmul_128_f32_rf32_rf32_rf32_rf32_r_NT_128_128_128_5_128_128_0_128_128_1faults in kernel
Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_..._WG32_8_1,which is one of the stale entries in
gfx950_Cijk_Ailk_Bjlk_S_B_BiasS_HAS_SAV_UserArgs.yaml.PR #11943 contains a 34-entry subset of this change (one file); expect an overlap when
merging.
Technical Details
The
KernArgsVersion: 2token is removed from each affectedInternalSupportParamsmapping so the entry inherits the default version 3.Nothing else changes: per file the diff has equal insertions and deletions, and the
solution count of every file is unchanged.
Logic/asm_full/gfx950/)gfx950/Equality/gfx950_Cijk_Ailk_Bjlk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yamlgfx950/Equality/gfx950_Cijk_Ailk_Bjlk_S_B_BiasS_HAS_SAV_UserArgs.yamlgfx950/Equality/gfx950_Cijk_Ailk_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yamlgfx950/Equality/gfx950_Cijk_Ailk_Bljk_S_B_BiasS_HAS_SAV_UserArgs.yamlgfx950/Equality/gfx950_Cijk_Alik_Bjlk_S_B_BiasS_HAS_SAV_UserArgs.yamlgfx950/Equality/gfx950_Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yamlgfx950/Equality/gfx950_Cijk_Alik_Bljk_S_B_BiasS_HAS_SAV_UserArgs.yamlgfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yamlgfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_S_MX_B_BiasS_HAS_SAV_UserArgs.yamlThese 9 files are the only shipped logic files in the repository that reference
KernArgsVersion; no other architecture is affected. The remaining v2 kernels in abuilt gfx950 library are all hand-written
Custom_*kernels, which are not sourcedfrom logic files and are out of scope here.
Test Plan
gfx950 (MI355X), build from source with
GPU_TARGETS=gfx950, A/B against the sametree with and without this change:
hipblaslt-bench --algo_method indexon entries that carriedthe override, plus unaffected entries from the same files as controls.
hipblaslt-bench --algo_method allover every supported solution forfp32 NN/NT/TN/TT and bf16 NN/NT/TN, checking both the fault count and that the
supported-solution count is unchanged.
precheckingtest above.Test Result
Memory access fault) and all pass after; 7 of 7 control entries pass in both arms.Covers 6 of the 9 files.
after, all 7 pass with exit 0 and no faults. Supported-solution counts identical in
both arms (645, 648, 781, 519, 1124, 802, 2689).
pre_checkin_matmul_128_f32_rf32_rf32_rf32_rf32_r_NT_128_128_128_5_128_128_0_128_128_1:exit 134 with a memory fault before,
[ PASSED ] 1 testafter.0 faults.
The two
gfx950_id75a3files (39 entries) are not built in this configuration and wereverified statically only: zero
KernArgsVersionreferences remain in the logic tree andtheir solution counts are unchanged.
Submission Checklist
Risk level
Low. Data-only change to gfx950 heuristic logic that removes an override and restores
the default already used by the rest of the file. No solution is added or removed, and
solution counts are unchanged.