Repository navigation
Revert "perf(hipblaslt): remove gfx950 equality custom kernels" - #14886
Merged
tony-davis merged 1 commit intoNov 2, 2026
Merged
tony-davis merged 1 commit into
tony-davis merged 1 commit into
Conversation
This reverts commit 2a155d5.
Contributor
There was a problem hiding this comment.
Note
Copilot was unable to run its full agentic suite in this review.
Copilot review overview
4 open issues
KernelNameMin/SolutionNameMin(...shortname1_gfx950).… · NewKernelNameMin/SolutionNameMin(...shortname1_gfx950).… · NewKernelNameMin/SolutionNameMin(e.g.,… · NewKernelNameMin/SolutionNameMin(e.g.,… · New
What changed in this PR
Reverts a prior ROCm libraries change by re-introducing gfx950 Equality custom-kernel logic entries for hipBLASLt/Tensile, restoring custom solutions and updating exact-logic mappings accordingly.
Changes:
- Re-adds custom kernel solution definitions for several gfx950 Equality problem types (including 256CU and 128CU variants).
- Updates existing Equality logic to insert a custom solution and re-map ExactLogic entries to the new solution indices.
- Removes an older non-custom solution block (likely superseded by the restored custom solution set).
| File | Description |
|---|---|
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_HHS_BH_BiasSH_HAS_SAV_UserArgs.yaml | Inserts a custom solution and shifts subsequent SolutionIndex values; updates ExactLogic mappings. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_F8BS_BH_SAB_UserArgs.yaml | Adds new Equality logic file defining multiple custom solutions for F8BS. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_F8B8BS_BH_SAB_UserArgs.yaml | Adds new Equality logic file defining multiple custom solutions for F8B8BS. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_BBS_BH_UserArgs.yaml | Adds new Equality logic file defining multiple custom solutions for BBS. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950_128cu/Equality/gfx950_Cijk_Alik_Bljk_HHS_BH_BiasSH_HAS_SAV_UserArgs.yaml | Adds a 128-CU-specific Equality logic file with a custom solution and ExactLogic. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950/Equality/gfx950_Cijk_Alik_Bljk_F8BS_BH_SAB_UserArgs.yaml | Adds 256-CU gfx950 Equality logic for F8BS custom kernels. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/gfx950/Equality/gfx950_Cijk_Alik_Bljk_F8B8BS_BH_SAB_UserArgs.yaml | Adds 256-CU gfx950 Equality logic for F8B8BS custom kernels. |
Review effort: Lite
Comment on lines
+412
to
+418
| KernelNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950 | ||
| SolutionNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950 | ||
| ActivationFuncCall: false | ||
| AssertSummationElementMultiple: 64 | ||
| AssignedDerivedParameters: true | ||
| AssignedProblemIndependentDerivedParameters: true | ||
| BaseName: Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950 |
Comment on lines
+570
to
+571
| KernelNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950 | ||
| SolutionNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950 |
Comment on lines
+412
to
+423
| KernelNameMin: Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950 | ||
| SolutionNameMin: Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950 | ||
| ActivationFuncCall: false | ||
| AssertSummationElementMultiple: 128 | ||
| AssignedDerivedParameters: true | ||
| AssignedProblemIndependentDerivedParameters: true | ||
| BaseName: Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950 | ||
| CUCount: null | ||
| CUOccupancy: -1 | ||
| ClusterLocalRead: 0 | ||
| CodeObjectVersion: 4 | ||
| CustomKernelName: Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950 |
| - SolutionIndex: 2 | ||
| SolutionUID: 0uR2gI270f1q | ||
| KernelNameMin: Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950 | ||
| SolutionNameMin: Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950 |
✅ 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. |
tony-davis
deleted the
revert-14281-users/m-kim/remove-gfx950-equality-custom-kernels
branch
November 2, 2026 19:15
Contributor
TheRock Submodule Bump ActivityNewest first
|
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.
ISSUE ID : #14887
Reverts #14281.
#14281 added explicit defaults to the gfx950 Equality logic YAMLs, which breaks
test_skinny_tn_range_logicin tensilelite for the BF16 and FP16 skinny-TN Range logic. The tensilelite CI was not triggered on #14281, so it merged with clean CI, and the failure is now blocking required tensilelite jobs on develop. Tracking in #14887.Revert check: the revert applies cleanly to current develop (git merge-tree dry run), no later commits touch the same files, and #14281 is not a TheRock bump. #14885 is a test-only fix-forward that can land independently so #14281 can be relanded later.