Skip to content

feat(hipblaslt-bench): tune through the library's tune mode - #12590

Closed
jinchen62 wants to merge 23 commits into
ROCm:developfrom
jinchen62:tuning-split/6-bench-tune-mode
Closed

jinchen62 wants to merge 23 commits into
ROCm:developfrom
jinchen62:tuning-split/6-bench-tune-mode

Conversation

@jinchen62

@jinchen62 jinchen62 commented Sep 25, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID: AIHPBLAS-3751

Stacked on #12588 (then #12587, #12586, #12585), and independent of #12589. Until those merge, this PR also shows their commits. Review only the changes on top of PR 4: Compare against PR 4

Motivation

Last of six PRs splitting #10424. hipblaslt-bench with HIPBLASLT_TUNING_FILE had its own tuning loop, which timed candidates differently from the library's tune mode. The bench and a tuning application could therefore pick different winners for the same problem. The bench now tunes through the library's tune mode, so there is one tuner. Users run the same commands. The bench now prints the winner once, and the file gets the rows tune mode writes.

Technical details

  • Setup. Before any hipBLASLt call, main points tune mode at the tuning file (HIPBLASLT_TUNING_MODE=tune, HIPBLASLT_TUNING_CACHE_PATH=<file>) and maps the bench's options to the tune-mode settings:

    • --cold_iters and --iters become the warm-up and timed launches;
    • --rotating becomes the rotation size;
    • --flush becomes the instruction-cache flush;
    • --requested_solution -1 means every kernel, and N means the top N ranked kernels.

    The per-shape budget is off unless HIPBLASLT_TUNING_BUDGET_MS_PER_SHAPE is set, since offline tuning is where a search is meant to finish. A conflicting HIPBLASLT_TUNING_MODE or HIPBLASLT_TUNING_CACHE_PATH is refused with an error.

  • The tuning pass. One untimed call through the C API tunes and records the winner. During that call, tuning_pass() hands the library the whole workspace limit rather than what the heuristic's pick needs, so the search covers every kernel that fits. The timed run afterwards is the winner alone, which the heuristic now returns like any cached problem. Its solution index is printed.

  • Removed: the bench's own row writer, its archName/CUs plumbing, the git-version preflight on the tuning file, and PR 1's kernel_name column. The file now gets the versioned rows the runtime cache writes, and those carry kernel_name themselves. HIPBLASLT_TUNING_OVERRIDE_FILE and cache mode read them.

  • Not tuned: grouped GEMM and pointer-array batched problems, with a message saying so.

  • QuickTune parses the single timed result for both baseline and tuning runs.

  • Docs: the bench section of the offline tuning how-to (what the options now control, and the new sample output), the runtime tuning intro and row-format paragraph, and HIPBLASLT_TUNING_FILE in the environment reference.

Interaction with PR 5. The two merge without conflicts. Whichever lands second should change one how-to sentence: "A problem the file already holds is not tuned again" becomes "is tuned again only by a run that searches more than the one that wrote it". With PR 5, a bench run with more candidates or iterations re-tunes the shape.

Test plan

  • hipblaslt-bench end to end on MI300X (1024x512x1024 fp16, --requested_solution 16 -i 20 -j 5):
    • the first run prints tuning-start, tuning-done ... persisted=yes, the timed winner with its solution index, and summary ... tuned=1, and the file gets one row;
    • an identical second run does not tune again;
    • HIPBLASLT_TUNING_OVERRIDE_FILE replays that index on both the C and the C++ API;
    • cache mode replays it;
    • HIPBLASLT_TUNING_FILE together with HIPBLASLT_TUNING_MODE=cache is refused with exit code 1.
  • testing_matmul.hpp drives the whole matmul test suite, so the smoke tier covers the non-tuning paths.
  • The combined PR 5 + PR 6 state was also built and tested locally:
    • TuningStore.* 31/31, and TuningCache_pre_checkin.*/TuningTune_pre_checkin.* 59 pass with the two-GPU case skipped on one visible GPU; that case passes with both GPUs visible;
    • the bench writes version 2 rows, skips an identical rerun, and re-tunes when --requested_solution grows from 16 to 32.

Test result

The GPU suites run in the standard tier and above, not in PR smoke gating, so these are local results on MI300X (gfx942).

  • Smoke: 3133/3133 pass. Every target builds.
  • TuningStore.* 21/21 and TuningCache_pre_checkin.*/TuningTune_pre_checkin.* 54 pass, with the two-GPU case skipped on the single visible GPU.

Rest of the stack

  1. fix(hipblaslt): repair the tuning override path and validate entries by kernel name #12585
  2. feat(hipblaslt): add a versioned tuning cache store with host tests #12586
  3. feat(hipblaslt): replay a managed tuning cache (HIPBLASLT_TUNING_MODE=cache) #12587
  4. feat(hipblaslt): tune unseen shapes at their first matmul (tune mode) #12588
  5. feat(hipblaslt): keep and revisit budget-stopped searches #12589
  6. hipblaslt-bench tunes through the library's tune mode (this PR)

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

@therock-pr-bot

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

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

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

🙋 Wish to Override Policy?

@therock-pr-bot

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

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@jinchen62
jinchen62 force-pushed the tuning-split/6-bench-tune-mode branch from d4eb9e1 to bd3d791 Compare September 26, 2026 04:09
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>

@newling newling left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Here is an agent review: review. It has 2 actionable items and 1 suggestion. Please take a look.

@jinchen62
jinchen62 force-pushed the tuning-split/6-bench-tune-mode branch from bd3d791 to 95a7e5f Compare October 7, 2026 17:01
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>
jinchen62 and others added 5 commits October 8, 2026 14:02
…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>
newling and others added 7 commits October 8, 2026 14:02
The extended profile named the kernel from algo, which on a call the tuner
served still holds default selection's index: a fresh winner, or one
another thread had just tuned, was logged under its own index with the
default kernel's names. Name it from the index launched, as cache replay's
profile fix did before replay began filling algo itself.

Co-authored-by: Cursor <cursoragent@cursor.com>
The scratch plan decides how much memory a candidate can touch, and the
tune tests reached it only through contiguous single-batch problems.
Expose it and tensorSpanBytes through test hooks, with the rotating budget
and the cap passed in rather than read from the environment, and check
the numbers directly: padded and overlapping strides, broadcast inputs
that must not be copied for rotation, unequal and in-place C and D,
gradient-bias spans, trimming to fit the cap, the bound on rotation
blocks, and arithmetic overflow.

Co-authored-by: Cursor <cursoragent@cursor.com>
hipblaslt-bench with HIPBLASLT_TUNING_FILE no longer times every candidate
itself. It points the library's tune mode at that file and runs one untimed
call, so the bench and a tuning application measure candidates the same way
and cannot choose different winners; the timed run is then the winner alone.

- --iters, --cold_iters, --requested_solution and --rotating become the
  tune-mode settings, the instruction cache is flushed between launches,
  and the search has no time limit unless one is set.
- The tuning pass hands the library the whole workspace limit, so the search
  covers every kernel that fits.
- The file gets the versioned rows the runtime cache writes, which the
  override and cache modes read; the bench's own writer, including the
  kernel_name column added earlier, is removed.
- HIPBLASLT_TUNING_MODE or HIPBLASLT_TUNING_CACHE_PATH set to anything else
  is refused.
- QuickTune parses the single timed result.

Co-authored-by: Cursor <cursoragent@cursor.com>
@jinchen62
jinchen62 force-pushed the tuning-split/6-bench-tune-mode branch from 95a7e5f to 130c670 Compare October 8, 2026 19:11
@codecov

codecov Bot commented Oct 8, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 75.36601% with 387 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp 74.23% 150 Missing and 93 partials ⚠️
...md_detail/rocblaslt/src/UserDrivenTuningParser.cpp 72.64% 54 Missing and 36 partials ⚠️
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp 81.90% 27 Missing and 11 partials ⚠️
...y/src/amd_detail/rocblaslt/src/include/utility.hpp 58.82% 5 Missing and 2 partials ⚠️
...l/rocblaslt/src/include/UserDrivenTuningParser.hpp 81.25% 6 Missing ⚠️
.../src/amd_detail/rocblaslt/src/TuningCacheStore.cpp 89.47% 0 Missing and 2 partials ⚠️
..._detail/rocblaslt/src/include/TuningCacheStore.hpp 94.12% 0 Missing and 1 partial ⚠️
Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #12590      +/-   ##
===========================================
+ Coverage    71.98%   72.23%   +0.25%     
===========================================
  Files         2879     2874       -5     
  Lines       478113   476740    -1373     
  Branches     70856    70467     -389     
===========================================
+ Hits        344166   344347     +181     
+ Misses      109976   108399    -1577     
- Partials     23971    23994      +23     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 49.33% <ø> (-0.35%) ⬇️ Carriedforward from a0b7b28
TensileLite-Unit 83.33% <ø> (-0.22%) ⬇️ Carriedforward from a0b7b28
hipBLAS 90.62% <ø> (ø) Carriedforward from a0b7b28
hipBLASLt 45.63% <75.37%> (+2.90%) ⬆️
hipCUB 82.68% <ø> (ø) Carriedforward from a0b7b28
hipDNN 86.66% <ø> (-0.15%) ⬇️ Carriedforward from a0b7b28
hipFFT 42.52% <ø> (+0.16%) ⬆️ Carriedforward from a0b7b28
hipRAND 76.12% <ø> (ø) Carriedforward from a0b7b28
hipSOLVER 68.96% <ø> (ø) Carriedforward from a0b7b28
hipSPARSE 86.99% <ø> (ø) Carriedforward from a0b7b28
rocBLAS 48.17% <ø> (-0.04%) ⬇️ Carriedforward from a0b7b28
rocFFT 51.43% <ø> (+4.18%) ⬆️ Carriedforward from a0b7b28
rocRAND 57.45% <ø> (-0.14%) ⬇️ Carriedforward from a0b7b28
rocSOLVER 76.83% <ø> (ø) Carriedforward from a0b7b28
rocSPARSE 74.07% <ø> (-0.23%) ⬇️ Carriedforward from a0b7b28
rocThrust 93.15% <ø> (ø) Carriedforward from a0b7b28

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

Files with missing lines Coverage Δ
...slt/clients/common/src/hipblaslt_bench_options.cpp 29.63% <100.00%> (+5.63%) ⬆️
.../amd_detail/rocblaslt/src/include/tensile_host.hpp 28.57% <ø> (ø)
..._detail/rocblaslt/src/include/TuningCacheStore.hpp 95.03% <94.12%> (+10.31%) ⬆️
.../src/amd_detail/rocblaslt/src/TuningCacheStore.cpp 76.68% <89.47%> (+36.86%) ⬆️
...l/rocblaslt/src/include/UserDrivenTuningParser.hpp 88.68% <81.25%> (-11.32%) ⬇️
...y/src/amd_detail/rocblaslt/src/include/utility.hpp 49.82% <58.82%> (+0.58%) ⬆️
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp 76.98% <81.90%> (+0.32%) ⬆️
...md_detail/rocblaslt/src/UserDrivenTuningParser.cpp 73.33% <72.64%> (-7.31%) ⬇️
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp 50.38% <74.23%> (+5.46%) ⬆️

... and 503 files with indirect coverage changes

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

jinchen62 and others added 11 commits October 12, 2026 00:57
The scratch plan gave every rotation block its own copy of the caller's
workspace limit. hipblaslt-bench hands the library its whole 128 MiB
limit, so 512 MiB of rotation held only two or three blocks, and
candidates were timed with their inputs still in the last-level cache.
The bench shares one workspace across its blocks and rotates only the
tensors. On gfx950 the difference made the tuner's winner up to 25%
slower than the bench's when the bench timed both, and at 256x256x256
slower than the untuned default.

Every block now shares one workspace, so the rotating budget goes to
the tensors. The plan tests only ever passed a zero workspace; a new
one plans with 128 MiB and expects the same blocks.

Co-authored-by: Cursor <cursoragent@cursor.com>
…ache-mode

Co-authored-by: Cursor <cursoragent@cursor.com>

# Conflicts:
#	projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp
ROCm#13354 made the tuning file writers return a GoogleTest assertion
result and checked it in every existing test. The cache-mode tests
added here still ignored it, so a write that failed surfaced later as
the wrong kernel. They now stop at the write, as the others do.

Co-authored-by: Cursor <cursoragent@cursor.com>
ROCm#13354 moved the cache-mode tests' tuning files into GoogleTest's
temporary directory and checked every write, because the installed
hipblaslt-test runs from a directory the test user cannot write. It
noted that the tune-mode tests needed the same change. Their files now
go to ::testing::TempDir() too, and the two places that rewrite a file
check that the write succeeded.

Co-authored-by: Cursor <cursoragent@cursor.com>

@Alex-Vasile Alex-Vasile left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approving CMakeList changes

@jinchen62
jinchen62 marked this pull request as draft October 13, 2026 19:47
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

Copy link
Copy Markdown
Contributor Author

Closing: #13568 replaces this. It moves hipblaslt-bench's candidate search into the library, and #12588's tune mode now runs that same search, so the bench and the library share one tuning implementation without routing the bench through tune mode.

@jinchen62 jinchen62 closed this Oct 13, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants