Repository navigation
feat(hipblaslt): keep and revisit budget-stopped searches - #12589
Merged
jinchen62 merged 88 commits intoOct 26, 2026
Merged
Conversation
This was referenced Sep 25, 2026
✅ 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. |
jinchen62
requested review from
babakpst,
bethune-bryant,
newling,
rithwik-amd,
talumbau and
tony-davis
September 25, 2026 06:58
jinchen62
force-pushed
the
tuning-split/5-partial-searches
branch
from
September 26, 2026 04:09
1dc9610 to
8130bdd
Compare
jinchen62
added a commit
that referenced
this pull request
Oct 1, 2026
…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>
jinchen62
force-pushed
the
tuning-split/5-partial-searches
branch
from
October 7, 2026 17:01
8130bdd to
fef0aeb
Compare
jinchen62
added a commit
that referenced
this pull request
Oct 8, 2026
…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>
HIPBLASLT_TUNING_MODE=tune benchmarks the candidate kernels for a shape the cache has no usable entry for, at that shape's first hipblasLtMatmul, appends the winner to HIPBLASLT_TUNING_CACHE_PATH and launches it. - Candidates are measured the way hipblaslt-bench measures them: warm-up and timed launches over rotating copies of every buffer, with the bench client's instruction-cache flush, on library-owned scratch. The fastest wins outright. - Tuning is serialised; a thread that waited re-checks the cache under the lock and launches the winner instead of searching again. - HIPBLASLT_TUNING_BUDGET_MS_PER_SHAPE (five minutes by default) bounds a search. A search it stops records nothing. - An attempt that reached tuning-start and left the shape wanting another search is latched for the process. - tuning-start and one terminal line per tuned shape are written without a log level, bounded per shape and reason; tuned= joins the summary. - A failure-injection test hook covers setup, enumeration, a throw mid-search and an expiring budget. Adds the TuningTune_pre_checkin GPU suite, including the key tests that need tune mode to write rows. Co-authored-by: Cursor <cursoragent@cursor.com>
… clock SearchTheBudgetStopsAfterMeasuringRecordsNothing needs a second candidate for the injected stop to be asked again; it now skips when the heuristic offers only one. BudgetExhaustedShapeIsNotRetried assumed a 1 ms budget always expires during the search; a new injection stage, 5, makes a configured budget read as already spent, so the stop no longer depends on how fast the machine is. Co-authored-by: Cursor <cursoragent@cursor.com>
The ranked search took the ranked libraries' best alone, so where they return few kernels tune mode measured fewer candidates than hipblasLtMatmulAlgoGetHeuristic offers and hipblaslt-bench measures. The heuristic's top-up from the size-independent libraries is moved unchanged into topUpWithSizeIndependentSolutions, which the C heuristic entry point now calls. Tune mode's ranked enumeration is getBestSolutions followed by that top-up, so it is the heuristic's list in its order, without the tuning file's entry. A test hook exposes the enumeration, and a test compares it with the heuristic's list at caps of 1, 16 and 256. Co-authored-by: Cursor <cursoragent@cursor.com>
The two passed different footprints to the shared block-count formula, so one problem could plan different rotations: tune mode counted C even at a zero beta and left out the scale vectors and a read bias, which the bench rotates. With C and D apart, beta 0, fp16 1024^3 and 512 MiB, the bench plans 86 blocks and tune mode planned 64. 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 it is read. The bench computes the same numbers through them, and tune mode now takes its block count from them. Tune mode also rotates what the bench rotates: a bias the epilogue reads, scalar and vector scales and the alpha vector, seeded from the caller's like A and B. Its blocks keep their 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, as only the caller knows their layout. Co-authored-by: Cursor <cursoragent@cursor.com>
…earches Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/hipblaslt/clients/tests/src/tuning_tune_test.cpp # projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp
getAllSolutions has returned solutions in index order since ROCm#11340, so tune mode's own sort of the all-kernels list repeated it, and its comment, that the order differs from one process to the next, no longer held. Co-authored-by: Cursor <cursoragent@cursor.com>
…earches Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp
The test called ContractionSolution::calculateInternalArgs<false> to read back the packed XCC value. That member template is instantiated inside the tensilelite library and is not exported, so hipblaslt-test.exe failed to link on Windows with an undefined symbol. Check the parameter prepareCandidateXcc sets instead. The packer uses a nonzero value as the mapping and zero as the candidate's own, so the parameter is what carries the choice. The test still starts from a stale value, alternates fallback and ordinary candidates, and fails if the assignment is removed.
isFusedAllToAll treated any fused epilogue descriptor as an all-to-all call. The descriptor says whether it has the prefix, and every other place that acts on it, including the Tensile problem setup, resolves it and checks hasA2APrefix. Use the same test so tune mode declines the calls it documents and leaves other descriptors on the normal tuning path. Checked by compiling tensile_host.cpp with fused all-to-all enabled.
The benchmark counts fp16 C with sizeof(hipDataType), giving 52 rather than 64 blocks for beta=1 at 1024 cubed and 512 MiB. Preserve both implementations and describe the runtime test as actual-byte planning. Remove the incorrect claim that neither caller counts block-scaled scales.
A DGELU call with no algo is tuned in tune mode and must still compute the gradient of the caller's E and leave it unchanged. With A and B all ones and E alternating between +8 and -8 by row, the GELU derivative is 1 or 0 to well within fp16 in either GELU form, so D is exactly k or zero; a product scaled by a zeroed E would be k / 2. Zeroing the caller's E during tuning fails the case. Co-authored-by: Cursor <cursoragent@cursor.com>
…annot serve The latch comment said a completed tune's entry closes the gate. It does only for calls the winner supports; one it does not, such as a call with less workspace, tunes the shape again, as the how-to now says. Co-authored-by: Cursor <cursoragent@cursor.com>
…earches Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/hipblaslt/docs/how-to/how-to-use-hipblaslt-offline-tuning.rst # projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp
jinchen62
added a commit
that referenced
this pull request
Oct 23, 2026
…#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>
newling
self-requested a review
October 23, 2026 20:00
newling
requested changes
Oct 23, 2026
newling
left a comment
newling
left a comment
Contributor
There was a problem hiding this comment.
I'll take a look at this once it is rebased (currently I see over 6000 lines added in this PR). For now, an agent review suggests 2 improvements, PTAL https://github.com/newling/agent-public-reviews/blob/main/rocm-libraries/review_pr12589_v3_20261023.md
Its changes are already on this branch, which carries ROCm#12588's own history; this records the squash so develop merges cleanly.
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task
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-3751
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
tuning-skippedand nothing is written.completeandbudget_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_versionto 2 and are strictly required in version 2 rows, like the key columns.OverrideMap::needsRetune):find()returns complete rows before partial ones, newest first within each group. Without that, a key holding both would replay the superseded partial row whileneedsRetunerefused 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.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.Test plan
14 new host cases in
TuningStore.*(39 in total). PR 4'sReplaceAllInstallsTheWinnerAlonegoes withreplaceAll, which this PR removes:8 new GPU cases in
TuningTune_pre_checkin.*(38 in total):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
SearchTheBudgetStopsAfterMeasuringRecordsNothingis removed, since such a search now records a partial row.Mutation check. Each of these fails at least one targeted case:
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*andCandidatePreparation*: 154 pass. Three skip: the FP8 FNUZ case, which has no solution on gfx950; the two-GPU case with one GPU visible; andNumericLocaleDoesNotChangeReadTiming, since the machine has no comma-decimal locale installed. The two-GPU case passes with two GPUs visible.hipblaslt-benchin 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:
hipblaslt-benchend to end at 2048x1024x2048 fp16 with default search settings (every kernel, 1000/1000 iterations):tuning-partial winner=219936 elapsed=15.1s persisted=yesand writes a row withcomplete=0 budget_ms=15000.Rest of the stack