Repository navigation
perf(hipblaslt): Correct gfx1201 HHS NN tile selection for large-M prefill GEMMs - #10222
Merged
Merged
Conversation
✅ 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 #10222 +/- ##
========================================
Coverage 69.77% 69.77%
========================================
Files 2810 2810
Lines 463930 463930
Branches 68433 68433
========================================
Hits 323682 323682
Misses 116804 116804
Partials 23444 23444
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
amd-maradosa
changed the title
amd-maradosa
force-pushed
the
users/amd-maradosa/gfx1201-hhs-nn-logic-fix
branch
from
August 27, 2026 12:02
2e86444 to
a167300
Compare
KKyang
approved these changes
Sep 1, 2026
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Sep 1, 2026
perf(hipblaslt): Correct gfx1201 HHS NN tile selection for large-M prefill GEMMs (#10222) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation [Jira issue](https://amd-hub.atlassian.net/browse/ROCM-21133) JIRA ID: ROCM-21133 Several ExactLogic entries in the gfx1201 HHS NN (`Ailk_Bljk`) Bias+SAV logic file select `MT16x16x32` / `MT16x32x32` for M=4096 prefill shapes, while every neighbouring entry selects `MT128x128x32`. Because `GridBasedMatching` rounds *up* rather than picking the nearest neighbour (`PropertyMatching.hpp`: `entryM >= queryM && entryN >= queryN`), one bad entry becomes the mandatory target for a whole query range. The file also had no entries at K=14336, so Llama-3-8B `down_proj` shapes fell through to unrelated neighbours Effect: these GEMMs ran at ~5% of peak on gfx1201. ## Technical Details Data-only change to one logic file: 6 entries repointed, 11 added. All reference solutions already present in the file — no kernel codegen changes, solution count unchanged (517), no entries removed. ## Test Plan - `hipblaslt-test --gtest_filter='*quick*:*smoke*'` against the rebuilt gfx1201 device library, to confirm no existing kernel selection is broken. - `hipblaslt-bench` on gfx1201, HHS NN + bias + scaleAlphaVec, `--rotating 1024`, with `--print_kernel_info` to verify the selected tile. - End-to-end: vLLM 0.21 Llama-3-8B-Instruct-GPTQ latency (tp1, fp16, out=32, mp executor). ## Test Result hipblaslt-test: **21698 tests, 21698 passed, 0 failed.** hipblaslt-bench: M=4096 N=2048 K=4096 MT16x16x32 11.8 TFLOP/s -> MT128x128x32 140.3 TFLOP/s M=4096 N=2048 K=14336 MT16x16x32 7.6 TFLOP/s -> MT128x128x32 122.6 TFLOP/s vLLM latency: - bs1 in2048 1.934 s -> 0.783 s - bs2 in2048 3.475 s -> 1.152 s - bs4 in2048 3.032 s -> 1.836 s Co-authored-by: Marko Radosavljevic <276607479+amd-maradosa@users.noreply.github.com>
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.
Motivation
Jira issue
JIRA ID: ROCM-21133
Several ExactLogic entries in the gfx1201 HHS NN (
Ailk_Bljk) Bias+SAV logic file selectMT16x16x32/MT16x32x32for M=4096 prefill shapes, while every neighbouring entry selectsMT128x128x32. BecauseGridBasedMatchingrounds up rather than picking the nearest neighbour (PropertyMatching.hpp:entryM >= queryM && entryN >= queryN), one bad entry becomes the mandatory target for a whole query range. The file also had no entries at K=14336, so Llama-3-8Bdown_projshapes fell through to unrelated neighboursEffect: these GEMMs ran at ~5% of peak on gfx1201.
Technical Details
Data-only change to one logic file: 6 entries repointed, 11 added. All reference solutions already present in the file — no kernel codegen changes, solution count unchanged (517), no entries removed.
Test Plan
hipblaslt-test --gtest_filter='*quick*:*smoke*'against the rebuilt gfx1201 device library, to confirm no existing kernel selection is broken.hipblaslt-benchon gfx1201, HHS NN + bias + scaleAlphaVec,--rotating 1024, with--print_kernel_infoto verify the selected tile.Test Result
hipblaslt-test: 21698 tests, 21698 passed, 0 failed.
hipblaslt-bench:
M=4096 N=2048 K=4096 MT16x16x32 11.8 TFLOP/s -> MT128x128x32 140.3 TFLOP/s
M=4096 N=2048 K=14336 MT16x16x32 7.6 TFLOP/s -> MT128x128x32 122.6 TFLOP/s
vLLM latency: