feat(hipblaslt): replay a managed tuning cache (HIPBLASLT_TUNING_MODE=cache) by jinchen62 · Pull Request #12587 · ROCm/rocm-libraries

jinchen62 · 2026-09-25T06:34:52Z

JIRA ID: AIHPBLAS-3751

Stacked on #12586 and #12585, both merged. The diff here is this PR's alone.

Motivation

Third of six PRs splitting #10424. It adds the replay half of the runtime tuning cache. HIPBLASLT_TUNING_MODE=cache with HIPBLASLT_TUNING_CACHE_PATH=<file> replays a tuning file. Tune mode, which writes to the same path, comes next.

Compared with HIPBLASLT_TUNING_OVERRIDE_FILE, cache mode:

also serves hipblasLtMatmul calls that pass no algo, which never reach the heuristic entry point;
reports what it loaded and a closing summary of what it served, without a log level, so a deployed cache can be confirmed to work;
is refused in a secure execution context (set-uid/set-gid or file capabilities), since the variables choose a file to read.

It works with any existing file, including one written by hipblaslt-bench with HIPBLASLT_TUNING_FILE.

Technical details

Mode and file selection. TuningModeConfig::fromEnvironment lives in the store and is host-tested. It uses the existing rocblaslt_secure_getenv_impl. TuningModeSingleton reads it once per process. selectTuningFile() chooses the file:
- no mode: the override file, exactly as before;
- a mode with a path: only the cache;
- a mode without a path: nothing, with a startup line saying so.
The two files are never merged. A cache file that does not exist yet is looked for again on the next query, since tune mode creates it; the startup line reports load=not-found once.
One launch rule on both paths.
- The heuristic returns the cached entry first, and an explicit algo is launched as given.
- runContractionProblem with no algo looks up the cache and launches the first entry that still validates and is supported for this call. The lookup uses tuning_cache_find_valid_entry, a pure probe built on a new isSolutionSupportedNoMutation.
- Grouped GEMM and pointer-array batch stay on default selection for no-algo calls. A lookup that throws leaves the call on default selection.
Notices.
- log_tuning_lifecycle routes a notice through log_info at level 4 and above. Otherwise it goes to the stream logging opened, or to stderr, and it never opens the log file itself.
- The startup line reports mode, path, load status and entries loaded, and names an override file it is ignoring.
- The closing summary reports distinct shapes, matched, fell back and invalidated. It is written from a static destructor and suppressed when nothing happened.
- Hit, miss and invalidation lines appear only at info level and at most once per problem.
- With no mode set, the override file keeps its existing messages and prints no tuning-cache: lines.
Counting. Hits and misses count lookups, and an invalidated entry is counted once even when both paths meet it. The summary tally counts distinct keys.
Test hooks (exported unconditionally, like hipblaslt_debug_reload()): the reset hook now also resets the mode, diagnostics and counters. hipblaslt_tuning_last_launch_for_test() reports which solution the thread's last matmul launched. hipblaslt_tuning_stats_for_test() reports the counters; new values will only be appended.
Separate commit: runContractionProblem's exception handlers left status untouched. For a call with no algo, that meant returning getBestSolutions' success for a GEMM that never launched.
Docs:
- a cache section in the offline tuning how-to;
- the new variables in the environment reference;
- corrections to the logging docs: levels are cumulative, the mask is read only without a level, the default sink is stderr (not stdout), and HIPBLASLT_LOG_FILE does nothing without a level.

Test plan

17 new GPU cases in TuningCache_pre_checkin.* (31 in total). They use hand-written hipblaslt-bench-format rows, since nothing writes current-schema rows until tune mode. They cover:
- replay on both APIs;
- no path is inert and does not fall back to the override file;
- the override file is ignored in cache mode;
- rejected stale entries, and one stale entry counted once across both paths;
- an explicit algo launches as given and is not counted;
- a null algo launches the cached kernel and verifies the product;
- a null algo without an entry launches default selection;
- a null algo falls back from an entry that cannot run the call;
- a cache file created after the first query is read, through the heuristic and by a null algo;
- a null algo replays an entry loaded after its shape fell back;
- shape tallies, including shapes that miss an empty cache;
- the startup line's content;
- silence on repeated hits;
- no notices in override mode.
2 new host cases in TuningStore.* (25 in total): privileged processes ignore the variables, and a mode needs a known value and a path.
Mutation check, two rounds. Each of these fails at least one targeted case:
- not replacing the launch index;
- counting invalidations per sighting;
- falling back to the override without a path;
- letting the cache override an explicit algo;
- consulting the override alongside a cache path;
- marking a missing cache file loaded.

Test result

TuningCache_pre_checkin.* runs in the standard tier and above, not in PR smoke gating, so these are local results.

On MI355X (gfx950), at this head:

TuningCache_pre_checkin.*: 31/31 pass.
TuningStore.*: 24 pass. NumericLocaleDoesNotChangeReadTiming skips because the machine has no comma-decimal locale installed.

On MI300X (gfx942), earlier in review:

Smoke: 3133/3133 pass. Every target builds.
hipblaslt-bench end to end on both APIs: a file written with HIPBLASLT_TUNING_FILE replays in cache mode. The run prints tuning-cache: mode=cache path=... load=ok loaded=1, then tuning-cache: summary shapes=1 matched=1 fellback=0 invalidated=0 at exit, and cache-hit index=... once at level 4. The same file through the override variable prints no tuning-cache: lines.

Rest of the stack

fix(hipblaslt): repair the tuning override path and validate entries by kernel name #12585 (merged)
feat(hipblaslt): add a versioned tuning cache store with host tests #12586 (merged)
HIPBLASLT_TUNING_MODE=cache (this PR)
feat(hipblaslt): tune unseen shapes at their first matmul (tune mode) #12588, which also builds on refactor(hipblaslt): move hipblaslt-bench's measurement and candidate search into the library source tree #13568. That PR moves hipblaslt-bench's candidate search into the library, so tune mode and the bench share one implementation.
feat(hipblaslt): keep and revisit budget-stopped searches #12589
feat(hipblaslt-bench): tune through the library's tune mode #12590 (draft)

PRs 5 and 6 both build on PR 4 and are independent of each other.

therock-pr-bot · 2026-09-25T06:39:06Z

⏳ Waiting on Required Checks

⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	⏳ Pending	⏳ Still running…
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⏳ Required checks are still running or have not started yet.
This PR is not ready for review until they complete.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-09-25T06:39:08Z

⏳ Waiting for required checks to start or finish:

⏳ pre-commit

…by kernel name (#12585) JIRA ID: AIHPBLAS-3751 ## Motivation First of six PRs that split #10424 (runtime tuning cache) into reviewable pieces. This one only touches the existing offline tuning path (`HIPBLASLT_TUNING_FILE` / `HIPBLASLT_TUNING_OVERRIDE_FILE`). It adds no new modes or environment variables. A tuning file records `solution_index`, which is a position in one build's kernel library. Today: - The C API drops the whole file when its `Git Version:` line differs from the running build, so one differing byte discards every entry. - The C++ extension API never checks, so an index written by another build can launch an unrelated kernel. ## Technical details - `hipblaslt-bench` writes a `kernel_name` column next to `solution_index`, whether or not `print_solution_found` is set. The name comes from the algo accessor, so multi-kernel solutions record their main kernel. Grouped GEMM runs no longer write rows: the file has no grouped key, so a single-group winner read back as a single GEMM of the same shape. - Replay on both APIs resolves each entry's index in the running library and uses it only if `getKernelNameFromAlgoIndex` still returns the recorded name. A `solution_name` column is checked the same way. A row with neither is used only when the file's `Git Version:` line matches the running build. A build made outside a git checkout has no version, so it uses no such rows. - `override_path_compare_git_version` is removed, so a version-line mismatch no longer switches the file off. The loader reports it once per file instead, at error level when it drops unnamed rows and at info level otherwise. - A file that does not open is looked for again on the next query, as before. Once a file is loaded, a query takes only a read lock. - The loader pairs cells by column name, and survives a header left without values by an interrupted append. - Fixes on the same path: - Replay reused one result vector across entries and read `[0]`, so later entries for a problem were checked against the first entry's solution. - Solution index `0` was rejected. - `*returnAlgoCount` was read uninitialised when the file served a single-algo request. - On the C++ API, an entry that failed both the XF32 check and its FP32 fallback check left an XF32 problem in FP32 mode, so default selection picked FP32 kernels. - A grouped GEMM heuristic query segfaulted whenever the override file held an entry, because the lookup read the grouped problem as a single GEMM. Grouped GEMMs now skip the file. - A file with no usable rows was re-parsed on every heuristic query. - `OverrideMap::find` returned iterators after releasing its lock (latent today, since the map is not written after loading). - `getSolutionNameFromData` dereferenced a null solution for RocRoller indices. - The logger destructor could throw from `close()` at exit. - `hipblaslt_tuning_reset_for_test()` is exported unconditionally, like `hipblaslt_debug_reload()`. ## Test plan - New GPU tests `TuningCache_pre_checkin.*` (13 cases). The fixture name puts them in the existing `*pre_checkin*` category, standard tier and above, so `test_categories.yaml` is unchanged. - Mutation check. Each of these makes at least one case fail: - reintroducing the `[0]` bug; - reintroducing the whole-file disable; - removing the unnamed-row build check; - skipping the name check; - dropping the XF32 restore; - removing the grouped GEMM guard, which segfaults; - marking a file loaded when it failed to open; - trusting unnamed rows when the running build has no stamp, checked by forcing an empty stamp. - End to end on MI300X: a `hipblaslt-bench` tuning run writes the file, and the override replays the tuned index. The same file with another build's version line still replays, and a tampered `kernel_name` is rejected and falls back to the heuristic. A grouped run writes no row, and a run with `print_solution_found` off still records `kernel_name`. - Smoke tier for regressions. ## Test result `TuningCache_pre_checkin.*` runs in the standard tier and above, not in PR smoke gating, so these are local results on MI300X (gfx942). - `TuningCache_pre_checkin.*`: 13/13 pass. - Smoke: 3133/3133 pass. - Not covered by an automated test: index `0` acceptance, single parse of an empty file, RocRoller `getSolutionName()`, logger destructor. ## Rest of the stack 1. **Override-path fixes and per-entry validation (this PR)** 2. #12586 3. #12587 4. #12588 5. #12589 6. #12590 PRs 5 and 6 both build on PR 4 and are independent of each other. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

newling

Here is an agent review: review. It has 1 actionable item and no additional suggestions. Please take a look.

…12586) JIRA ID: AIHPBLAS-3751 > **Stacked on #12585.** Until that merges, this PR also shows its commits. Review only the top commit: `feat(hipblaslt): add a versioned tuning cache store with host tests`. [Compare against PR 1](jinchen62/rocm-libraries@tuning-split/1-override-validation...tuning-split/2-cache-store) ## Motivation Second of six PRs splitting #10424. The runtime tuning cache needs a row format that identifies a problem precisely and a store it can be read into. This PR adds both, with host-only tests. The existing `HIPBLASLT_TUNING_OVERRIDE_FILE` path loads through the store from now on, so it has a user from the start. No new environment variables or modes yet; those come in the next PRs. ## Technical details - **`TuningCacheStore.{hpp,cpp}`**, built as a static library (`hipblaslt-tuning-store`) linked into `hipblaslt`. It uses no device code, solution library or logger: it compiles against TensileLite's, rocisa's and HIP's headers (`hip::host`), and links nothing else from hipBLASLt. - `ProblemOverride` is widened from ten fields to the whole problem: strides, epilogue, bias, scaling formats, swizzle, stream-K mode, SM-count target, summation order, `d_type` separate from `c_type`, conjugate transpose, and the device (full `gcnArchName` and CU count). The header documents every `RocblasltContractionProblem` field that is deliberately left out, and why. - Rows with a `schema_version` column are parsed strictly: every key column is required and must be a well-formed number or known type, a current row must carry a name, and a row cut short by an interrupted append is rejected. - Rows without `schema_version` (the format `hipblaslt-bench` writes) go into a separate map and are matched on the historical ten fields only, exactly as before. - `formatTuningRow` / `appendTuningRow` write the current format. Nothing in the library writes yet; tune mode will. - **One key builder.** `RocblasltContractionProblem2ProblemOverride` builds the full key. The C++ extension builds it once at `gemmCreate`, only while a tuning file is in play, and keeps it on `TensileDataGemm`. The stream-K and summation-order preferences applied later update it. Previously the C++ path rebuilt a key from the Tensile problem with its own rules. - **Type-string helpers** moved into headers the store can use without HIP compilation, with no change in behavior: `string_to_hip_datatype` and `string_to_hipblas_computetype` go to `hipblaslt_type_strings.hpp`, and `hipDataType_to_tensile_type` and `rocComputeType_to_tensile_type` go to `tensile_type_map.hpp`. - The per-entry trust rules from PR 1 are unchanged. Named rows are validated at replay, and unnamed rows need the running build's stamp, which a build with no stamp never has. `loadTuningRows` reports how many unnamed rows it left out and the file's version line, so the loader can log them. ## Test plan - - New host-only test binary `hipblaslt-test-tuning-store` (`TuningStore.*`, 23 cases), run by TheRock CI's `test_hipblaslt.py` on the first shard and labelled for every CTest tier, quick included, via `tuning_store_test_categories.yaml`, which selects the whole binary rather than a fixture name. It covers: - round trip of every column; - FNUZ FP8 spellings; - newer schema, non-GEMM types, blank cells, trailing text, missing key columns and cut-short rows all rejected; - orphan headers; - each widened column keying only its own problem; - unnamed-row build trust, including a build with no stamp; - a version line read trimmed; - index 0; - legacy rows matching on the historical fields, in file order; - repeat and rename handling; - append and re-read; - numbers written and read the same way whatever the application's locale; - a row naming both its kernel and its solution; - an append after an interrupted line. - - Mutation check: removing the cut-row check fails two cases, and letting legacy rows key on the full key fails three. Letting an empty build stamp match an empty file stamp fails the no-stamp case. - The 14 GPU tests from PR 1 (`TuningCache_pre_checkin.*`) exercise the override path through the new store on both APIs. ## Test result - `TuningStore.*`: 23/23 pass (no GPU needed), and clean under ASAN. - `TuningCache_pre_checkin.*`: 14/14 pass on MI355X (gfx950). - Smoke: 3133/3133 pass. - `hipblaslt-bench` end to end on both APIs: a tuning run's file replays through the override, still replays with another build's version line, and a tampered `kernel_name` is rejected. The widened key's GPU-side coverage lands with tune mode (PR 4), whose tests write rows through this builder. Those tests check conjugate transpose, the stream-K preference on the C++ path, and scalar scaling without scales matching across APIs. ## Rest of the stack 1. #12585 2. **Versioned tuning cache store with host tests (this PR)** 3. #12587 4. #12588 5. #12589 6. #12590 PRs 5 and 6 both build on PR 4 and are independent of each other. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

…ts algo runContractionProblem's exception handlers left status as it was. For a call that passed no algo, status already held getBestSolutions' success, so an exception thrown while solving or launching returned success for a GEMM that never ran. Co-authored-by: Cursor <cursoragent@cursor.com>

…=cache) HIPBLASLT_TUNING_MODE=cache with HIPBLASLT_TUNING_CACHE_PATH replays a tuning file through the store, validating each entry as the override file does. - Mode and path are read once per process through TuningModeConfig, and a process in a secure execution context ignores both. With a mode set only the cache is consulted; HIPBLASLT_TUNING_OVERRIDE_FILE is ignored. - One launch rule on both paths: the heuristic hands out the cached entry first, an explicit algo launches as given, and a matmul with no algo launches the first entry that still validates and supports the call. - A startup line and a closing summary report what loaded and what was served, without a log level; per-problem hit, miss and invalidation lines appear at info level, once per key. - Test hooks report which solution a matmul launched and the tallies. - Docs: cache section in the offline tuning how-to, the variables in the environment reference, and corrections to the logging docs. Co-authored-by: Cursor <cursoragent@cursor.com>

Replaying a cached entry made hipblasLtMatmul with no algo about 2.7 us slower than default selection on a 256^3 fp16 GEMM, and a cache file that was not there added 1.5 to 2.5 us to every call. Replay ran the heuristic before looking at the cache, then on every call checked each entry's kernel name against its index and probed support on a copy of the whole Tensile problem, and a missing file was opened and stat'ed each time. Replay now runs first, and default selection only when no entry serves the call. Name checks depend only on the loaded entries and the build, so they are made once per shape and kept until an entry is added, which OverrideMap::generation now reports. Support depends on the call and is still checked each time, but on the problem the launch goes on to use, putting back the parameters the check sets rather than copying every tensor. The lookup tally hears about a shape only when its answer improves, and a missing file is looked for again at most once a second; the heuristic still looks on every query. The extended profile names the kernel from algo again, which now always holds the index launched. Measured on gfx950 as host time spent in hipblasLtMatmul over 200k calls with no algo, 256^3 fp16 with f32 compute, against 7.0 us with the cache off: replaying a row went from +2.74 to +0.02 us, a file with no row for the shape from +0.23 to +0.18, and a missing file to +0.22. Co-authored-by: Cursor <cursoragent@cursor.com>

codecov · 2026-10-09T00:33:42Z

Codecov Report

❌ Patch coverage is 73.07692% with 140 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...md_detail/rocblaslt/src/UserDrivenTuningParser.cpp	74.63%	36 Missing and 15 partials ⚠️
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp	77.58%	27 Missing and 10 partials ⚠️
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp	57.65%	33 Missing and 3 partials ⚠️
...y/src/amd_detail/rocblaslt/src/include/utility.hpp	58.82%	5 Missing and 2 partials ⚠️
...l/rocblaslt/src/include/UserDrivenTuningParser.hpp	82.14%	5 Missing ⚠️
.../src/amd_detail/rocblaslt/src/TuningCacheStore.cpp	81.25%	0 Missing and 3 partials ⚠️
..._detail/rocblaslt/src/include/TuningCacheStore.hpp	87.50%	0 Missing and 1 partial ⚠️

❌ Your project check has failed because the head coverage (43.34%) 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   #12587    +/-   ##
=========================================
  Coverage    71.98%   71.99%            
=========================================
  Files         2879     2879            
  Lines       478113   478578   +465     
  Branches     70856    70921    +65     
=========================================
+ Hits        344166   344512   +346     
- Misses      109976   110070    +94     
- Partials     23971    23996    +25

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`49.67% <ø> (ø)`		Carriedforward from 214ad34
TensileLite-Unit	`83.55% <ø> (+0.01%)`	⬆️	Carriedforward from 214ad34
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 214ad34
hipBLASLt	`43.34% <73.08%> (+0.61%)`	⬆️
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 214ad34
hipDNN	`86.80% <ø> (ø)`		Carriedforward from 214ad34
hipFFT	`42.35% <ø> (ø)`		Carriedforward from 214ad34
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 214ad34
hipSOLVER	`68.96% <ø> (ø)`		Carriedforward from 214ad34
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 214ad34
rocBLAS	`48.21% <ø> (ø)`		Carriedforward from 214ad34
rocFFT	`47.25% <ø> (ø)`		Carriedforward from 214ad34
rocRAND	`57.59% <ø> (ø)`		Carriedforward from 214ad34
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 214ad34
rocSPARSE	`74.31% <ø> (ø)`		Carriedforward from 214ad34
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 214ad34

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
.../amd_detail/rocblaslt/src/include/tensile_host.hpp	`28.57% <ø> (ø)`
..._detail/rocblaslt/src/include/TuningCacheStore.hpp	`88.16% <87.50%> (+3.44%)`	⬆️
.../src/amd_detail/rocblaslt/src/TuningCacheStore.cpp	`41.23% <81.25%> (+1.40%)`	⬆️
...l/rocblaslt/src/include/UserDrivenTuningParser.hpp	`89.80% <82.14%> (-10.20%)`	⬇️
...y/src/amd_detail/rocblaslt/src/include/utility.hpp	`49.82% <58.82%> (+0.58%)`	⬆️
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp	`45.15% <57.65%> (+0.23%)`	⬆️
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp	`76.60% <77.58%> (-0.06%)`	⬇️
...md_detail/rocblaslt/src/UserDrivenTuningParser.cpp	`75.58% <74.63%> (-5.07%)`	⬇️

... and 2 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

…ache-mode Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp

ROCm#13354 made the tuning file writers return a GoogleTest assertion result and checked it in every existing test. The cache-mode tests added here still ignored it, so a write that failed surfaced later as the wrong kernel. They now stop at the write, as the others do. Co-authored-by: Cursor <cursoragent@cursor.com>

…ache-mode

newling

Thanks for this PR, it looks good to me. There are 2 small commits in an agent review here that look legit to me, please consider absorbing them into this PR before landing.

…/4-tune-mode Brings in e6b03ec (report cache events on direct and empty lookups), 232b503 (count invalidations by saved entry identity), bdc9952 (state the rocRoller cache replay limitation) and the fix to their info-level death test. The no-path notice for a direct call moves ahead of the replay block, which this branch skips whenever no path is set. The invalidation comment keeps the recheck under the tuning lock among the paths that can meet one stale row. Co-authored-by: Cursor <cursoragent@cursor.com>

assistant-librarian · 2026-10-16T13:04:32Z

TheRock Submodule Bump Activity

Newest first

2026-10-16 — Included in TheRock via ROCm/TheRock#9032.

…#12588) JIRA ID: AIHPBLAS-3751 > #12585, #12586, #12587 and #13568 have merged, so the diff here is this PR's alone. ## Motivation Fourth PR in the split of #10424. It adds the tuner. With `HIPBLASLT_TUNING_MODE=tune`, the first `hipblasLtMatmul` of a shape the cache has no usable entry for benchmarks the candidate kernels, records the winner in `HIPBLASLT_TUNING_CACHE_PATH`, and launches it. Later calls and later processes replay it through the cache mode from PR 3. Most of #10424's review discussion is about this part. One deliberate difference from #10424: **a search the per-shape budget stops records nothing**. #10424 keeps such a search's best candidate as an "incomplete" row and revisits it later. That policy, with its extra columns and retune rule, is PR 5, so it can be reviewed on its own. The tuner measures candidates with `hipblaslt-bench`'s own candidate search, which #13568 moved into the library, and takes its candidate list and its rotation rules from the code the bench uses, so the two share one implementation. Their inputs still differ where the callers do: the workspace each candidate is given, the bench's estimate of C, and block-scaled scales (see Rotation and Workspace below). ## Technical details - **Measurement** (`benchmarkAndSelectWinner`) runs the bench's candidate search from #13568, so it measures as `hipblaslt-bench` does: - `HIPBLASLT_TUNING_COLD_ITERS` untimed launches, then `HIPBLASLT_TUNING_HOT_ITERS` launches timed as one span, on HIP events as with the bench's `--use_gpu_timer`; - the bench's instruction-cache flush between timed launches, with its per-device cost measured once and subtracted; - rotating copies of the tensors and vectors the bench rotates, over a block count planned by the bench's rules, with one workspace shared by every block (see Rotation below); - the fastest candidate wins, and a tie keeps the earlier candidate. Tune mode supplies the rest: library-owned scratch, the candidates, support checks, the budget and the cache write. - **Candidates.** Every kernel that can run the problem by default. With `HIPBLASLT_TUNING_ALL_KERNELS=0`, the first `HIPBLASLT_TUNING_MAX_CANDIDATES` kernels `hipblasLtMatmulAlgoGetHeuristic` returns, in its order, without the tuning file's entry: the ranked libraries' best, topped up from the size-independent libraries when they return fewer. The top-up is moved unchanged out of the heuristic entry point into `topUpWithSizeIndependentSolutions`, which both call. The kernel the call would otherwise launch is measured first. Each candidate gets the XCC setting the normal launch gives it. - **Rotation.** The bench's rules move into `hipblaslt_benchmark`, beside the block-count formula: a strided tensor's elements per block, a bias's, and which tensors count, with C only when beta reads it. The bench computes the same numbers through them, and tune mode takes its block count from them. Tune mode rotates what the bench rotates: A, B, C, D, E, the bias, scalar and vector scales and the alpha vector, inputs seeded from the caller's data. Its copies keep their own safe sizes: a read bias's block is sized like a gradient's, and only the caller's extent is copied into it. Block-scaled and w4a16 scales stay in place and out of tune mode's footprint, as only the caller knows their layout; the bench sizes, counts and rotates its own. One more difference is in the bench: it sizes C with `sizeof(To)`, the size of the `hipDataType` rather than of an element, so for nonzero beta it counts C at 4 bytes an element. That is more than a 16-bit or 8-bit C takes, so it plans fewer blocks (52 rather than 64 for fp16 1024³ and 512 MiB), and less than a 64-bit C takes. Fixing it changes the bench's plan, so it is left for a separate change. - **Workspace.** Each candidate is prepared with the caller's whole workspace limit, where the bench gives every candidate the largest requirement among them. A kernel whose launch depends on workspace, such as Stream-K, can therefore launch differently in the two for the same solution index. - **Isolation.** Candidates run on library-owned scratch, capped by `HIPBLASLT_TUNING_SCRATCH_MAX_BYTES` and allocated per device. The scratch is zeroed first, and a gradient epilogue's E, which it reads, is seeded with the caller's in every block. The caller's buffers are never written except by the winner's real launch. Declined: - an in-place problem with nonzero beta, because repeated launches would compound C; - a call whose fused epilogue carries an all-to-all stage, because a candidate would still signal and write its peer ranks, and the ranks do not coordinate a search. The check resolves the descriptor as the Tensile problem setup does. - **Budget.** `HIPBLASLT_TUNING_BUDGET_MS_PER_SHAPE` defaults to five minutes and is checked between candidates. A search it stops records nothing, and the shape is not searched again in this process. A failure the search reports after measuring a candidate ends it before that candidate can win. - **Concurrency.** Tuning is serialised by one lock. A thread that meets a shape being tuned waits, re-checks the cache under the lock, and launches the winner rather than repeating the search. - **Latch.** An attempt that reached `tuning-start` and left the shape wanting another search is not repeated in the process: a budget stop, setup or enumeration failure, no winner, or an exception. A completed tune is not latched, so a later call its winner cannot serve, such as one with less workspace, tunes the shape again. Declines made before `tuning-start` cost microseconds and are retried. - **One launch rule.** The winner is launched by the call that tuned only when that call passed no `algo`. An explicit `algo` still launches as given, and a reused algo keeps launching its own kernel. - **Lifecycle notices**, without a log level: `tuning-start` (the call will block), then exactly one of `tuning-done` (with `persisted=`), `tuning-skipped` or `tuning-fallback` with a reason, and `tuned=` in the summary. They are bounded per shape and per reason at the default level. At level 4 you also get the setup, a 10-second progress heartbeat, and the winner against the baseline. - **Writing rows.** `appendTunedEntry` writes the PR 2 format and uses distinct spellings for the FNUZ FP8 types, so those rows read back as the types they were tuned for. The winner replaces any dead rows for its key in memory. - **Test hooks**, exported unconditionally like PR 1's: - `hipblaslt_tuning_inject_failure_for_test(stage)` fails setup, fails enumeration, throws mid-search, stops the search after one candidate, or makes a configured budget read as already spent; - `hipblaslt_tuning_candidates_for_test` returns the candidates tune mode would measure, and `hipblaslt_tuning_scratch_plan_for_test` the scratch it would lay out; - the stats hook appends `tuned`, `skipped`, `attempts` and tuned shapes after PR 3's values. - **From James Newling's review of this PR**, his four commits: - seed a gradient epilogue's E in every scratch block (`tuning_scratch.hpp`), and decline overlapping layouts where the copy length is uncertain; - share the normal launch's XCC preparation with tune mode (`candidate_preparation.hpp`); - end the search on a failure reported after measuring a candidate; - let the cache diagnostics test run on a one-algorithm library. Also from that review: the fused all-to-all exclusion, the heuristic's candidate list and the bench's rotation plan above, and two tests made independent of the library and the clock. And his four commits from the follow-up review: - keep the XCC preparation test off a TensileLite template the Windows DLL does not export, which broke the Windows link; - decline only fused calls whose descriptor carries the all-to-all stage; - qualify the rotation comment and name the block-count test for what it checks; - say in the how-to that a completed search can be repeated, and only failed attempts are latched. Also from that review: the numerical DGELU case below. - **Cache-mode fixes suggested on #12587's review** (James Newling's commits; #12587 landed without them): - direct and empty-cache lookups report hits and misses at info level, once per problem; - a call with no `algo` reports a mode set without a path, which otherwise only the heuristic query does. This branch extends that from cache mode to tune mode, which has the same gap; - invalidations are counted per saved entry rather than per index; - the how-to states that rocRoller calls bypass cache replay. - **Docs:** a runtime tuning section in the offline tuning how-to (modes, launch rule, measurement, budget, output, limitations), and the tuning knobs in the environment reference. ## Test plan - New `TuningTune_pre_checkin.*` GPU suite (31 cases, `tuning_tune_test.cpp`). The fixture name puts it in the existing `*pre_checkin*` category, standard tier and above. It covers: - off mode writes nothing; - rows carry the key and identity columns; - exhaustive enumeration; - with `HIPBLASLT_TUNING_ALL_KERNELS=0`, the candidates are the heuristic's list in its order, at caps of 1, 16 and 256, the last reaching the size-independent top-up; - tuned launches compute the right result; - a tuned DGELU call computes the gradient of the caller's E and leaves it unchanged: with A and B all ones and E alternating between +8 and -8, D must be exactly k or zero; - FP8 FNUZ round trip; - in-place nonzero beta is declined; - replay in cache mode; - tampered names are rejected, retuned and replayed; - a stale row does not hide its replacement; - scratch is per device; - lifecycle lines appear without a log level and skips are reported once; - a call with no algo and no path reports why nothing is tuned; - a tune whose winner cannot be written is still counted; - a tuned entry is not retuned; - a budget-stopped search records nothing and is not retried, whether it stopped before or after measuring. The first case reads the budget as spent through a test hook rather than racing a 1 ms clock, and the second skips when the heuristic offers one algorithm; - one stale entry is counted once; - reused and null algos; - failed and throwing searches are latched and counted; - concurrent shapes; - a waiting thread launches the winner. It also has the three key tests deferred from PR 2: a conjugate-transpose row doesn't serve a plain transpose, the C++ path keys on the preference's stream-K mode, and scalar scaling without scales matches across the C and C++ APIs. - Host test `ReplaceAllInstallsTheWinnerAlone`, and the mode tests extended to `tune`. - 14 cases in `TuningScratchPlan_pre_checkin.*` and `TuningScratchSpan_pre_checkin.*` check the scratch plan: - spans follow the given strides and each tensor's leading dimension; overlapping strides are rounded up and flagged; a shape whose span overflows is declined; - broadcast inputs are not rotated; gradient bias, gradient E overlap and in-place C/D are covered; - the cap trims the rotation; the rotation is bounded by the launch count; every block shares one workspace; - the block count comes from actual tensor bytes: with C and D apart, fp16 1024³ and 512 MiB, 86 blocks at beta 0, as the bench plans, and 64 at beta 1, where the bench's 4-byte C gives 52; - a read bias, vector scales and the alpha vector rotate and count, and only the caller's extent of the bias is copied, per batch at its stride. - 5 host cases in `RotatingFootprint.*` cover the shared rules: whole columns and batches, a bias's stride per batch, C only when read, general batched, and saturation. - 4 cases in `CandidateSearch_pre_checkin.*` cover the hooks tune mode adds to the shared search: skipping a candidate while preparing it, ruling one out after its launches, stopping between candidates, and a failure after measurement ending the search. - `CandidatePreparation_pre_checkin` checks the per-candidate XCC setting, and `TuningScratchSeed_pre_checkin` the gradient E copy into every block. - 3 cache cases from #12587's review: direct and empty-cache lookups log once at info level, distinct stale rows sharing an index are counted separately, and a call with no algo and no path reports why the cache is inert. - Mutation check. Each of these fails at least one targeted case: - not latching a budget stop; - recording a truncated search; - not latching a search that threw; - dropping the re-check under the lock; - leaving out the heuristic's top-up (the candidate list case, at a cap of 256); - counting C at a zero beta, or leaving the vectors out of the footprint (the two block-count cases); - zeroing the caller's E during tuning (the DGELU case). - `hipblaslt-bench` is unchanged. The footprint rules it now calls were checked against its former arithmetic on a million random problems with no mismatch. A rocprofv3 dispatch trace of the bench before and after the candidate-list and rotation changes was compared on 26 configurations. Those are the 16 used for #13568, each with and without `--cotenant-cus 16`, plus a 64-candidate request reaching the heuristic's top-up, and beta 1, a bias, the alpha vector, strided batches, a bias batch stride, C equal to D and pointer-array batch. In every one, the kernel sequence, the output and the planned block count were identical. ## Test result The GPU suites run in the standard tier and above, not in PR smoke gating, so these are local results. On MI355X (gfx950), at this head: - `TuningCache*`, `TuningTune*`, `TuningScratch*`, `CandidateSearch*`, `TuningStore*`, `Rotating*` and `CandidatePreparation*`: 134 pass. Three skip: the FP8 FNUZ case, which has no solution on gfx950; the two-GPU case with one GPU visible; and `NumericLocaleDoesNotChangeReadTiming`, since the machine has no comma-decimal locale installed. The two-GPU case passes with two GPUs visible. - `hipblaslt-bench` in tune mode and then in cache mode, with `--verify`: plain, with a bias and the alpha vector, and with a bias batch stride over three batches. Each tunes once, verifies with a norm error near 1e-5, and replays as a cache hit. - `tensile_host.cpp` and `rocblaslt_auxiliary.cpp` compile with fused all-to-all enabled, including the descriptor check. - Windows: the follow-up review found `hipblaslt-test.exe` failing to link on the XCC test's call to an unexported TensileLite template. The test now checks the value the preparation sets instead, so the link has no such dependency; the hosted Windows build reruns on this head. - Tune mode against `hipblaslt-bench`'s own search with the same settings, with differing picks re-timed by the bench, before this review: - 256x256x256 TN and 2048x1024x2048 TN fp16: the same kernel. - 128x4096x4096 TN fp16: tune mode's pick is 0.8% slower. - 512x512x512 NN fp16: 3.1% slower, which is noise among near-tied kernels. Run twice, the bench itself picked the fastest kernel once and a 3.1% slower one once. On MI300X (gfx942), before the measurement moved to the shared search: - Smoke: 3133/3133 pass. Every target builds. - `hipblaslt-bench` end to end at 2048x1024x2048 fp16: - A 1 ms budget prints `tuning-start` then `tuning-skipped time budget stopped the search ...`. It writes no row, and the bench's later iterations don't retry. - A 32-candidate search prints `tuning-done winner=... elapsed=1.1s persisted=yes` and `summary ... tuned=1`. - The row that search wrote through the C API is a `cache-hit` for the C++ API in cache mode. ## Rest of the stack 1. #12585 (merged) 2. #12586 (merged) 3. #12587 (merged) 4. **`HIPBLASLT_TUNING_MODE=tune` (this PR)**, built on #13568 (merged) 5. #12589 --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

JIRA ID: AIHPBLAS-3751 > #12585, #12586, #12587, #13568 and #12588 have merged, so the diff here is this PR's alone. ## Motivation Fifth PR in the split of #10424. This is the partial-search policy that #10424's review asked to see on its own. In PR 4, a search the per-shape budget stops records nothing. A shape too large to finish inside the budget therefore never benefits from tune mode, and every process that meets it pays the budget again (once per process, thanks to the latch). This PR keeps such a search's best candidate as an **incomplete** row, and tune mode revisits it when a later run can finish or widen the search. ## Technical details - **When a truncated search is recorded.** The kernel the call would otherwise launch (the caller's algo, the cached entry, or default selection's pick) is moved to the front of the candidate list, as in PR 4. A truncated search is recorded only if that kernel was measured, so the incomplete winner is never slower than running untuned. Otherwise it is still `tuning-skipped` and nothing is written. - **What rows record.** Every tuned row now carries `complete` and `budget_ms`, and the search that produced it: all kernels or a ranked prefix and its length, the caller's workspace limit, the iteration counts, flush, and rotation size. - **Schema version 2.** These columns bump `schema_version` to 2 and are strictly required in version 2 rows, like the key columns. - Version 1 rows, as written by PR 4 builds, keep reading as finished searches with none recorded. - An older build ignores version 2 rows instead of replaying a partial row as final. - The cut-row check covers the new trailing columns. - **Retune rule** (`OverrideMap::needsRetune`): - A complete row is final if its recorded search covers this run's, meaning at least as many candidates, workspace and iterations, with flush and rotation on if this run has them on. A row with no recorded search is final. - A partial row is final if it searched exactly this way under a ceiling at least as generous. Zero means unlimited. - Otherwise the shape is tuned again. The winner is added beside the rows already there, in memory as in the append-only file, so the precedence below picks the same entry in the process that tuned it and in every later one. A partial search that tries to widen a finished one therefore leaves the finished winner in use. - Rerunning a workload with the same settings costs nothing and appends nothing. - **Precedence.** `find()` returns complete rows before partial ones, newest first within each group. Without that, a key holding both would replay the superseded partial row while `needsRetune` refused to tune again. A row that repeats an algorithm already stored moves to the newest place in its group, so a search that picks it again outranks the rows written since. A kernel's complete and partial rows are kept apart, so a later partial search that picks the same kernel cannot overwrite its finished row. Live entries and a reloaded file behave alike, and legacy rows keep file order. Those two rules are James Newling's commits from his review. - **Notices.** A partial tune prints `tuning-partial winner=...` with a hint to raise the budget. It is latched for the process like a budget skip, since a retry would stop in the same place. - **Docs:** the budget section of the how-to now covers incomplete rows, what each row records and when a shape is revisited, including a call that cannot use the winner, which a later process tunes again. The budget entry in the environment reference is updated to match. ## Test plan - 14 new host cases in `TuningStore.*` (39 in total). PR 4's `ReplaceAllInstallsTheWinnerAlone` goes with `replaceAll`, which this PR removes: - version 1 rows read as finished searches; - a row without its completion column is rejected; - rows cut before the completion or search columns are rejected; - the round trip includes the new fields; - a complete row wins over partial rows; - the newest row wins within each group; - a partial row is retuned only when this run can get further; - a partial row beside a complete one doesn't retune; - a finished search is widened but not repeated; - rows without a recorded search are final; - live entries choose as the reloaded file does, and refreshing an entry moves the map's generation; - a winner picked again becomes the newest in its group, live and reloaded, while legacy rows keep file order; - a partial result for a kernel keeps that kernel's completed row and its search; - the budget and search-coverage comparisons. - 8 new GPU cases in `TuningTune_pre_checkin.*` (38 in total): - a partial entry replays and is retuned under a more generous budget, and not under the same one; - a complete row wins over the older partial row; - a partial search measures the caller's algo first, even one past the candidates searched; - a narrower finished search is widened but not repeated; - a partial search leaves the finished winner in use; - a stale completed row does not finalize a valid partial one. The budget-exhaustion case reads the budget as spent through PR 4's test hook, so it stops before measuring anything and is an exact skip that records no row. PR 4's `SearchTheBudgetStopsAfterMeasuringRecordsNothing` is removed, since such a search now records a partial row. - Mutation check. Each of these fails at least one targeted case: - dropping complete-first ordering; - rejecting version 1 rows; - treating partial rows as final; - not measuring the untuned kernel first. ## Test result The GPU suites run in the standard tier and above, not in PR smoke gating, so these are local results. On MI355X (gfx950), at this head: - `TuningCache*`, `TuningTune*`, `TuningScratch*`, `CandidateSearch*`, `TuningStore*`, `Rotating*` and `CandidatePreparation*`: 154 pass. Three skip: the FP8 FNUZ case, which has no solution on gfx950; the two-GPU case with one GPU visible; and `NumericLocaleDoesNotChangeReadTiming`, since the machine has no comma-decimal locale installed. The two-GPU case passes with two GPUs visible. - `hipblaslt-bench` in tune mode and then in cache mode, with `--verify`, plain and with a bias and the alpha vector: each tunes once, verifies, and replays as a cache hit. On MI300X (gfx942), before #12588 moved to the shared search: - Smoke: 3133/3133 pass. Every target builds. - `hipblaslt-bench` end to end at 2048x1024x2048 fp16 with default search settings (every kernel, 1000/1000 iterations): - 3 s and 6 s budgets expire during setup, which takes about 7 s, so nothing is measured and the search is skipped with no row. - A 15 s budget prints `tuning-partial winner=219936 elapsed=15.1s persisted=yes` and writes a row with `complete=0 budget_ms=15000`. - Rerunning with 15 s does not re-tune. A 25 s budget re-tunes and appends a second partial row. - Cache mode serves the entry. ## Rest of the stack 1. #12585 (merged) 2. #12586 (merged) 3. #12587 (merged) 4. #12588 (merged) 5. **Keep and revisit budget-stopped searches (this PR)** --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

jinchen62 requested review from a team as code owners September 25, 2026 06:34

github-actions Bot added documentation project: hipblaslt labels Sep 25, 2026

jinchen62 mentioned this pull request Sep 25, 2026

feat(hipblaslt): tune unseen shapes at their first matmul (tune mode) #12588

Merged

assistant-librarian Bot added the organization: ROCm label Sep 25, 2026

jinchen62 requested review from babakpst, bethune-bryant, newling, rithwik-amd, talumbau and tony-davis September 25, 2026 06:57

jinchen62 force-pushed the tuning-split/3-cache-mode branch from 527d788 to 4313b77 Compare September 26, 2026 04:09

newling reviewed Oct 6, 2026

View reviewed changes

jinchen62 force-pushed the tuning-split/3-cache-mode branch from 4313b77 to 663ac54 Compare October 7, 2026 17:01

jinchen62 and others added 4 commits October 8, 2026 14:02

fix(hipblaslt): profile the kernel selected by cache replay

031d4b0

jinchen62 force-pushed the tuning-split/3-cache-mode branch from 663ac54 to 5a8dc64 Compare October 8, 2026 19:11

jinchen62 and others added 4 commits October 12, 2026 01:10

Merge remote-tracking branch 'upstream/develop' into tuning-split/3-c…

f4972c6

…ache-mode Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp

Merge branch 'develop' into tuning-split/3-cache-mode

1195bfa

Merge remote-tracking branch 'upstream/develop' into tuning-split/3-c…

8020b67

…ache-mode

newling approved these changes Oct 13, 2026

View reviewed changes

jinchen62 merged commit 115e23c into ROCm:develop Oct 13, 2026
132 of 134 checks passed

jinchen62 deleted the tuning-split/3-cache-mode branch October 13, 2026 20:14

yzhang93 mentioned this pull request Oct 23, 2026

feat(tensilelite): add a local kernel library checked before the shipped one #13359

Merged

1 task done

Navigation Menu

Repository navigation

feat(hipblaslt): replay a managed tuning cache (HIPBLASLT_TUNING_MODE=cache) - #12587

feat(hipblaslt): replay a managed tuning cache (HIPBLASLT_TUNING_MODE=cache)#12587
jinchen62 merged 8 commits into
ROCm:developfrom
jinchen62:tuning-split/3-cache-mode

jinchen62 commented Sep 25, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 25, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 25, 2026 •

edited

Loading

Uh oh!

newling left a comment

Uh oh!

codecov Bot commented Oct 9, 2026 •

edited

Loading

Uh oh!

newling left a comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 16, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

jinchen62 commented Sep 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical details

Test plan

Test result

Rest of the stack

Uh oh!

therock-pr-bot Bot commented Sep 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

⏳ Waiting on Required Checks

Uh oh!

therock-pr-bot Bot commented Sep 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

newling left a comment

Choose a reason for hiding this comment

Uh oh!

codecov Bot commented Oct 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

newling left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 16, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

jinchen62 commented Sep 25, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 25, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 25, 2026 •

edited

Loading

codecov Bot commented Oct 9, 2026 •

edited

Loading