Repository navigation
fix(hipblaslt): repair the tuning override path and validate entries by kernel name - #12585
Merged
jinchen62 merged 8 commits intoOct 1, 2026
Merged
Conversation
…by kernel name A tuning file records solution indices, which are positions in one build's kernel library. The C API discarded the whole file when its "Git Version:" line did not match the running build, while the C++ extension API applied it without checking at all. Trust is now decided per entry on both APIs. hipblaslt-bench records a kernel_name beside each solution_index in HIPBLASLT_TUNING_FILE. At replay a named entry's index is resolved in the running library and used only if it still names the recorded kernel. An entry without a name is used only when the file was written by the running build. The header mismatch is now an info message rather than a switch that disables the file. 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. - A file with no usable rows was re-parsed on every heuristic query. - OverrideMap::find returned iterators after releasing its lock. - getSolutionNameFromData dereferenced a null solution for RocRoller indices. - The logger destructor could throw from close() at exit. Adds GPU tests (TuningCache.*, standard tier) behind a test-only reset hook that is compiled in only when the client tests are built. Co-authored-by: Cursor <cursoragent@cursor.com>
… entry When an override entry failed both the XF32 support check and the FP32 fallback check, the C++ extension path left the problem in FP32 mode. The remaining entries were then checked as FP32, and default selection picked FP32 kernels for an XF32 request. The C API already restores its compute type after the check. Adds TuningCache.ExtApiXf32ProblemStaysXf32AfterAnUnusableEntry, which fails without the restore. Co-authored-by: Cursor <cursoragent@cursor.com>
A grouped GEMM heuristic query crashed whenever HIPBLASLT_TUNING_OVERRIDE_FILE held an entry. problem_override_from_file_cpp built its key with TensileDataGemm2ProblemOverride, which reads gemmData as a TensileDataGemm, while a grouped GEMM carries a TensileDataGroupedGemm. Tuning file rows describe single GEMMs, so grouped GEMMs now skip the file and use default selection. Adds TuningCache.GroupedGemmUsesDefaultSelection, which segfaults without the guard. Co-authored-by: Cursor <cursoragent@cursor.com>
✅ 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. |
This was referenced Sep 25, 2026
jinchen62
requested review from
babakpst,
bethune-bryant,
newling,
rithwik-amd,
talumbau and
tony-davis
September 25, 2026 06:56
talumbau
approved these changes
Sep 25, 2026
talumbau
left a comment
talumbau
left a comment
Collaborator
There was a problem hiding this comment.
Nice piece of work -- the problem statement is sharp, the stack is well sequenced, and the fix list is honest about what it's fixing. The overrideResults[0] reuse bug in particular is the kind of thing that would have quietly poisoned multi-entry files indefinitely. Approving with comments; nothing below is a blocker.
One thing I want to clear explicitly: the new else { problem->setF32XdlMathOp(rocisa::DataType::XFloat32); } in problem_override_from_file_cpp looks at first glance like it's missing a restore on the fallback-success path. I chased it and it's correct as written -- leaving the persistent problem in Float after a successful FP32 fallback is the established idiom at tensile_host.cpp:4702, :4752 and :5421, and it matters, because makeArgument re-resolves the algo against that same problem. Restoring XF32 there would hand an FP32 kernel an XF32-tagged problem. Flagging so the next reviewer doesn't re-litigate it.
Four inline comments, in rough priority order: the bad()-vs-fail() check on the loaded-path marker (one-liner, worth taking now); the empty-build-stamp gate failing open in no-git builds (fine as a follow-up, but I'd like it tracked); kernel_name written empty for clients that don't set print_solution_found; and the hardcoded TuningCache.* category entries.
A few smaller things not worth inlining:
getContractionProblemsFromFiletakes the globalm_guardbefore checkingisLoaded(path), and runs on every heuristic query whileenv_modeis set. The critical section after load is just astd::set::count, so it's a short serialization point rather than a query-length one -- but it's permanent and process-wide. A double-checkedisLoadedbefore takingm_guard(re-checked inside) would remove it. Worth noting the PR slightly widens this, since droppingenv_mode = falsemeans a stale-stamp file no longer turns the path off.override_path_compare_git_versionis now effect-free -- with theenv_modeassignment gone, its return value only selects between twolog_infostrings, yet it still opens the file and re-parses the version line on every C-API heuristic call, with different trimming than the loader. Two parsers of the same header line that can disagree. Either delete it or emit the message once from the loader, which already hasfileBuildStamp. (Thelog_error->log_infodowngrade is invisible at the default level, where both were silent, but it does lose the diagnostic atHIPBLASLT_LOG_LEVEL1-3.)hipblaslt_strip_custom_tuning_suffixcan't fire on the value it's applied to: the" (Custom tuning: ...)"suffix comes only fromgetSolutionNameFromData, whilekernel_namecomes fromgetKernelNameFromAlgo->getKernelNameFromAlgoIndex, which returnssolution->kernelNamebare -- on both the ext and non-ext writer paths. Fine as defence-in-depth, but the comment currently asserts the opposite about which accessor appends.- The grouped-GEMM early return is the right call -- the pre-PR path was live UB (
static_pointer_cast<TensileDataGemm>over aTensileDataGroupedGemm), so this removes a crash rather than a feature. But the writer still appends grouped winners toHIPBLASLT_TUNING_FILE, and since the parser ignores thegrouped_gemmcolumn those rows land under a key indistinguishable from a single GEMM of the same shape (reachable atgrouped_gemm == 1;>= 2self-rejects on the(m0,m1)spelling). Pre-existing and not worsened here, but suppressing the write would close the loop. Longer term, pushing the type check down intoTensileDataGemm2ProblemOverride-- which still casts ashared_ptr<void>with no tag -- would stop the next caller from reintroducing it. HIPBLASLT_ENABLE_TUNING_TEST_HOOKSmakes an exported symbol conditional on build config, while the directly analogoushipblaslt_debug_reloadis unconditional and consumed the same way bystreamk5_test.cpp. An out-of-tree client build against a packagedlibhipblaslt.sowould fail to linktuning_cache_test.cppbut not its neighbour. Also,reloadForTest()andresetForTest()are public and unconditionally compiled, so the#ifdefonly guards the C wrapper.- Given four of this PR's fixes each had to be written twice across
problem_override_from_fileand_cpp, and the XF32 restore now exists in only one of them, it may be worth factoring the shared skeleton before the rest of the stack lands on top of it.
A file that failed to open was marked loaded, because a failed open sets failbit rather than badbit, so a missing or unreadable override file was written off for the rest of the process. It is now looked for again on the next query, as before this change. A build made outside a git checkout has an empty build stamp, and a file with no "Git Version:" line matched it, so its unnamed rows were trusted. A build with no stamp now trusts no unnamed row. The version line is matched without its trailing space, and a missing HIPBLASLT_VERSION_TWEAK is a compile error rather than an empty stamp. override_path_compare_git_version no longer affected anything and parsed the version line separately from the loader. It is removed; the loader logs once per file instead, at error level when it ignores rows. Also: - A query checks isLoaded before taking the map's mutex, so a loaded file costs it only a read lock. - hipblaslt-bench records kernel_name whether or not print_solution_found is set, no longer records grouped GEMM winners, which read back as single GEMMs, and no longer strips a custom-tuning suffix that kernel names never carry. - hipblaslt_tuning_reset_for_test is exported unconditionally, like hipblaslt_debug_reload. - The fixture is named TuningCache_pre_checkin, so the existing *pre_checkin* pattern selects it and test_categories.yaml is unchanged. Adds TuningCache_pre_checkin.FileCreatedAfterTheFirstQueryIsRead and UnnamedEntryWithoutAVersionLineIsIgnored. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
Thanks for the review. The four inline comments are answered in their threads; replies to the smaller notes:
The matching changes are carried up the stack: #12586 to #12590 are rebased on 1bade4e. |
jinchen62
enabled auto-merge (squash)
September 26, 2026 04:29
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (48.82%) 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 #12585 +/- ##
===========================================
+ Coverage 70.51% 70.55% +0.03%
===========================================
Files 2932 2932
Lines 482388 482160 -228
Branches 69538 69493 -45
===========================================
+ Hits 340134 340140 +6
+ Misses 118405 118165 -240
- Partials 23849 23855 +6
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
newling
approved these changes
Sep 28, 2026
newling
left a comment
newling
left a comment
Contributor
There was a problem hiding this comment.
LGTM. An agent review suggests one very defensive improvement (see here) but it's really an edge case.
The 2 main things that this PR does in my mind are
- don't blindly accept a solution index (C++ API)
- don't blindly reject a solution index if the commit hash changes (C API)
This is good, but we should keep in mind that the current solution here is not perfect, because it now accepts an index if the kernel name is correct, but some runtime settings are not captured in the kernel name, these are summarized here: https://github.com/newling/agent-public-reviews/blob/main/rocm-libraries/review_pr12585.md#enumerating-solution-fields-is-a-smaller-alternative
But this PR is a step in the right direction IMO, so I'm happy for it to land as is!
…rnels GroupedGemmUsesDefaultSelection asserted that a grouped GEMM heuristic query finds a solution before the test writes its tuning file. The library logic for gfx11 and gfx12 has no grouped GEMM kernels, so there the query finds none and the test failed. With no default selection to compare against, the test now skips. Co-authored-by: Cursor <cursoragent@cursor.com>
…header The loader paired a row's cells with its header only up to the shorter of the two. A row cut short by an interrupted append lost its kernel_name cell and read as a row that records no name, so in a file stamped by the running build its index was trusted without the name check, and a cut inside the index trusted a truncated one. A row whose cell count differs from its header's is now skipped. Adds TuningCache_pre_checkin.TruncatedNamedEntryIsIgnored, which cuts the row just before and just after the comma that precedes the name. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
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>
1 task
jinchen62
added a commit
that referenced
this pull request
Oct 9, 2026
…13354) JIRA ID : ROCM-32394 ## Motivation hipBLASLt multiplies matrices on AMD GPUs. A tuning file can tell it which kernel, or GPU function, to run for a given problem. The `TuningCache_pre_checkin` tests, added in #12585, write such a file and check which kernel the library picks. They wrote the file into the current directory and never checked that the write worked. The 10.2.0-20261007 nightly validation runs the installed `hipblaslt-test` from `/opt/rocm-10.2.0/bin`, which the test user cannot write, so no file was created. The seven tests that expect the recorded kernel got the default one and failed on all six reported systems. The seven tests that expect the file to be ignored passed without a file to ignore. The 10.2.0-20261006 nightly passed because it was built from a rocm-libraries commit that predates #12585. ## Technical Details - `tempTuningPath` now puts `hipblaslt_<test name>_<process id>.tuning` in GoogleTest's temporary directory, `::testing::TempDir()`, as `tuning_store_test.cpp` already does. On Linux that is normally `/tmp/`. - `writeRawTuningFile` and `writeTuningFile` return a GoogleTest assertion result. The result is a failure that names the path when the file cannot be opened, or when writing or closing it fails. - All 14 tests check that result with `ASSERT_TRUE`, so a failed write stops the test at the write instead of at a later check on the selected kernel. - Only `clients/tests/src/tuning_cache_test.cpp` changes. The library is unchanged. Not addressed here: `tuning_tune_test.cpp`, added by #12588, also names its tuning files relative to the current directory and needs the same change. ## Test Plan On one AMD Instinct MI355X (gfx950), build `hipblaslt-test` from `develop` (`faea76a5cde`) with and without this change. Run `--gtest_filter='TuningCache_pre_checkin.*'` from `/tmp` and from `/opt/rocm/bin`, which the test user cannot write. To check the new error message, make the write fail by placing a non-empty directory at the tuning file's path. ## Test Result | Test file | From `/tmp` | From `/opt/rocm/bin` | |---|---|---| | `develop` | 14 pass | 7 pass, 7 fail at lines 512, 547, 562, 579, 661, 724 and 745 | | This change | 14 pass | 14 pass | The seven failures on `develop` are the same cases, at the same lines, as the nightly report. When the write is forced to fail, `NamedEntryReplays` and `UnnamedEntryFromAnotherBuildIsIgnored` now fail at the write with `could not open tuning file <path>`. On `develop`, the second test passes when no file is written. The automated checks run these tests from a writable directory, which is why they passed on #12585. They cannot reproduce the nightly failure, and they are pending for this change. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test-only change. The library and the kernel it selects for applications are unchanged. The tests now need a writable temporary directory instead of a writable current directory. Co-authored-by: Cursor <cursoragent@cursor.com>
jinchen62
added a commit
that referenced
this pull request
Oct 13, 2026
…=cache) (#12587) 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 1. #12585 (merged) 2. #12586 (merged) 3. **`HIPBLASLT_TUNING_MODE=cache` (this PR)** 4. #12588, which also builds on #13568. That PR moves `hipblaslt-bench`'s candidate search into the library, so tune mode and the bench share one implementation. 5. #12589 6. #12590 (draft) 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>
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>
jinchen62
added a commit
that referenced
this pull request
Oct 26, 2026
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>
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
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:Git Version:line differs from the running build, so one differing byte discards every entry.Technical details
hipblaslt-benchwrites akernel_namecolumn next tosolution_index, whether or notprint_solution_foundis 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.getKernelNameFromAlgoIndexstill returns the recorded name. Asolution_namecolumn is checked the same way. A row with neither is used only when the file'sGit 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_versionis 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.[0], so later entries for a problem were checked against the first entry's solution.0was rejected.*returnAlgoCountwas read uninitialised when the file served a single-algo request.OverrideMap::findreturned iterators after releasing its lock (latent today, since the map is not written after loading).getSolutionNameFromDatadereferenced a null solution for RocRoller indices.close()at exit.hipblaslt_tuning_reset_for_test()is exported unconditionally, likehipblaslt_debug_reload().Test plan
TuningCache_pre_checkin.*(13 cases). The fixture name puts them in the existing*pre_checkin*category, standard tier and above, sotest_categories.yamlis unchanged.[0]bug;hipblaslt-benchtuning run writes the file, and the override replays the tuned index. The same file with another build's version line still replays, and a tamperedkernel_nameis rejected and falls back to the heuristic. A grouped run writes no row, and a run withprint_solution_foundoff still recordskernel_name.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.0acceptance, single parse of an empty file, RocRollergetSolutionName(), logger destructor.Rest of the stack
PRs 5 and 6 both build on PR 4 and are independent of each other.