Repository navigation
feat(hipblaslt): fp32 GridBased logic + Equality retune for gfx1200/g… - #12116
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. |
42dc2f5 to
1c34df4
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #12116 +/- ##
===========================================
- Coverage 71.93% 71.93% -0.01%
===========================================
Files 2861 2863 +2
Lines 472441 472670 +229
Branches 69773 69824 +51
===========================================
+ Hits 339832 339969 +137
- Misses 109043 109118 +75
- Partials 23566 23583 +17
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
1f6f78d to
cf287d1
Compare
This sounds bad! Can you file a Github issue/ Jira ticket with some more details - thanks! |
There was a problem hiding this comment.
Thanks for fixing the "gaps" here! hipBLASLt kernel library on gfx120x is already by far our largest in RDNA, so we should quantify what we are doing to library size with this change. It's likely that we need it regardless, but the data should be captured for future efforts in library size reduction. Can you please add before/after sizes on kernel library with this PR? thank you!
bea8f00 to
a77652e
Compare
@talumbau I built the libraries locally and measured this: Same develop commit and identical CMake config for both builds (
So +747.5 KB total, under 0.2% — roughly 93% code objects, 7% No build artifacts are added or removed (282 / 292 files before and after). The 8 new GridBased YAMLs don't create new outputs, since those fp32 problem types already existed from the Equality files — the new logic just folds into the same per-problem-type Kernel count is the bigger relative jump: 16 → 302 on gfx1200 and 50 → 323 on gfx1201. They're non-MI ThreadTile MAC kernels though, so ~1.5 KB each. If size becomes a concern the easiest lever is grid density — 216 points (6×6×6) is already a subsample, and 5×5×5 would roughly halve the GridBased pool at some cost to interpolation quality on off-grid shapes. Happy to do that if you'd prefer. |
a77652e to
29bf045
Compare
29bf045 to
d5a2e12
Compare
|
Thanks for the numbers @stsokolo! I didn't realize that the fp32 coverage was so small. Although it absolute terms the number of kernels goes up a lot, this is obviously not a significant bump in size. Thanks for confirming! LGTM |
|
Cancelling the remaining gfx950 jobs on this run so the TheRock submodule bump PRs can be prioritized on the gfx950 queue. This change does not target gfx950, and gfx950 was the only work still outstanding. Jobs that already finished are unchanged. Re-run failed jobs if you need gfx950 results before merge. |
d5a2e12 to
893c1e7
Compare
…12866) JIRA ID : ROCM-29721 ## Motivation `getAllSolutions()` used to return solutions in heap-address order. For problem sizes missing from the tuned tables, `hipblasLtMatmulAlgoGetHeuristic` returns the first supported entry of that list, so identical processes could run different kernels for the same GEMM. On gfx1201 this made the ranks of a PyTorch FSDP job compute different gradients, so `TestStateDict.test_fsdp` failed intermittently (ROCM-29721), and it caused a random 2x slowdown reported in #12126. #11340 fixed the order by sorting on solution index but added no test, so nothing stops the order from regressing. ## Technical Details - New test `AlgoOrder_pre_checkin.GetAllAlgosIsSortedBySolutionIndex` checks that `hipblaslt_ext::getAllAlgos()`, which returns the same list the heuristic fallback reads, gives strictly increasing solution indices for all four transpose combinations in FP32, FP16 and BF16. - Because the list is sorted, the duplicate check in `getAllSolutions()` compares each solution with the last kept result instead of scanning every kept result. The returned list is unchanged. - `row_uniformity_gtest.cpp` no longer sorts `getAllAlgos()` results itself; its comment now states that it relies on the index order the new test checks. - Production kernel selection does not change, so there is no flag. Related: #11340 (the ordering fix this test locks), #11090 and #12126 (issues it fixed), #12116 (FP32 size-grid tables for gfx120x, which reduce how often gfx12 FP32 GEMMs reach the fallback). ## Test Plan On gfx1201 (RX 9070 XT), with hipBLASLt host libraries built from the 10.1.0a20260814 nightly source without the sort, from that source with the sort, and from this branch merged with develop: - the new gtest; - a dump of every `getAllAlgos()` index for the 12 transpose and type combinations, compared across builds. ## Test Result - The new gtest fails on all 5 runs against the build without the sort (10 to 12 of the 12 combinations out of order) and passes on all 5 runs with it. Against this branch merged with develop it passes 2/2. - `getAllAlgos()` returns the same 4,435 algorithms in the same order with the quadratic and the linear duplicate check, and in this branch merged with develop. - Automated checks are pending: TheRock CI (Linux gfx94X and gfx950, Windows gfx110X) and Math CI pre-checkin, whose gfx12 stage runs the `*_pre_checkin` gtest suites, including the new test. ## Risk level Low (2/5): a host-only change to the duplicate check that returns identical results, plus test changes. Production kernel selection is unchanged. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
TheRock Submodule Bump ActivityNewest first
|
Motivation
Neither gfx1200 nor gfx1201 ships any fp32
GridBasedlogic (0 of 108 files on gfx1201, 0 of 112 on gfx1200), soEqualityis the only matching row for fp32. An Equality miss is a hard miss, so any fp32 shape not in those exact tables gets no solution at all.HIPBLAS_COMPUTE_32F_FAST_TF32hits the same wall, since it falls back to fp32 and misses again.Reported from PyTorch as:
JIRA ID : ROCM-28540
Technical Details
Adds fp32
GridBasedlogic for both archs and retunes the fp32Equalitytables. 16 files, fp32 (_S_B_Bias_HA_S_SAV_UserArgs) only, no other datatype touched.GridBased/Equality/GridBased/Equality/GridBased cannot miss, so adding the row removes the no-solution failure. Grid is a 6-point grid, 216 points per transpose:
Equalityis row 0, so it is consulted before GridBased on every shape it covers. Those tables are retuned here so the exact-match path benefits from the same solution pool. Shape sets are unchanged, verified per file; only the selected kernel changes.Test Plan
Host with both cards (gfx1200 + gfx1201), ROCm nightly. Per
CONTRIBUTING.md, tested against an existing ROCm install viaHIPBLASLT_TENSILE_LIBPATH, run once per arch.hipblaslt-test, fp32 filter, against both the stock and the new library:hipblaslt-test --gtest_filter='*f32_rf32_rf32_rf32_rf32_r*'python -m pytest inductor/test_flex_attention.py \ -k "test_debug_flag_disables_internal_compilation_cuda" -vhipblaslt-bench, 22 shapes per arch, stock vs new. Both arms get a discarded warmup then best-of-3.TensileLogic --check-allandTensileCreateLibrary.LibraryType, arch/device names, entry counts, contiguousSolutionIndex, allExactLogicreferences in range, 100%UseDualFMAC, and the non-MI fp32 constraints (PrefetchLocalRead <= 1,ScheduleIterAlg == 1, noUseE).Test Result
Correctness: 3862 fp32 tests per arch, 0 failures. Same pass count as stock, so no coverage was lost.
quickfp32 (2108 tests)pre_checkinfp32 (1754 tests)Reported failure: fixed on both archs.
Performance: no regressions, slowest shape on either arch is ~+40%.
Representative shapes (gfx1201, TN unless noted):
Largest gains are fat-K / small-MN shapes, where
GlobalSplitUwas previously unavailable.Notes
There is also a bug in the xf32 to fp32 fallback path in
tensile_host.cpp. It is not included here and all results above were measured without it. This PR fixes the underlying cause, but the two are independent and the fallback fix should follow in its own PR.