Repository navigation
fix(tensilelite): Do not overwrite logic-file tiles with a zero custom-kernel macrotile - #12599
Merged
dsaffars merged 1 commit intoSep 28, 2026
Merged
Conversation
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✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (48.80%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #12599 +/- ##
========================================
Coverage 69.94% 69.94%
========================================
Files 2933 2933
Lines 482194 482200 +6
Branches 69514 69516 +2
========================================
+ Hits 337261 337272 +11
+ Misses 120723 120720 -3
+ Partials 24210 24208 -2
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
talumbau
approved these changes
Sep 25, 2026
talumbau
left a comment
talumbau
left a comment
Collaborator
There was a problem hiding this comment.
Thanks for adding test coverage LGTM 👍
dsaffars
deleted the
users/dsaffars/algomethod_floatingpoint_exception
branch
September 28, 2026 09:17
1 task done
Contributor
TheRock Submodule Bump ActivityNewest first
|
jinchen62
added a commit
that referenced
this pull request
Oct 7, 2026
#12797) JIRA ID : ROCM-31692 ## Motivation Since #11155, hand-written custom kernels launch through `generateCustomCall`. When a kernel's `custom.config` has no `CustomKernel:` block, `CustomKernels.py` infers one, and for kernels that take no `NumWorkGroups` argument 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) and `Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_AS_SAB_SAV_shortname0/1_gfx942` (Equality logic, GSU 16 and 2). Both failure modes return `HIPBLAS_STATUS_SUCCESS`: - The `CustomGSUs_*` kernels use `MultipleBufferSingleKernel`, 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_error` 1, or NaN). - `shortname0`/`shortname1` use `MultipleBuffer`, where the PostGSU kernel sums every workspace slice, but only one was written. Since #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 in `internalArgs` but launches `Num work groups: 64, 1, 1` where `64, 16, 1` is needed. ## Technical Details - **Grid:** Add `CustomGridSize::TilesYGSU`, which launches `tiles.y × GSU` work groups using the runtime-effective GSU (the same value encoded in `internalArgs` and the `SplitK` argument). `CustomKernels.py` now infers `[TilesX, TilesYGSU, Batch]` for kernels without a `NumWorkGroups` argument. That is the grid `generateSingleCall` gave them before #11155, and it is identical to `TilesY` at GSU 1. - **Build-time guard:** `Solution._assignCustomKernelParameters` fails library generation when a hand-written kernel can run with a `GlobalSplitU` other than 0 or 1 (including -1, auto) but its grid is sized from tile counts alone (no GSU or persistent term), the same rule `generateCustomCall` applies at launch. No hipBLASLt or hipSPARSELt logic file trips it. For such grids it also clears `SupportUserGSU`, so `checkInternalArgumentsSupport` rejects 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. - **Runtime guard:** `generateCustomCall` throws when the runtime GSU is above 1 and the grid is sized from tile counts alone, so a library built before this change returns `HIPBLAS_STATUS_INTERNAL_ERROR` instead of success. - **Logging:** The generic `std::exception` handlers in `runContractionProblem` and `makeArgument` now `log_error` the message (shown with `HIPBLASLT_LOG_LEVEL=1`). They previously discarded it. Device libraries must be rebuilt to pick up the new grid, as with #12599. A host without `TilesYGSU` fails to load a library that uses it (the msgpack reader reports "Enum not found") rather than misreading it. ## Test Plan - TensileLite Python unit tests for custom kernels, and the full `tensilelite-tests` suite. New `CustomKernelTest` cases cover `TilesYGSU` and the runtime guard. - MI300X (gfx942), hipBLASLt built with `HIPBLASLT_LIBLOGIC_PATH` limited to the five affected logic files: - every `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 - default heuristic at both tuned fp8 × fp16 shapes, with bias and with `--scaleA 1 --scaleB 1` - `--algo_method all` at 4096×16×8192, and a 4096-solution heuristic request at 4096³ - this host against a library built from `develop` - The same fp16 solutions on a library from before #11155, which still launches them through `generateSingleCall`. ## Test Result - Python: 1002 passed, 3 skipped, 26 snapshots unchanged. `tensilelite-tests`: 503 passed, 8 skipped (existing skips). - The rebuilt library serializes all 198 split-K solutions as `[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 for `shortname0`. | Case | `develop` | This PR | |---|---|---| | fp8 × fp16 `CustomGSUs_*` by index (188) | all fail (`norm_error` 1 or NaN) | all pass | | fp16 → fp32 `CustomGSUs_*` by index (8) | all fail (`norm_error` 1) | 6 pass; the two GSU 17 entries fail (see below) | | Default heuristic, 8192×16×65536, bias / scale A/B (`shortname0`) | 0.963 / 0.966 | 1.51e-05 / 1.62e-05 | | Default heuristic, 16384×32×4096, bias / scale A/B (`shortname1`) | 0.682 / 0.706 | 1.07e-05 / 1.43e-05 | - The fp16 GSU 16, 18 and 19 results are bit-identical to the pre-#11155 launch path, and `shortname0` at 8192×16×65536 matches the pre-#11155 library exactly (1.51036e-05). - Enumeration and the 4096-solution heuristic request complete. Against a `develop` library, this host returns `HIPBLAS_STATUS_INTERNAL_ERROR` and logs `... runs with GSU 16 but its custom-kernel grid [TilesX, TilesY, Batch] launches one GSU slice per tile`. ### Not addressed here - **GSU 17 kernel bug (ROCM-32355):** The GSU 17 entries of the two fp16 `CustomGSUs_*` kernels give a different wrong answer on each run, and do the same on the pre-#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. - **`MultipleBufferSingleKernel` workspace with a large `ldd` (ROCM-32356):** `AddressD` points at the workspace while `StrideD0` stays `ldd`, and the workspace is sized from tiles × GSU. Once `ldd` exceeds 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 #11155 and is reachable again now that every slice launches. - **Bench harness (ROCM-32357):** `hipblaslt-bench --algo_method all` reuses D across algorithms, which hid this regression. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- 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.
JIRA ID : AIHPBLAS-4871
Motivation
Dense solution enumeration (
hipblaslt-bench --algo_method 1/ geko search) dies withFloating point exceptionwhen the walked library contains a handwritten custom kernel whose macrotile cannot be inferred. Pinned-index (algo_method 2) and heuristic (algo_method 0) paths work because they never cost the zero-tile custom solution.The crash is host-side integer divide-by-zero in
ContractionProblemGemm::getNumTiles():CustomKernels.pyderivesmacrotilefromMatrixInstruction/MIWaveTileincustom.config, then from anMTxxxtoken in the kernel name. Any custom kernel that has neither infers zeros, and_assignCustomKernelParametersthen copied that over the tile already present on the solution from the logic YAML. The built device library storedsizeMapping.macroTile: [0, 0, 1](or[0, 0, DepthU]when onlyDepthUis incustom.config).Technical Details
In
Solution._assignCustomKernelParameters, keep usingCustomKernel.macrotilewhen every component is positive. If a component is missing or<= 0, keep the value already onstatefrom the consuming logic file. Write the resolved[MacroTile0, MacroTile1, DepthU]back into both:state["MacroTile0" / "MacroTile1" / "DepthU"]— whatSizeMapping.ReadOriginalMacroTileserializesstate["CustomKernel"]["macrotile"]— what the C++ runtime reads incalculateTiles/ workspace sizingIf neither source supplies a positive integer, raise at library-build time rather than emit a solution that divides by zero at selection time. Logic-file default
DepthU: -1is rejected as a tile.Kernels that already declare a real macrotile (BBS / HHS / MX Stream-K custom kernels) are unchanged: a positive inferred tile still wins over the logic file.
This is a host-library-generation fix. Existing device libraries must be rebuilt for the affected custom solutions to pick up non-zero tiles.
Test Plan
PYTHONPATH=$PWD python3 -m pytest Tensile/Tests/unit/test_custom_kernel_host.py --noconftestin thedsgemmcontainer (39 passed, including the four new cases below).hipblaslt-bench --yamlF8BS config (transA: T, transB: N, M: 21504, N: 4720, K: 5376, a_type/b_type: f8_r, c_type/d_type: bf16_r, algo_method: 1) — enumeration completes.New unit coverage:
[0,0,0]falls back to logic-file256×256×128and writes that back intoCustomKernel.macrotileRuntimeErrorDepthU: -1is not accepted as a tileTest Result
Unit suite: 39 passed, 0 failed.
After device-lib rebuild, F8BS dense enumeration no longer SIGFPEs (same config that previously dumped core).