Repository navigation
fix(hipblaslt): launch every GSU slice of hand-written split-K kernels - #12797
Conversation
Since #11155, hand-written custom kernels launch through generateCustomCall, and for those without an explicit CustomKernel block the grid is inferred. Kernels that take no NumWorkGroups argument were given [TilesX, TilesY, Batch], which drops the GSU factor generateSingleCall multiplies into Y, so each tile launched one GSU slice. On gfx942 this hits the nine CustomGSUs_* kernels and the two fp8 x fp16 Equality shortname kernels: MultipleBufferSingleKernel kernels never reach their last-workgroup reduction and leave D unwritten, and MultipleBuffer kernels reduce a single valid slice into D. Both return HIPBLAS_STATUS_SUCCESS. Since #12599 fixed the zero macro tile, the default heuristic reaches this at the two tuned fp8 x fp16 shapes instead of crashing. - Add CustomGridSize::TilesYGSU (tiles.y * GSU) and infer [TilesX, TilesYGSU, Batch] for these kernels. - Fail library generation when a hand-written kernel can run with a GlobalSplitU other than 0 or 1 but its grid has no GSU term, and clear SupportUserGSU for such grids so a runtime GSU override is rejected during solution selection. - Throw from generateCustomCall when a split-K launch would use a grid with no GSU term, and log exceptions caught in runContractionProblem and makeArgument, so a library built before this fix returns an error instead of success. Tested on MI300X with a library built from the five affected gfx942 logic files. All 188 fp8 CustomGSUs_* solutions and 6 of the 8 fp16 ones pass CPU verification when run by index, and the default heuristic passes at both tuned fp8 x fp16 shapes. The two fp16 GSU 17 entries fail the same way on the pre-#11155 launch path; that is a separate kernel issue. TensileLite Python unit tests and tensilelite-tests pass. JIRA ID: ROCM-31692 Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ 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❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #12797 +/- ##
===========================================
- Coverage 71.93% 70.92% -1.01%
===========================================
Files 2861 2944 +83
Lines 472441 485487 +13046
Branches 69773 70098 +325
===========================================
+ Hits 339832 344316 +4484
- Misses 109043 117231 +8188
- Partials 23566 23940 +374
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Resolve the CustomGridSize conflicts with the persistent grid values added on develop (PersistentGrid, PersistentWithBatch, PersistentNoBatch), keeping their enum positions and appending TilesYGSU after them. Adapt both GSU guards to develop's execution-policy model. They now fire only for grids sized from tile counts alone, so persistent grids never trip them, and the build-time guard detects persistent kernels with isPersistent() instead of the legacy StreamK key. Co-authored-by: Cursor <cursoragent@cursor.com>
…unch The build-time GSU guard also exempted persistent kernels by execution policy, while generateCustomCall looks only at the grid, so a persistent kernel declared with a tile-count grid and a GlobalSplitU above 1 would build and then fail at launch. Every persistent grid already fails the tile-count test, so both guards now judge the grid alone. Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
left a comment
tony-davis
left a comment
There was a problem hiding this comment.
Hi Jinchen! This looks great to me. The fix lines up with the old generateSingleCall grid, the two guards use the same rule, and the tests would actually fail if someone reverted any piece of it. I also checked the red CI lanes: the ASAN gfx90a hang, the FFM gfx1250 failures, and the codecov project checks show up on other open PRs too, so I don't think they're yours. The rocSOLVER SYEVD_64 NaN on gfx950 is probably worth one rerun, just to be sure.
I'm approving, but I'd like two small things before merge (no need to rework the description):
-
Trackers for the "Not addressed here" items. Could you file tickets for the GSU 17 kernel bug, the
MultipleBufferSingleKernelworkspace with a largeldd, and the bench harness reusing D, and link them in that section? The GSU 17 entries are still in the shipped FreeSize logic, so I'd like them to have a home. For thelddone, my worry is that it could be an out-of-bounds write into the caller's workspace, not just a wrong answer, so it'd be good to capture that in the ticket (or tell me I'm wrong!). -
One sentence about gfx950. The inferred grid change also reaches 36 gfx950 custom-kernel solutions (the
Custom_Cijk_Alik_Bljk_*_shortname*_gfx950kernels). They're compiled with GSU 1 and advertiseSupportUserGSU, so with a runtime GSU override they now launch GSU slices along Y, like they did before #11155. I bet that's a fix there too. Something like this in Technical Details would cover it: "On gfx950, 36 custom-kernel solutions compiled with GSU 1 advertiseSupportUserGSU; with a runtime GSU override they now launch GSU slices along Y, as before #11155 (not verified on gfx950)." If you do get a quick gfx950 run with--gsu 4 -v, even better.
Thanks for the really thorough writeup and test evidence. It made this easy to review.
|
Thanks @tony-davis !
|
TheRock Submodule Bump ActivityNewest first
|
…client hipblaslt-bench -v launches every candidate solution into the same D and checks each one in turn, and so does the hipblaslt-test loop outside fast_check, but D was only restored between solutions under --c_equal_d. A solution that never stores D was therefore checked against the previous solution's result and reported that solution's error. On gfx942 before #12797, three CustomGSUs split-K kernels (indices 57 to 59) passed --algo_method all -v with the 1.8065e-05 norm_error of the generated solution checked before them, and failed with norm_error 1 when run alone. Fill D with all ones, NaN in the IEEE floating types, before each solution is launched whenever a result check is on and D is not C: in the candidate-search prepare hook and in the hipblaslt-test loop. fast_check keeps its own sentinel, and --c_equal_d keeps restoring C from the host. JIRA ID: ROCM-32357 Co-authored-by: Cursor <cursoragent@cursor.com>
…14743) ROCM-32357 ## Motivation All solutions in one hipblaslt-bench run write the same D buffer. D was filled with a sentinel between solutions only under fast_check, and restored only under c_equal_d. With -v (norm_check) and separate C and D, a solution that leaves D unwritten was checked against the previous solution's result and could pass; this hid the CustomGSUs_* kernels of ROCM-31692 under --algo_method all. ## Technical Details Fill D with the fast_check sentinel before every solution whenever unit_check, norm_check, allclose_check or fast_check is on, in both the verification loop and the timing path's prepareCandidate. Fill every D buffer over its full size, so pointer-array batches and rotating-buffer blocks are covered, and fall back to all-ones bytes for D types the sentinel helper does not support instead of aborting the run. ## Test Plan CI, tox and check before #12797 ## Test Result pass and as expected ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. Co-authored-by: victorwu <victor.wu2@amd.com>
…t never launches (#14647) JIRA ID : ROCM-32357 ## Motivation `hipblaslt-bench` and `hipblaslt-test` share one test harness, `testing_matmul.hpp`. It can run several solutions (kernel configurations) for one matrix multiply and compare each result, the output matrix D, with a CPU reference. All solutions write into the same D buffer, so a solution that never stored D used to be checked against the previous solution's result and pass. On gfx942 before #12797, `hipblaslt-bench --algo_method all -v` passed three hand-written `CustomGSUs_*` kernels this way. #14743 now fills D with a NaN sentinel before each solution, but no test fails if that fill stops working. ## Technical Details - A new test argument, `skip_launch_solution`, leaves the solution at that position in the list unlaunched (C interface with strided batches only), so that solution never writes D. After each solution's check, the harness fails the case if the allclose check reported a solution other than the unlaunched one or missed the unlaunched one, and also if the unlaunched solution was never checked. - The new per-PR (`pre_checkin`) case `matmul_check_reports_unlaunched_solution` sets it for f32 and bf16 at 131×67×200, in the `hipblaslt-test` loop and in the bench's loop over candidate solutions (`timing: 1`). It uses allclose alone because `unit_check` and `norm_check` fail a case outright when D holds NaN. Cases with `timing: 1` now get a `_TIMING` name suffix. - Only test code changes; the library, its kernels and the clients' behavior without the new argument do not. `Arguments`, the clients' parameter struct, gains one `int32_t` field, added at the same position in `hipblaslt_common.yaml`, from which `hipblaslt_gtest.data` is generated at build time. - This branch first carried its own fill of D. #14743 landed an equivalent fill on develop, and the merge of develop keeps develop's `fillDWithSentinel`, so this pull request now adds only the self-test. ## Test Plan On MI300X (gfx942), with hipBLASLt built from this branch after the merge of develop and `HIPBLASLT_LIBLOGIC_PATH` limited to six logic files: fp8 × fp16 → f16, f16 → f32, grouped GEMM, aux output, and the f32 and bf16 N/T grid-based files cut to the 8 solutions nearest 131×67×200. - `matmul_check_reports_unlaunched_solution` with develop's fill and with the fill disabled, plus the existing `fast_check_inject` self-test. - Every solution with `-v`, to check that the fill passes correct kernels: fp8 × fp16 at 4096×16×8192 with `--verify_strict`, `--c_equal_d`, `--rotating 256` with timed iterations, `--api_method mix` and `--HMM`; f16 → f32 at 4096×16×2048 with the C and C++ interfaces; f16 → f32 with pointer-array batches; grouped GEMM; aux output. ## Test Result | Case | Fill disabled | Develop's fill | |---|---|---| | `matmul_check_reports_unlaunched_solution`, 4 cases | 4 fail: "solution 1 of 4 … was never launched, but D passed its checks" | 4 pass, 134 ms in total | - With every solution launched, every check passes with develop's fill: 53 fp8 × fp16 solutions in each bench configuration (`--verify_strict` exits 0), 39 f16 → f32 solutions with each interface, 37 with pointer-array batches, and the single grouped-GEMM and aux-output solutions. The `fast_check_inject` self-test passes. - On the head before the merge, all four new cases ran and passed in TheRock's gfx942 and gfx110X hipBLASLt lanes. - Automated GitHub checks for the merged head are pending. ### Not addressed here - Under `unit_check` and `norm_check`, the harness still does not reset E (the aux output), the bias gradient or amaxD between solutions; only `fast_check` fills them. [ROCM-33792](https://amd-hub.atlassian.net/browse/ROCM-33792) tracks this. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. [ROCM-33792]: https://amd-hub.atlassian.net/browse/ROCM-33792?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Cursor <cursoragent@cursor.com>
JIRA ID : ROCM-31692
Motivation
Since #11155, hand-written custom kernels launch through
generateCustomCall. When a kernel'scustom.confighas noCustomKernel:block,CustomKernels.pyinfers one, and for kernels that take noNumWorkGroupsargument it inferred the grid[TilesX, TilesY, Batch].generateSingleCall, which launched these kernels before #11155, multiplies Y by GSU. The inferred grid does not, so every tile launched only one of its GSU slices.This affects 11 gfx942 kernels (198 solutions): the nine
CustomGSUs_*kernels (FreeSize logic, GSU 2–32) andCustom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_AS_SAB_SAV_shortname0/1_gfx942(Equality logic, GSU 16 and 2). Both failure modes returnHIPBLAS_STATUS_SUCCESS:CustomGSUs_*kernels useMultipleBufferSingleKernel, where the last workgroup to reach a tile reduces into D. With one slice launched that never happens, so D keeps whatever it held (norm_error1, or NaN).shortname0/shortname1useMultipleBuffer, where the PostGSU kernel sums every workspace slice, but only one was written. Since fix(tensilelite): Do not overwrite logic-file tiles with a zero custom-kernel macrotile #12599 fixed their zero macro tile, the default heuristic selects them at 8192×16×65536 and 16384×32×4096 (fp8 × fp16 with bias or scale A/B) and returns a wrong D (norm_error≈ 0.96 and ≈ 0.7) where it used to SIGFPE.With
TENSILE_DB=0x40, a GSU 16 fp16 kernel at 4096×16×2048 receives GSU 16 ininternalArgsbut launchesNum work groups: 64, 1, 1where64, 16, 1is needed.Technical Details
CustomGridSize::TilesYGSU, which launchestiles.y × GSUwork groups using the runtime-effective GSU (the same value encoded ininternalArgsand theSplitKargument).CustomKernels.pynow infers[TilesX, TilesYGSU, Batch]for kernels without aNumWorkGroupsargument. That is the gridgenerateSingleCallgave them before feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1 #11155, and it is identical toTilesYat GSU 1.Solution._assignCustomKernelParametersfails library generation when a hand-written kernel can run with aGlobalSplitUother than 0 or 1 (including -1, auto) but its grid is sized from tile counts alone (no GSU or persistent term), the same rulegenerateCustomCallapplies at launch. No hipBLASLt or hipSPARSELt logic file trips it. For such grids it also clearsSupportUserGSU, socheckInternalArgumentsSupportrejects a runtime GSU override during solution selection. This only affects the explicit-grid external kernels (aiter, CK, RRGEMM, Triton, wave), none of which is referenced by shipped logic.generateCustomCallthrows when the runtime GSU is above 1 and the grid is sized from tile counts alone, so a library built before this change returnsHIPBLAS_STATUS_INTERNAL_ERRORinstead of success.std::exceptionhandlers inrunContractionProblemandmakeArgumentnowlog_errorthe message (shown withHIPBLASLT_LOG_LEVEL=1). They previously discarded it.Device libraries must be rebuilt to pick up the new grid, as with #12599. A host without
TilesYGSUfails to load a library that uses it (the msgpack reader reports "Enum not found") rather than misreading it.Test Plan
tensilelite-testssuite. NewCustomKernelTestcases coverTilesYGSUand the runtime guard.HIPBLASLT_LIBLOGIC_PATHlimited to the five affected logic files:CustomGSUs_*solution by index, each in its own process with CPU verification (-v): fp16 → fp32 with bias at 4096×16×2048, and fp8 × fp16 at 4096×16×8192 with and without bias--scaleA 1 --scaleB 1--algo_method allat 4096×16×8192, and a 4096-solution heuristic request at 4096³developgenerateSingleCall.Test Result
tensilelite-tests: 503 passed, 8 skipped (existing skips).[TilesX, TilesYGSU, Batch], with no zero macro tile and user GSU still enabled. The launch grid is now 64×16×1 for the fp16 GSU 16 kernel and 16×16×1 forshortname0.developCustomGSUs_*by index (188)norm_error1 or NaN)CustomGSUs_*by index (8)norm_error1)shortname0)shortname1)shortname0at 8192×16×65536 matches the pre-feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1 #11155 library exactly (1.51036e-05).developlibrary, this host returnsHIPBLAS_STATUS_INTERNAL_ERRORand logs... runs with GSU 16 but its custom-kernel grid [TilesX, TilesY, Batch] launches one GSU slice per tile.Not addressed here
CustomGSUs_*kernels give a different wrong answer on each run, and do the same on the pre-feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1 #11155 launch path. They still fail when K divides evenly into 17 slices and pass with a single tile, which points at the kernel. ROCM-31692's first acceptance criterion depends on that separate fix.MultipleBufferSingleKernelworkspace with a largeldd(ROCM-32356):AddressDpoints at the workspace whileStrideD0staysldd, and the workspace is sized from tiles × GSU. Oncelddexceeds M, the partial-sum stores run past the end of the workspace: a fault if that memory is unmapped, silent corruption of whatever follows it otherwise. This predates feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1 #11155 and is reachable again now that every slice launches.hipblaslt-bench --algo_method allreuses D across algorithms, which hid this regression.Submission Checklist