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

caio96 · 2026-10-17T05:26:57Z

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 feat(hipblaslt): add adaptive sm_count_target mode #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

therock-pr-bot · 2026-10-17T05:32:35Z

❌ 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-17T05:32:38Z

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

codecov · 2026-10-17T08:35:32Z

Codecov Report

❌ Patch coverage is 92.34234% with 17 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...blaslt/src/include/rocblaslt_adaptive_sm_count.hpp	92.31%	9 Missing and 7 partials ⚠️
...lt/library/src/amd_detail/rocblaslt/src/handle.cpp	66.67%	1 Missing ⚠️

Additional details and impacted files

@@                        Coverage Diff                         @@
##           users/csalvado/occupancy-probe   #13892      +/-   ##
==================================================================
+ Coverage                           76.11%   76.12%   +0.01%     
==================================================================
  Files                                2886     2887       +1     
  Lines                              484025   484232     +207     
  Branches                            73007    73039      +32     
==================================================================
+ Hits                               368399   368591     +192     
- Misses                              90392    90401       +9     
- Partials                            25234    25240       +6

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`50.81% <ø> (ø)`
TensileLite-Unit	`84.07% <ø> (ø)`
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from e90c3ba
hipBLASLt	`66.45% <92.34%> (+0.26%)`	⬆️
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 Δ
...ail/rocblaslt/src/include/rocblaslt_secure_env.hpp	`100.00% <100.00%> (ø)`
...lt/library/src/amd_detail/rocblaslt/src/handle.cpp	`56.00% <66.67%> (-1.63%)`	⬇️
...blaslt/src/include/rocblaslt_adaptive_sm_count.hpp	`92.31% <92.31%> (ø)`

🚀 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.

jaopaulolc

LGTM.

Alex-Vasile

Approving CMakeLists changes

The adaptive sm_count_target settings parse the same on/off tokens in a HIP-free header that clients also include, so the trim-and-lowercase helper moves out of handle.cpp into rocblaslt_secure_env.hpp.

A HIP-free per-stream estimator turns occupancy-probe readings into an sm_count_target hint: environment config, probe cadence and epochs, the probe-window hint rule and miss handling. It is not called yet; host unit tests cover 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) → #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 #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 #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 #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 #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](https://claude.com/claude-code)

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

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 the project: hipblaslt label Oct 17, 2026

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

This was referenced Oct 17, 2026

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

Merged

feat(hipblaslt): add adaptive sm_count_target mode #13263

Merged

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

Merged

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

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from d71678b to 49d6d34 Compare October 19, 2026 23:46

caio96 marked this pull request as ready for review October 19, 2026 23:47

caio96 requested review from a team as code owners October 19, 2026 23:47

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 49d6d34 to 2a8751b Compare October 20, 2026 23:57

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch 2 times, most recently from 0efa2bc to 4e54ec4 Compare October 22, 2026 04:18

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 4e54ec4 to 8f15efb Compare October 22, 2026 04:21

jaopaulolc reviewed Oct 22, 2026

View reviewed changes

Comment thread projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp Resolved

jaopaulolc approved these changes Oct 22, 2026

View reviewed changes

AlexBrownAMD approved these changes Oct 22, 2026

View reviewed changes

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 8f15efb to 2420ce4 Compare October 22, 2026 21:02

Alex-Vasile approved these changes Oct 23, 2026

View reviewed changes

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 2420ce4 to 50eb4fe Compare October 23, 2026 21:50

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 50eb4fe to 004d8e2 Compare October 26, 2026 20:11

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 004d8e2 to 98c5b51 Compare October 27, 2026 15:28

caio96 added 2 commits October 27, 2026 12:47

Move the environment value normalizer next to the secure getenv

9836cde

The adaptive sm_count_target settings parse the same on/off tokens in a HIP-free header that clients also include, so the trim-and-lowercase helper moves out of handle.cpp into rocblaslt_secure_env.hpp.

Add the adaptive sm_count_target estimator

e9243b7

A HIP-free per-stream estimator turns occupancy-probe readings into an sm_count_target hint: environment config, probe cadence and epochs, the probe-window hint rule and miss handling. It is not called yet; host unit tests cover it.

caio96 force-pushed the users/csalvado/adaptive-sm-estimator branch from 98c5b51 to e9243b7 Compare October 27, 2026 17:52

caio96 merged commit bf40f0b into develop Oct 28, 2026
142 of 159 checks passed

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

Navigation Menu

Repository navigation

feat(hipblaslt): add adaptive sm_count_target estimator - #13892

feat(hipblaslt): add adaptive sm_count_target estimator#13892
caio96 merged 2 commits into
users/csalvado/occupancy-probefrom
users/csalvado/adaptive-sm-estimator

caio96 commented Oct 17, 2026 •

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 17, 2026

Uh oh!

codecov Bot commented Oct 17, 2026 •

edited

Loading

Uh oh!

Uh oh!

jaopaulolc left a comment

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!

4 participants

Conversation

caio96 commented Oct 17, 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

Per-launch decisions (one stream)

Reading a probe

Readings → hint (window = last PROBE_WINDOW readings, default 3)

Knobs (prefix HIPBLASLT_ADAPTIVE_SM_COUNT)

Uh oh!

therock-pr-bot Bot commented Oct 17, 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 17, 2026

Uh oh!

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

Uh oh!

Codecov Report

Uh oh!

Uh oh!

jaopaulolc left a comment

Choose a reason for hiding this comment

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!

4 participants

caio96 commented Oct 17, 2026 •

edited

Loading

Readings → hint (window = last `PROBE_WINDOW` readings, default 3)

Knobs (prefix `HIPBLASLT_ADAPTIVE_SM_COUNT`)

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

edited

Loading

codecov Bot commented Oct 17, 2026 •

edited

Loading