feat(hipblaslt): add adaptive sm_count_target mode by caio96 · Pull Request #13263 · ROCm/rocm-libraries

caio96 · 2026-10-07T17:42:47Z

JIRA ID : N/A

Stack: #13692 (kernel probe) → #13892 (estimator) → this PR (adaptive launches) → #13893 (re-selection). Review in that order.

Motivation

sm_count_target helps when part of the GPU is busy with another workload, but users rarely know what value to pass, and the right value changes over time. This PR measures how many CUs GEMMs can actually use (the probe from #13692, the estimator from #13892) and applies that as the hint automatically.

Technical Details

New opt-in mode: HIPBLASLT_ADAPTIVE_SM_COUNT=1, with StreamK tile scheduling AUTO or ON and no explicit sm_count_target. The mode applies whatever the device or library; the hint only moves when the launched kernels carry the probe (gfx950, probe-built library).
Every 8th launch on a stream probes the GPU; the host reads finished probes and keeps a per-stream hint.
The hint is used like a user-given sm_count_target: the heuristic query ranks under it, and each launch uses it for the StreamK hybrid-assignment grid and static/dynamic mode. Swapping the solution at launch is feat(hipblaslt): re-select the top solution when the adaptive hint changes #13893.
The top heuristic result reports a workspace large enough for a regridded launch (see Appendix).
Host side, adaptive_sm::Controller (one per handle, null with the mode off) owns the probe slots in host-mapped memory and, per stream, a lock, a hipEvent, the feat(hipblaslt): add adaptive sm_count_target estimator #13892 estimator and the published hint. _FORCE is applied there only.
Each adaptive launch takes a Launch from Controller::begin() (hint, probe target). Its destructor records the probe's event, so no return path can leave a probe reserved but unrecorded.
TensileLite marks the kernel invocations that carry the probe (KernelInvocation::occupancyProbe), so the host learns from TensileLite whether a launch probed rather than re-deriving its padding rules.
A launch on a capturing stream is not adaptive: it neither probes nor takes the hint (not even _FORCE) and leaves the stream's state alone, since a graph replays the captured kernel arguments.
Probe events are created with the controller; if that fails, the handle runs without adaptive mode and logs why.
Timing tools (hipblaslt-bench, -sequence, -api-overhead) refuse unforced adaptive mode; _FORCE / _FORCE_PROBE give fixed, measurable behaviour.
Docs: Stream-K how-to and environment variable reference; CHANGELOG entry (the mode, the workspace raise, the timing-client guard).

Test Plan

hipblaslt-test-hostunit: the benchmark timing guard (estimator tests are in feat(hipblaslt): add adaptive sm_count_target estimator #13892).
tensilelite-tests (PersistentArgumentLayout.*): the probe flag on kernel invocations: set exactly when the probe address is packed; unset for grouped GEMM and for a multi-dimensional grid too small to pad.
hipblaslt-test GPU tests: query workspace, forced mode (AUTO and ON), a probe producing a reading, a cotenant holding n CUs (the reading is at most CUs − n), graph capture (a captured launch leaves the stream's state alone), the stream lock held through a probe launch and released after a non-probe launch. Mode-specific tests skip when their mode is not set, so CI runs them with the mode off.
Manually: the adaptive, SK and tuning suites with HIPBLASLT_ADAPTIVE_SM_COUNT=1, _FORCE=128 and _FORCE=128 _FORCE_PROBE=1.
Manual, opt-in (HIPBLASLT_TEST_COTENANT_PROBE=1, needs an otherwise idle GPU since other processes taking CUs change the count): the reading beside a cotenant equals exactly CUs − n.
MI355X with hipblaslt-cotenant: correctness (-v, including batched GEMM) with and without contention, a driver where the cotenant appears and leaves, forced-mode bench runs.

Test Result

All suites pass; mode off leaves behaviour unchanged.
Under contention the hint converges to the free-CU count, follows the cotenant appearing and leaving, and results match the baseline.
Forced mode times the same kernel as an explicit sm_count_target.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Risk level

Low. Opt-in at runtime (and needs a probe-built library); with the mode off, behaviour is unchanged.

Appendix

Block flow (one `hipblasLtMatmul` call)

hipblasLtMatmul ── adaptive mode applies? (C API, StreamK AUTO/ON,
   │               no explicit sm_count_target, stream not capturing)
   │                                   no ──► normal launch, no hint, no probe
   ▼ yes
┌─ Controller::begin(): before the launch (per-stream state) ──────────────────┐
│ 1. Count the launch. A probe is "due" from the first launch, and again at    │
│    every PROBE_EVERY-th (8) launch, a tick.                                  │
│ 2. Poll this stream's pending probe (at most 1, one hipEvent):               │
│      finished → count CUs whose slot holds its epoch → reading               │
│      readings → new hint? (rules in #13892; PROBE_WINDOW = 3 readings)       │
│ 3. Probe due? Take it unless the GEMM is tiny (< TINY_MFLOPS) or the         │
│    previous probe is not read yet (then it waits; that is not a miss).       │
│    A tick that finds the probe still due (not taken since the previous       │
│    tick, e.g. only tiny GEMMs) and none pending is a miss;                   │
│    2 misses in a row → hint 0.                                               │
│    A probing Launch holds the stream's lock until it is destroyed,           │
│    so probes from several threads reach the GPU in epoch order.              │
└──────────────────────────────────────────────────────────────────────────────┘
   ▼
┌─ Kernel launch ──────────────────────────────────────────────────────────────┐
│ • sm_count_target = hint → StreamK grid size and static/dynamic mode         │
│   (the solution itself is the caller's; re-selection is #13893)              │
│ • Probe launch: kernel gets ProbeAddr/ProbeEpoch; probe-only workgroups      │
│   pad the grid to all CUs and exit after marking                             │
│ • TensileLite reports whether a kernel carried the probe → Launch::probed    │
│   (only gfx950 StreamK hybrid kernels from a probe-built library do;         │
│   if none do, the hint stays 0)                                              │
└──────────────────────────────────────────────────────────────────────────────┘
   ▼
┌─ ~Launch(): after a probe launch only ───────────────────────────────────────┐
│ • Record the probe's hipEvent on the stream (read in a later begin(), once   │
│   the GPU has run it)                                                        │
│ • Kernel did not probe, or the record failed → the probe stays due (a miss   │
│   if the next tick still finds it due)                                       │
└──────────────────────────────────────────────────────────────────────────────┘

Timeline (one stream, a cotenant holds 64 of 256 CUs)

The host enqueues launches long before the GPU runs them, so a probe's result is read many launches after it was launched. PROBE_EVERY = 8, PROBE_WINDOW = 3.

host launch   on the host                                  pending probe / GPU
───────────   ──────────────────────────────────────────   ─────────────────────────
L1            first probe due → probe A launched, event      [A]
              recorded; launches use hint 0 (whole GPU)
L2 … L7       poll: A not finished yet
L8            tick: probe due, A not read yet → waits       [A]
                                                            GPU runs L1: 192 CUs mark A
L13           poll: A done → reading 192 (1 of 3)            [ ]
              due probe taken → probe B                     [B]
                                                            GPU runs L13: 192 CUs mark B
L22           poll: B done → reading 192 (2 of 3)            [ ]
              due probe taken → probe C                     [C]
                                                            GPU runs L22: 192 CUs mark C
L31           poll: C done → reading 192 (3 of 3)            [ ]
              → hint 192; due probe taken → probe D         [D]
L31 …         launches use hint 192: StreamK grid and mode
              for 192 CUs

cotenant leaves → the next probe (padded to all 256 CUs) reads 256 → hint 0 at once

At most one probe per stream is pending, so no newer probe can overwrite its slots before it is read: each reading is exactly one probe.
A hint change needs PROBE_WINDOW readings on the same side of the hint, so one odd reading does not move it. The exception is a reading of the full CU count, which clears the hint at once.

Workspace

At launch the hint can change the StreamK grid of the queried solution (and, with feat(hipblaslt): re-select the top solution when the adaptive hint changes #13893, the solution itself), so the workspace must cover what the hint could lead to.
The top heuristic result therefore reports min(MAX_WORKSPACE_BYTES, max(required, 128 MiB)). 128 MiB matches origami's own partial-tile limit for StreamK grids; a launch that needs more falls back to a slower path, never fails.
Callers that pass a fixed workspace buffer (most frameworks) are unaffected; set MAX_WORKSPACE_BYTES to your budget.

What adapts

Case	Adapts?
C API, gfx950, probe-built library, StreamK AUTO/ON	✅
Strided and general batched GEMM	✅
Explicit `sm_count_target` (desc / handle)	❌ explicit value wins
Ext API, grouped GEMM, rocRoller, graph capture	❌
Mode on, but not gfx950 or not a probe-built library	❌ hint stays 0 (workspace and timing rules still apply)
Cotenant that takes whole CUs	✅
Cotenant that shares CUs	❌ not visible to the probe

Knobs

See #13892 (prefix HIPBLASLT_ADAPTIVE_SM_COUNT: _PROBE_EVERY, _TINY_MFLOPS, _PROBE_WINDOW, _LOG, _FORCE, _FORCE_PROBE).

Timing rule

A timed solution must never change mid-measurement, so benchmark clients refuse unforced adaptive mode. _FORCE=<cus> times the solution adaptive mode picks at that hint; adding _FORCE_PROBE=1 includes the probe cost.

🤖 Generated with Claude Code

therock-pr-bot · 2026-10-07T17:47:31Z

❌ PR Check — Action Required

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	❌ Fail	Error: Check concluded with `cancelled`.
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

❌ pre-commit

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

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

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-07T17:47:34Z

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

codecov · 2026-10-07T20:29:21Z

Codecov Report

❌ Patch coverage is 13.26531% with 255 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...rocblaslt/src/rocblaslt_adaptive_sm_controller.cpp	5.53%	185 Missing and 3 partials ⚠️
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp	0.00%	34 Missing and 2 partials ⚠️
...t/src/include/rocblaslt_adaptive_sm_controller.hpp	0.00%	10 Missing ⚠️
...ary/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp	23.08%	9 Missing and 1 partial ⚠️
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp	40.00%	6 Missing and 3 partials ⚠️
...rary/src/amd_detail/rocblaslt/src/include/handle.h	89.47%	0 Missing and 2 partials ⚠️

❌ Your project check has failed because the head coverage (65.71%) 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                            @@
##           users/csalvado/adaptive-sm-estimator   #13263      +/-   ##
========================================================================
- Coverage                                 76.12%   76.08%   -0.04%     
========================================================================
  Files                                      2887     2889       +2     
  Lines                                    484232   484520     +288     
  Branches                                  73039    73093      +54     
========================================================================
+ Hits                                     368591   368627      +36     
- Misses                                    90401    90642     +241     
- Partials                                  25240    25251      +11

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`50.80% <ø> (-<0.01%)`	⬇️
TensileLite-Unit	`84.07% <ø> (ø)`
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from e90c3ba
hipBLASLt	`65.71% <13.27%> (-0.73%)`	⬇️
hipCUB	`82.68% <ø> (ø)`		Carriedforward from e90c3ba
hipDNN	`86.85% <ø> (ø)`		Carriedforward from e90c3ba
hipFFT	`38.02% <ø> (ø)`		Carriedforward from e90c3ba
hipRAND	`76.12% <ø> (ø)`		Carriedforward from e90c3ba
hipSOLVER	`69.67% <ø> (ø)`		Carriedforward from e90c3ba
hipSPARSE	`87.03% <ø> (ø)`		Carriedforward from e90c3ba
rocBLAS	`68.73% <ø> (ø)`		Carriedforward from e90c3ba
rocFFT	`47.34% <ø> (ø)`		Carriedforward from e90c3ba
rocRAND	`57.31% <ø> (ø)`		Carriedforward from e90c3ba
rocSOLVER	`76.74% <ø> (ø)`		Carriedforward from e90c3ba
rocSPARSE	`77.16% <ø> (ø)`		Carriedforward from e90c3ba
rocThrust	`93.15% <ø> (ø)`		Carriedforward from e90c3ba

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

Files with missing lines	Coverage Δ
...lt/library/src/amd_detail/rocblaslt/src/handle.cpp	`56.35% <100.00%> (+0.35%)`	⬆️
...blaslt/src/include/rocblaslt_adaptive_sm_count.hpp	`93.78% <ø> (+1.47%)`	⬆️
...rary/src/amd_detail/rocblaslt/src/include/handle.h	`78.38% <89.47%> (+0.83%)`	⬆️
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp	`58.78% <40.00%> (-0.06%)`	⬇️
...t/src/include/rocblaslt_adaptive_sm_controller.hpp	`0.00% <0.00%> (ø)`
...ary/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp	`89.59% <23.08%> (-0.69%)`	⬇️
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp	`81.13% <0.00%> (-1.04%)`	⬇️
...rocblaslt/src/rocblaslt_adaptive_sm_controller.cpp	`5.53% <5.53%> (ø)`

... and 1 file 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.

Alex-Vasile

Approving CMakeLists changes

With HIPBLASLT_ADAPTIVE_SM_COUNT set, a co-tenant that takes whole CUs otherwise leaves SK5 grids sized for the full device. Every few C API launches on a stream now carry the occupancy probe; its reading feeds the stream's estimator, whose hint becomes the implicit sm_count_target for the heuristic query and for launches. The top query result reports enough workspace for a regrid under a later hint. A per-handle controller owns the probe buffer, the per-stream estimators and the published hints. Each launch holds its adaptive state in an RAII object whose end records the probe's completion event, and TensileLite marks the kernel invocations that carry the probe. The probe events are created with the controller; if creating one fails, the handle runs without adaptive mode and logs why.

A varying adaptive hint can change the timed kernel mid-run, so the timing clients exit unless HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE fixes it.

…rnels (#13692) JIRA ID : N/A ## Motivation hipBLASLt can steer kernel choice with an `sm_count_target` hint when part of the GPU is busy, but nothing measures how many CUs are actually free. This PR adds the measuring part: an opt-in probe in StreamK kernels with hybrid work assignment that records which CUs a launch ran on. It is the first PR of a stack: #13892 (estimator), #13263 (adaptive launches) and #13893 (re-selection) use it to set the hint automatically. ## Technical Details - New codegen option `--emit-occupancy-probe` (`TENSILELITE_EMIT_OCCUPANCY_PROBE=ON`, `invoke build --emit-occupancy-probe`). Off by default. - With it on, StreamK hybrid-assignment kernels on gfx950 get three extra kernel arguments at the end (`ProbeAddr`, `ProbeEpoch`, `ProbeGrid`) and a few instructions at the start that write the launch's epoch into the slot of the CU they run on. - The host can pad a launch with extra "probe-only" workgroups so it sees the whole GPU; those exit right after marking. - Nothing calls this yet: without a caller every probe address is null and kernels behave as before. - CHANGELOG entry for the build option. ## Test Plan - TensileLite unit tests: gate, probe emission, argument layout, end-to-end codegen (metadata offsets match the probe loads; nothing else stores before the probe-only exit), library metadata. - C++ `tensilelite-tests`: argument packing and padding (batched, 2-D grid, grouped GEMM), serialization. - Flag-off build compared against develop. - MI355X, with the rest of the stack as the caller: probe readings with and without `hipblaslt-cotenant`; `-v` correctness, including strided and general batched GEMM. ## Test Result - All tests pass. - With the flag off, generated kernels and library files are byte-identical to develop. - With the flag on, register usage of every kernel is unchanged. - Probe readings match the number of free CUs; results are correct with the probe on every launch. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. Opt-in at build time; off by default with identical output, and inert without a host caller. --- ## Appendix ### What the probe writes ``` host-visible buffer (512 × u32 per stream) slot = XCC << 6 | SE << 4 | CU (unique per CU on gfx950) 512 slots = 8 XCC × 4 SE × 16 CU ids; MI355X uses 256 (CU ids 0–8 per SE) launch with epoch 1042: buffer after the launch WG on XCC0 SE0 CU3 ──► slot 3 = 1042 WG on XCC0 SE1 CU0 ──► slot 16 = 1042 WG on XCC5 SE2 CU7 ──► slot 359 = 1042 CU busy with another app → slot untouched (older epoch) ``` Count of slots holding the epoch = CUs this launch could use. ### Kernel prologue (every workgroup; thread 0 stores) ``` load ProbeAddr, ProbeEpoch, ProbeGrid read CU id (HW_ID, XCC_ID) → slot (while the loads are in flight) wait for the loads ProbeAddr == 0 ? ──yes──► continue as normal kernel │ no thread 0 stores epoch to ProbeAddr[slot] raw workgroup id ≥ ProbeGrid ? ──yes──► exit (probe-only workgroup) │ no continue as normal kernel ``` ### Probe-only workgroups (padding) ``` real grid = 224 WGs (< 256 CUs) dispatch = 256 WGs ┌──────────── 224 real WGs ───────────┐┌─ 32 probe-only ─┐ do the GEMM mark CU, exit ``` - Only the dispatch size grows; every grid-derived kernel argument (`numWorkGroups`, `SKGrid`, workgroup mapping) keeps the real grid. - Probe-only workgroups exit before touching memory, flags or the work queue, using the raw (pre-remap) workgroup id. ### Notes - Gated to StreamK hybrid-assignment kernels on gfx950, non-grouped, non-custom; gfx942 is off until its CU layout is verified on hardware. - CUs whose id falls outside the 512-slot layout skip the mark instead of aliasing. - A probe-built library needs a runtime that passes the probe arguments (this PR's host code does). 🤖 Generated with [Claude Code](https://claude.com/claude-code)

JIRA ID : N/A **Stack:** #13692 (kernel probe) → **this PR** (estimator) → #13263 (adaptive launches) → #13893 (re-selection). Review in that order. ## Motivation The probe from #13692 tells the host which CUs a launch ran on. This PR adds the policy that turns those readings into a per-stream CU hint: when to probe, how to read a probe, and when the hint changes. It is pure logic (no HIP calls), so it can be reviewed and tested on its own; #13263 wires it into launches. ## Technical Details - New header `rocblaslt_adaptive_sm_count.hpp` (HIP-free), namespace `rocblaslt::adaptive_sm`: - `Config`: the `HIPBLASLT_ADAPTIVE_SM_COUNT*` environment variables, read once, invalid values warn and keep the default. It also answers the two questions `_FORCE` raises for callers: `hintVaries()` (can timing see the hint change) and `tracksLaunches()` (do launches update the estimator: not under `_FORCE` unless `_FORCE_PROBE`). - `Estimator`: one stream's state over its probe slots. `beforeLaunch` reads a finished probe and returns the epoch to probe with (or none); `afterProbe` reports whether that probe was recorded. It does not know about `_FORCE`; the caller in #13263 substitutes the forced hint. - `applies` (does the adaptive hint apply to a matmul) and `tinyProblem` (too small to probe). - `countFrom`: CUs a probe marked; `ProbeEpoch`: probe epochs and the periodic slot clear. - `ReadingWindow` and `candidate`: the last readings, and the hint they give. - `normalizeEnv` moves to `rocblaslt_secure_env.hpp` so the header and `handle.cpp` share it. - Nothing calls the estimator yet; behaviour is unchanged. ## Test Plan - `hipblaslt-test-hostunit` (`HostUnitAdaptiveSmCount*`): the estimator driven launch by launch over a hand-filled slot array (probe ticks, misses and reset, a due probe waiting for the previous one, failed events, readings to hint, a full reading clearing the hint at once), epoch wrap and clear, environment parsing and the `Config` FORCE rules. ## Test Result - All host unit tests pass (runs on every CI tier, no GPU needed). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level None. New header and tests only; nothing calls it yet. --- ## Appendix ### Per-launch decisions (one stream) ``` every launch : pending probe finished? → reading → maybe new hint start, tick : a probe is due from the first launch, and again every PROBE_EVERY-th launch (a tick) due probe : taken by the first launch that can take it not tiny (< TINY_MFLOPS) and no probe pending a pending probe defers it (not a miss) missed : a tick that finds the probe still due (not taken since the previous tick) and no probe pending; 2 misses in a row → hint 0. A recorded probe restarts the count ``` At most one probe per stream is pending, so a newer probe never overwrites an unread one: each reading is exactly one probe. ### Reading a probe - Each probe has an epoch; its workgroups store it in the slot of the CU they run on. - Reading = slots holding that epoch or newer (signed 32-bit compare, so the wrap is safe); 0 means never written. - Every 2³⁰ epochs the stream's slots are zeroed before the next probe, so no stale mark can ever compare as newer. ### Readings → hint (window = last `PROBE_WINDOW` readings, default 3) | Window | Hint | |---|---| | a hint is set and one reading is at the full CU count | 0 at once; the window restarts from that reading | | all at the full CU count | 0 (no hint) | | all below the current hint | the smallest reading (exact) | | all above the current hint | the smallest reading (exact), 0 if it is the full count | | otherwise | unchanged | A reading of 0 is ignored. Requiring the whole window on one side filters one-off readings (e.g. a cotenant workgroup briefly preempted). ### Knobs (prefix `HIPBLASLT_ADAPTIVE_SM_COUNT`) | Var | Default | Meaning | |---|---|---| | (unsuffixed) | off | enable | | `_PROBE_EVERY` | 8 | launches between probes | | `_TINY_MFLOPS` | 256 | GEMMs below this never probe | | `_PROBE_WINDOW` | 3 | readings that must agree to change the hint | | `_LOG` | 0 | 1: config, hint changes, failed probe events, re-selections (#13893); 2: also readings. Timestamped lines to `HIPBLASLT_LOG_FILE` or stderr | | `_FORCE` | – | fixed hint (timing) | | `_FORCE_PROBE` | off | with `_FORCE`: probe but ignore readings | 🤖 Generated with [Claude Code](https://claude.com/claude-code)

assistant-librarian · 2026-10-30T04:11:34Z

TheRock Submodule Bump Activity

Newest first

2026-10-30 — Included in TheRock via ROCm/TheRock#9594.

@128

…anges (#13893) JIRA ID : AIHPBLAS-5323 **Stack:** #13692 (kernel probe) → #13892 (estimator) → #13263 (adaptive launches) → **this PR** (re-selection). ## Motivation With #13263 the adaptive hint ranks solutions at query time and sets the StreamK grid at launch. Frameworks usually query once and cache the algo, so when contention changes later the cached solution keeps running, only regridded. This PR lets the launch switch to the solution the heuristic would pick at the current hint. ## Technical Details - **Tag at query** (`rocblaslt_auxiliary.cpp`): the top heuristic result is tagged with the hint it was ranked under (`rocblaslt_matmul_algo::data[4..7]`; `tagAlgo` / `algoTag` in `rocblaslt_adaptive_sm_count.hpp`). - **Re-select at launch** (`tensile_host.cpp`, `reselectForAdaptiveHint`, called in `runContractionProblem` right after `updateTensileProblem`): when the hint of the launch's `adaptive_sm::Launch` (#13263) differs from the tag, the heuristic's pick at that hint runs if it fits the caller's workspace; otherwise the query's solution runs. Launches without a `Launch` (mode off, ext API, grouped GEMM) never re-select. - **Graph capture**: a launch on a capturing stream is not adaptive (#13263), so it has no `Launch` and is never re-selected: it runs the tagged solution as given. - **Memo** (`AdaptiveReselectionMemo` on `TensileDataGemm`, i.e. per matmul descriptor): the last 4 choices, keyed by the launch problem (shape, types, launch hint, caller workspace), the tagged solution, its query hint and the query's workspace limit. A repeat launch skips the ranking. - **Never replaced**: algos obtained by index, lower-ranked results and tuning-file or tuning-cache (`HIPBLASLT_TUNING_MODE=cache|tune`) entries (they are not tagged). - **Tune mode**: re-selection runs before a first-matmul benchmark; the query-hint fallback is re-applied after the benchmark rewrites the problem. The adaptive hint is not part of the tuning key. - **Uniform summation order**: re-selection runs too; the pick must also keep that order at the launch hint (`uniformSummationOrderSupported`, the check heuristic selection uses). - Docs: re-selection notes in the Stream-K how-to; `_LOG=1` also logs re-selections; CHANGELOG entry. ## Test Plan - `hipblaslt-test-hostunit`: tag round trip and version check. - `hipblaslt-test` GPU tests (FORCE-only cases skip otherwise): - a result tagged with another hint re-selects once, and a repeat launch hits the memo; - a result ranked at the full device is swapped for the solution picked at the forced hint, and that one launches; - a launch on a capturing stream runs such a result as given, and the next live launch re-selects it; - under uniform summation order it re-selects too and the launch passes the order's check; - with no workspace at all the launch still re-selects to the StreamK pick (it needs less workspace when given less); - StreamK OFF results are never tagged; - a tuning-file entry is never tagged, so never re-selected. - Manually: the adaptive, SK, tuning and uniform-summation suites with `HIPBLASLT_ADAPTIVE_SM_COUNT=1`, `_FORCE=128` and `_FORCE=128 _FORCE_PROBE=1`. - Launch overhead and ranking cost measured on MI355X (Test Result). ## Test Result - All suites pass; with the mode off behaviour is unchanged. - Under contention, launches from a cached algo run the solution chosen at the measured hint. - Once memoised, a re-selecting launch costs ~0.7 µs more host time than a plain one. - Capturing 16 GEMMs takes 211 µs of host time (172 µs with the mode off), down from 2643 µs when captured launches were re-selected. --- ## Appendix ### From query to launch ``` hipblasLtMatmulAlgoGetHeuristic (rocblaslt_auxiliary.cpp) rank under the query hint (#13263) → top result A adaptive, not from a tuning file → tagAlgo(A, query hint); workspace ≥ 128 MiB │ caller caches A ▼ hipblasLtMatmul(A) (rocblaslt_mat.cpp, #13263) Controller::begin → Launch { hint = the stream's current hint, probe } ▼ runContractionProblem (tensile_host.cpp) updateTensileProblem: data->problem holds the launch hint and workspace reselectForAdaptiveHint(A) → A, or a copy of A with another solution index look up the solution, build kernels, launch ``` The caller's algo is never modified: a swapped solution runs through a local copy of it. ### Launch-time decision ``` launch with algo A │ ├─ no Launch, or A not tagged (by index, lower rank, tuning file/cache) ──► run A │ ├─ launch hint == A's tag ────────────────────────────────────────► run A │ └─ hint differs memo hit? ──yes──► run the remembered choice │ no heuristic's pick P at the launch hint (ranked with the query's workspace limit) ├─ P fits the caller's workspace ──► run P └─ otherwise ──► run A, at the launch hint if A fits there, else at A's tag (the query hint) remember the choice ``` "Fits" means `requiredWorkspaceSize` at the caller's workspace is within it and, under uniform summation order, the order holds at the launch hint. StreamK solutions need less workspace when given less, so on gfx950 the pick fit in every case tried, down to zero workspace; the fallback covers solutions whose workspace is fixed. ### What is remembered Per matmul descriptor, the last 4 re-selections, the oldest replaced first: | Key | Value | |---|---| | launch problem: shape, types, launch hint, caller workspace | solution to run | | tagged solution | whether it runs at the query hint instead of the launch hint | | query hint | | | query's workspace limit | | Example: a framework creates descriptor D for bf16 512×8192×8192, queries once at startup (hint 0, solution A tagged 0) and caches A. | # | Situation | Stream hint | What happens | Memo for D | |---|---|---|---|---| | 1 | GPU idle | 0 | tag 0 = hint 0 → A | – | | 2 | a cotenant takes 64 CUs | 192 | miss → rank at 192 → B fits → B; logs `reselect … hint=0->192 solution=A->B` | {@192 → B} | | 3 | same contention | 192 | hit → B | {@192 → B} | | 4 | the cotenant leaves | 0 | tag 0 = hint 0 → A; memo not consulted | {@192 → B} | | 5 | it returns with 128 CUs | 128 | miss → rank at 128 → C | {@192 → B, @128 → C} | | 6 | back to 64 CUs | 192 | hit → B | unchanged | | 7 | same D, 1024×8192×8192 | 192 | new launch problem → miss → own entry | 3 entries | | 8 | a smaller workspace is passed | 192 | the workspace is in the key → miss → fit re-checked | 4 entries | ### Tag `data[0..3]` keeps the solution index. `data[4]` bit 0 marks the result as re-selectable, `data[5]` is the tag version (an unknown version reads as untagged), `data[6..7]` the hint at query. Untagged algos (all of them with the mode off) behave as before. ### Workspace The tagged result already reports `min(MAX_WORKSPACE_BYTES, max(required, 128 MiB))` (#13263). A re-selected solution must fit the workspace actually passed to `hipblasLtMatmul`; otherwise the query's solution runs. ### Limits - The memo is per descriptor: callers that create a descriptor per call start empty each time. They still hit TensileLite's ranking cache. - The tag holds one query hint: a query made while the hint is 192 tags its result 192, and its launches re-select only once the hint moves away from 192. - Graph capture: captured launches are not adaptive, so they are never re-selected. A replay runs the queried solution at hint 0 (#13263). - The query hint is the latest hint published on any stream (#13263). With per-call queries and different hints on two streams, launches on one of them keep mismatching and re-select through the memo or the ranking cache.

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Oct 7, 2026

therock-pr-bot Bot added the Not ready to Review label Oct 7, 2026

caio96 added the rocm:gemm algos Label to identify PRs of the GEMM Algorithms team. label Oct 7, 2026

assistant-librarian Bot added the organization: ROCm label Oct 7, 2026

caio96 force-pushed the users/csalvado/predict-sm-hint branch 2 times, most recently from 16ba5d4 to 0e6462f Compare October 14, 2026 22:49

github-actions Bot added the documentation label Oct 14, 2026

caio96 mentioned this pull request Oct 14, 2026

feat(tensilelite): add opt-in CU-occupancy probe to StreamK hybrid kernels #13692

Merged

1 task

caio96 changed the title ~~feat(hipblaslt): add adaptive sm_count_target via in-kernel CU probe~~ feat(hipblaslt): add adaptive sm_count_target mode Oct 14, 2026

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 0e6462f to 4186f5a Compare October 14, 2026 23:00

caio96 changed the base branch from develop to users/csalvado/occupancy-probe October 14, 2026 23:00

caio96 added this pull request to stack #13693 October 14, 2026 23:02

therock-pr-bot Bot removed the Not ready to Review label Oct 14, 2026

caio96 force-pushed the users/csalvado/predict-sm-hint branch 2 times, most recently from 9629960 to df88170 Compare October 16, 2026 22:31

caio96 force-pushed the users/csalvado/predict-sm-hint branch 2 times, most recently from 2cf9244 to f75f2fb Compare October 17, 2026 05:26

This was referenced Oct 17, 2026

feat(hipblaslt): add adaptive sm_count_target estimator #13892

Merged

feat(hipblaslt): re-select the top solution when the adaptive hint changes #13893

Merged

caio96 removed this pull request from stack #13693 October 17, 2026 05:27

caio96 changed the base branch from users/csalvado/occupancy-probe to users/csalvado/adaptive-sm-estimator October 17, 2026 05:28

caio96 added this pull request to stack #13894 October 17, 2026 05:28

caio96 force-pushed the users/csalvado/predict-sm-hint branch from f75f2fb to b367f23 Compare October 19, 2026 23:46

caio96 marked this pull request as ready for review October 21, 2026 21:41

caio96 requested review from a team as code owners October 21, 2026 21:41

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 50ada6d to 6a74b20 Compare October 22, 2026 04:18

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 6a74b20 to 824bd49 Compare October 22, 2026 04:21

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 824bd49 to 3b45434 Compare October 22, 2026 04:33

jaopaulolc approved these changes Oct 22, 2026

View reviewed changes

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 3b45434 to eee0c5e Compare October 22, 2026 21:02

Alex-Vasile approved these changes Oct 23, 2026

View reviewed changes

caio96 force-pushed the users/csalvado/predict-sm-hint branch from eee0c5e to 91a50a1 Compare October 23, 2026 21:50

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 91a50a1 to e874d92 Compare October 26, 2026 20:11

caio96 force-pushed the users/csalvado/predict-sm-hint branch from e874d92 to 026f1f9 Compare October 27, 2026 15:28

caio96 added 3 commits October 27, 2026 12:47

Refuse benchmark timing with an unforced adaptive sm_count_target

daab76b

A varying adaptive hint can change the timed kernel mid-run, so the timing clients exit unless HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE fixes it.

Document the adaptive sm_count_target mode

da2a1c1

caio96 force-pushed the users/csalvado/predict-sm-hint branch from 026f1f9 to da2a1c1 Compare October 27, 2026 17:52

caio96 merged commit 49fc36c into develop Oct 28, 2026
178 of 199 checks passed

caio96 deleted the users/csalvado/predict-sm-hint branch October 28, 2026 15:38

Navigation Menu

Repository navigation

feat(hipblaslt): add adaptive sm_count_target mode - #13263

feat(hipblaslt): add adaptive sm_count_target mode#13263
caio96 merged 3 commits into
users/csalvado/adaptive-sm-estimatorfrom
users/csalvado/predict-sm-hint

caio96 commented Oct 7, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 7, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 7, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Oct 7, 2026 •

edited

Loading

Uh oh!

Alex-Vasile left a comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 30, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

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

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Risk level

Appendix

Block flow (one hipblasLtMatmul call)

Timeline (one stream, a cotenant holds 64 of 256 CUs)

Workspace

What adapts

Knobs

Timing rule

Uh oh!

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

Uh oh!

❌ PR Check — Action Required

Uh oh!

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

Uh oh!

Uh oh!

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

Uh oh!

Codecov Report

Uh oh!

Alex-Vasile left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 30, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

caio96 commented Oct 7, 2026 •

edited

Loading

Block flow (one `hipblasLtMatmul` call)

therock-pr-bot Bot commented Oct 7, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 7, 2026 •

edited

Loading

codecov Bot commented Oct 7, 2026 •

edited

Loading