Repository navigation
Conversation
fjankovi
requested review from
a team,
Jinp800125,
KKyang,
Serge45,
TonyYHsieh,
aazz44ss,
hcman2,
imcarsonliao,
jichangjichang,
solaslin and
vin-huang
as code owners
June 23, 2025 08:57
ammallya
pushed a commit
that referenced
this pull request
Sep 18, 2025
) * disable sparse functions when building a static library * explicitly disable * add include guards as failsafe * remove reference to unused variable * build without sparse by default [ROCm/hipSOLVER commit: e3d1aa8]
rahulc-gh
pushed a commit
that referenced
this pull request
Sep 23, 2025
) * disable sparse functions when building a static library * explicitly disable * add include guards as failsafe * remove reference to unused variable * build without sparse by default [ROCm/hipsolver commit: e3d1aa8]
ammallya
pushed a commit
that referenced
this pull request
Nov 5, 2025
* Added ranges support for contraction, permutation and reduction tests
ammallya
pushed a commit
that referenced
this pull request
Nov 5, 2025
* Added ranges support for contraction, permutation and reduction tests [ROCm/hipTensor commit: 2725e1a]
evetsso
pushed a commit
to evetsso/rocm-libraries
that referenced
this pull request
Dec 31, 2025
…-gfx1250-from-develop-082125 cherrypick rocroller gfx1250 PRs (merge commit) from Users/nhenders/new gfx1250 from develop 082125
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. |
Contributor
|
This pull request has been automatically closed due to inactivity (30 days with no updates). If you'd like to continue working on this, feel free to reopen the PR or create a new one. |
1 task done
tony-davis
added a commit
that referenced
this pull request
Oct 12, 2026
JIRA ID : ROCM-32153 Targets `develop`. It removes the quarantine that #13005 (AIHPBLAS-5015, now merged) added for these cases. ## Motivation On gfx950, two f16 GEMM solutions return success with wrong values for integer inputs at 4,099 x 2,053 x 129: a TN custom-schedule MT256x160x64 direct-to-LDS kernel, and an NN MT32x192x64 kernel. Neither is a tail-loop or edge problem. The TN kernel fails whenever its custom main loop runs, and the NN kernel fails at every K from 64 to 256. The TN kernel matters most, because the heuristic picks it by default for f16 TN 1,536 x 6,144 x 3,072 on MI355X, and it gave wrong answers in 6 of 30 runs there. I didn't find a small, clearly correct codegen fix for either kernel, so this PR removes them from the gfx950 logic. The custom schedule itself is tracked in ROCM-32343, which also covers its bf16 copies. ## Technical Details - Three solution entries come out of the gfx950 logic files (under `library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/`): - the TN kernel in `gfx950/Equality/..._Alik_Bljk_HHS...yaml` (#135) and its copy in `gfx950_id75a3/Equality/` (#151); - the NN kernel in `gfx950/Equality/..._Ailk_Bljk_HHS...yaml` (#334). - Each removal renumbers the solutions after it and drops the one exact-size entry the solution was tuned for, so that size falls back to the remaining selection. Every other exact-size entry still points at the same solution, and the build's logic validation passes. No tests or tuning files reference these kernels. - `known_bugs.yaml`: the two f16 `M: 4099` entries come out. ## Why remove these kernels instead of keeping them I know these kernels are fast. I looked hard at keeping them, and I don't think either one can stay as it is. **TN (custom-schedule MT256x160x64).** It fails whenever its custom main loop runs (K > 128), not just at odd sizes. That includes the one size it was tuned for: at f16 TN 1,536 x 6,144 x 3,072 the heuristic picks it by default on MI355X, and it was wrong in 6 of 30 runs. The kernel the heuristic picks instead (MT192x192x64) takes 58.7 µs against 57.1 µs, so about 2.8% slower, and it was right every time. On MI350X the difference is smaller. There, 68.6 µs against 68.1 µs is about 1% slower (the old kernel happened to be right in 10 of 10 runs there, but it fails the 4,099 x 2,053 x 129 cases on MI350X too). I think a few percent on one shape is a fair price for not returning wrong answers. **NN (MT32x192x64, StreamK 3).** I hoped to keep this one and just restrict it to the size it was tuned for (24 x 168 x 96, batch 6,144). The data doesn't support that. At exactly that size it was wrong in 30 of 30 runs on MI355X, with integer inputs (max error 3,344 to 3,480 ULP, where a correct kernel gives 0). It's wrong on MI350X there too. It's correct for up to 256 batches and wrong from 304 batches up, so my guess is it breaks once there are more tiles than CUs. On MI350X it was the heuristic's default at that size, so users were getting wrong answers from it. The replacement default (MT32x192x32) takes 84.6 µs against 74.8 µs, so it's about 13% slower, but it's correct. On MI355X the default at that size is a different kernel, and it doesn't change. One thing that made the NN problem easy to miss: `hipblaslt-bench -v` doesn't flag it. The bench computes the allclose result and then discards it, and the atol and rtol columns it prints are tolerances, not a verdict (ROCM-32137). Here the norm error was about 108, but the bench still exited 0. `--ulp` shows the problem right away. ROCM-32343 tracks fixing the custom schedule, so the TN kernels can come back once it's right. If someone fixes the NN kernel's Stream-K path, it can come back the same way. ## Test Plan On gfx950 (MI355X and MI350X): the f16 every-solution cases at 4,099 x 2,053 x 129 with the same test binary against this branch's library and the stack tip's, repeated; the ticket's `hipblaslt-bench -v` repros; and the heuristic default and timing at the sizes the removed solutions were tuned for. ## Test Result ROCm 10.2.0 (TheRock 10.2.0a20260928): - MI355X: with the stack tip's library, 4 of 16 f16 every-solution tests fail (NN and TN at 4,099); with this branch, 16 of 16 pass, and 96 of 96 over 6 repeats. - MI350X: the same, 4 of 16 failing before and 16 of 16 passing after. - The ticket's bench repros fail on the stack tip with CPU-reference errors of 0.073 and 0.045 (TN, both copies) and 0.019 (NN). - The heuristic default at 4,099 x 2,053 x 129 is unchanged and correct (MT224x160x32 for TN, MT256x144x32 for NN). - f16 TN 1,536 x 6,144 x 3,072 on MI355X now uses MT192x192x64: 58.7 µs with 0 of 10 runs wrong, against 57.1 µs for the removed kernel, which was sometimes wrong. That's about 2.8% slower. - Math CI precheckin passed against develop. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): logic-file data only, no code. Three wrong solutions stop being offered on gfx950; one tuned size gets a correct kernel that is about 3% slower on MI355X. Co-authored-by: Cursor <cursoragent@cursor.com>
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.
No description provided.