Repository navigation
feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1 - #11155
Merged
Merged
Conversation
Add code object V6. Fix some build/packaging issues with new custom kernel format. Add DebugPattern argument which puts easy to spot data in the argument buffer for debugging.
Update build scripts to allow subdirectories in CustomKernels folder to help organize files.
Update the way CustomKernel info gets stored in a solution to improve default build, and avoid having to come up with equivalent tuning arguments to describe external kernels.
- Add padding argument semantic
e.g. { type: uint32, semantic: Padding}
- Add padding param to the arg
e.g. { type: address, semantic: AddressD, padding: 8 }
AIGECORE-1985
Fix `KeyError: 'MFMA_BF16_1K'` when building install.sh
Contributor
|
This PR is very large and will likely continue getting conflicts. I would strongly recommend you look into splitting this up into multiple PRs as a stack. The models are quite good at this task, ask them to split it in proper order (fixes, refactors, new features) and they'll churn away. |
Fix merge conflicts and update unit test
4 of 6 tasks
1 task done
… library wins (#12278) JIRA ID : AIHPBLAS-4086 Targets `GemmFromAnywhere` (#11155) rather than `develop`, so @jinchen62 can test it directly on that branch. ## Motivation The `preliminary(hipblaslt)` job on #11155 fails with **127 failed, 5 passed** in the tensilelite common tests, while the same tests pass locally on a gfx942 box. The entire 140MB CI log contains exactly one distinct error, repeated once per failed test: ``` .../caches/4463c0a054b3/source/library/gfx942/TensileLibrary.yaml:55:1: error: unknown key 'solutions' terminate called after throwing an instance of 'std::runtime_error' what(): Failed to load solution library → Aborted (core dumped), exit 134 ``` The `FileNotFoundError` on `00_Final.csv` that the `test_use_cache` tests report is purely downstream — the client aborted, so no benchmark CSV was ever written. **This is not a bug in #11155's logic.** It is a pre-existing library-shadowing bug in the test harness that #11155 is simply the first PR to expose, because it is the first to change the solution schema (`CustomKernelName` → `customKernel`). ## Technical Details 1. `tox.ini` builds `tensilelite-client` into `build_tmp/` and hands it to pytest via `--prebuilt-client`. 2. That client links `libtensilelite-host.so.1` — where the YAML deserializer (`src/llvm/Loading.cpp`) lives — and recorded a **`RUNPATH`**, not an `RPATH`. 3. **`RUNPATH` is searched after `LD_LIBRARY_PATH`.** `tox.ini` sets `LD_LIBRARY_PATH = {ROCM_PATH}/lib:...`, and the Jenkins job independently exports `LD_LIBRARY_PATH=/opt/rocm/lib`. 4. `tensilelite/CMakeLists.txt:170` installs the `tensilelite-host` target, so on a CI container with hipBLASLt installed, `/opt/rocm/lib/libtensilelite-host.so.1` exists. 5. The freshly built host library is therefore **silently bypassed**; the client runs against the older installed one, which does not understand the new schema. Why it cannot fail on a typical dev box: local `/opt/rocm/lib` has only `libhipblaslt.so*` and **no `libtensilelite-host.so`**, so the lookup falls through to `RUNPATH` and picks up the correct fresh build. The fix links the client with `-Wl,--disable-new-dtags`, emitting `RPATH`, which is searched *before* `LD_LIBRARY_PATH`. `tensilelite-client` has no `install()` rule — it is a build/test-only binary — so there is no packaging or install-time impact, and `CMAKE_INSTALL_RPATH` is untouched. ## Test Plan On a gfx942 box (ROCm 7.2.4), with a stale `libtensilelite-host.so.1` built from `develop`@merge-base placed first on `LD_LIBRARY_PATH` to emulate the CI container: 1. Direct client check: linker tag, library resolution, and a real run against the exact `ClientParameters.ini` from a failing cache dir. 2. End-to-end: CI's own pipeline — `tox -e py3` with `--cxx-flags-release=-O3`, sdist install, `invoke build-client` — scoped to the previously failing `use_cache` tests. 3. Negative control: revert the flag, relink, re-run with tox's exact `LD_LIBRARY_PATH`. ## Test Result Direct client check: | | linker tag | resolves to | client exit | |---|---|---|---| | before | RUNPATH | stale lib | **134**, `Failed to load solution library` | | after | RPATH | build-tree lib | **0** | End-to-end with the stale lib injected: **5 passed** (`test_use_cache.py::test_use_cache`, `test_compile_and_benchmark`, `test_use_cache_creates_cache_dir`, `test_use_cache_different_params`, `test_build_only.py::test_use_cache`) — all previously failing in CI. Negative control: with the flag reverted, the client picks the stale library and aborts 134 under tox's exact `LD_LIBRARY_PATH`. So the green run is attributable to the fix, not to the stale library never being on the path. Caveats, stated plainly: - Verification used a `develop`-built stand-in for the stale library, **not** the real ROCm-installed one from CI's container (which also runs Clang 24.0.0 vs 22.0.0 locally). The mechanism is proven; the exact CI binary was not reproducible off-box. Consequently my repro reports `unknown key 'customKernel'` where CI reports `unknown key 'solutions'` — same failure mode, different key, because the CI container's library is older than `develop`@merge-base. - The real confirmation is CI itself: if this hypothesis is right, `preliminary(hipblaslt)` on `GemmFromAnywhere` + this commit should go green. - **`cpu-gemm-driver` still has `RUNPATH`** and the identical latent hazard. Deliberately left alone to keep this change minimal; it is a one-line follow-up. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. Six lines in one `CMakeLists.txt`, affecting the link of a single non-installed test binary. The worst case is that `LD_LIBRARY_PATH` can no longer override the host library for `tensilelite-client` — which is precisely the behaviour being fixed, but worth noting if anyone deliberately relies on that override to A/B a library build. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude <noreply@anthropic.com>
Merged
1 task done
…match the linked LLVM (#12295) JIRA ID : AIHPBLAS-4086 Targets `GemmFromAnywhere` (#11155) rather than `develop`, so @jinchen62 can take it directly on that branch. ## Motivation `preliminary(hipblaslt)` on #11155 fails with **127 failed, 8 passed, 1144 skipped** in the tensilelite common tests, while the same tests pass on a local gfx942 box. The entire ~140MB CI log contains exactly one distinct error, repeated once per failed test: ``` .../caches/<hash>/source/library/gfx942/TensileLibrary.yaml:55:1: error: unknown key 'solutions' terminate called after throwing an instance of 'std::runtime_error' what(): Failed to load solution library → Aborted (core dumped), exit 134 ``` The `FileNotFoundError` on `00_Final.csv` that `test_use_cache` reports is a downstream symptom — the client aborted, so no benchmark CSV was written. **This is not a bug in the Gemm-From-Anywhere logic.** It is a build-configuration problem that this branch is the first to trigger. ## Technical Details `tensilelite-host` was being compiled against the **system LLVM 18** headers (`/usr/include/llvm` → `/usr/lib/llvm-18/include/llvm`) while linking **ROCm's LLVM 24** static libraries. LLVM 24 added `AliasMap` / `saveAliasHNode` to `llvm::yaml::Input`, so the class layouts differ. The compiled code touched `ValidKeys` / `Mapping` / `EC` at the wrong offsets, which is why the failure looked impossible against the source: `mapRequired("solutions")` runs, silently yields an empty vector, and the key is *then* reported as never-seen. How `-fopenmp` gets onto the host target: 1. This branch links `roc::mxDataGenerator` into `tensilelite-host`. 2. `roc::mxDataGenerator` links `OpenMP::OpenMP_CXX` (`shared/mxdatagenerator/CMakeLists.txt:14`), whose `INTERFACE_COMPILE_OPTIONS` carry `-fopenmp=libomp`. 3. `-fopenmp` makes clang add `${ROCM_PATH}/lib/llvm/bin/../include` as an **implicit system** include dir in order to locate `omp.h`. That is the same directory as the explicit `-I/opt/rocm/lib/llvm/include`, so the explicit user `-I` is demoted, and `<llvm/...>` includes fall through to the system LLVM 18 tree. develop's `tensilelite-host` has no `-fopenmp` (verified: count 0 vs 1), which is why only this branch is affected. Evidence — preprocessed TU comparison of `src/llvm/Loading.cpp`, identical compiler (Clang 24) for both: | build | `/usr/include/llvm/` markers in TU | `YAMLTraits.h` resolved to | |---|---|---| | develop | 0 | `/opt/rocm/lib/llvm/include/...` (LLVM 24) | | this branch | **356** | `/usr/include/llvm/...` (**LLVM 18**) | Flag isolation via `-E` alone: ``` PR flags (-I + -fopenmp) -> 356 keep -fopenmp, drop explicit -I -> 356 keep -I, drop -fopenmp -> 0 <-- only correct configuration drop both -> 356 ``` **The fix:** consume mxDataGenerator's include directories explicitly, and take the target itself as a `$<LINK_ONLY:...>` dependency so its `INTERFACE_COMPILE_OPTIONS` never reach our compiles. The pre-swizzle helper is correct serially, so dropping `-fopenmp` there is safe. This branch also **reverts #12278** (`5b3aae7f8d5`). That commit was my earlier, incorrect diagnosis of this same failure — it claimed a `RUNPATH`/`LD_LIBRARY_PATH` shadowing of `libtensilelite-host.so`. The CI image contains no `libtensilelite-host` and no `libhipblaslt` at all, so there was nothing to shadow, and it did not fix anything. The RUNPATH issue it addressed is real but latent and unrelated; it can be re-proposed separately against `develop` on its own merits. ## Test Plan Reproduce and verify inside the actual CI container image (`mci-base/ubu24-stg1`, Clang 24, gfx942), which is required — the failure cannot reproduce on a ROCm 7.2.4 / Clang 22 host. 1. Confirm `-fopenmp` is gone from `tensilelite-host` (`grep -c fopenmp .../tensilelite-host.dir/flags.make`) while the mxdatagen include dir and `HIPBLASLT_ENABLE_MXDATAGENERATOR` are retained. 2. Run the previously failing client invocation directly against the generated `TensileLibrary.yaml`. 3. Clean build, full CI scope: `tox -e py3 -- -m common --durations=0`. ## Test Result 1. `-fopenmp` count on `tensilelite-host`: **0**; mxdatagen include dir present; `HIPBLASLT_ENABLE_MXDATAGENERATOR` still defined. 2. Client exits **0** on the previously failing `TensileLibrary.yaml` (the one containing `customKernel`). 3. Clean build, full CI scope: ``` ========= 132 passed, 1144 skipped, 64 warnings in 2592.60s (0:43:12) ========== congratulations :) ``` `unknown key 'solutions'` occurrences: **0**. `FAILED` lines: **0**. Same 1144 skip count as the CI run, so the same test scope — versus CI's `127 failed, 8 passed, 1144 skipped`. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. One `CMakeLists.txt` hunk changing how a single target consumes mxDataGenerator's usage requirements, plus a revert of a commit that was verified to have no effect on this failure. No source or kernel-generation changes. Two follow-ups deliberately **not** included here, to keep the change minimal: - `tensilelite-client-common` still receives `-fopenmp`. It does on `develop` too, and `develop` passes, so this matches known-good behaviour — but it is worth a hardening pass. - More broadly, **any** target on this image that picks up `-fopenmp` will compile `<llvm/...>` from LLVM 18. A systemic fix would stop relying on an explicit `-I` outranking clang's implicit system include dirs. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
tony-davis
added a commit
that referenced
this pull request
Sep 18, 2026
…rid-overflow Resolves a DECISIONS.md conflict: the Gemm-From-Anywhere custom kernel framework (#11155, 1833a35) landed D24 and D25 on develop while this branch was adding its own D24. Renumbered this branch's entry to D26 and placed it after develop's two, which are taken unchanged. Only the numbering collided; the entries are unrelated. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
This was referenced Sep 24, 2026
1 task done
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.
Re-land #9304 with the fix of #11105.
JIRA ID: AIHPBLAS-4086