Repository navigation
feat(hipblaslt): re-select the top solution when the adaptive hint changes - #13893
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #13893 +/- ##
===========================================
- Coverage 77.29% 77.24% -0.05%
===========================================
Files 2889 2892 +3
Lines 484853 485458 +605
Branches 73234 73340 +106
===========================================
+ Hits 374740 374989 +249
- Misses 85381 85717 +336
- Partials 24732 24752 +20
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
1e80b5e to
9255d2b
Compare
9255d2b to
fddf430
Compare
056944d to
fb27400
Compare
fb27400 to
75f1b07
Compare
75f1b07 to
5d68c49
Compare
5d68c49 to
84f0157
Compare
84f0157 to
0e47afe
Compare
0e47afe to
dde7db0
Compare
dde7db0 to
ffd9338
Compare
ffd9338 to
ae8c8be
Compare
ae8c8be to
33da47f
Compare
The heuristic query ranks solutions under the hint current at query time, which can differ by launch. The top Tensile result is tagged with its query hint; a launch under a different hint swaps it for the heuristic's pick at the launch hint when that fits the caller's workspace, memoised per matmul desc. A tuning-file result is never tagged, so its recorded solution still wins.
…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)
33da47f to
133031c
Compare
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)
PR Pundit review(testing PR Pundit Code Review agent. Not all comments may be relevant. Please reach out to @peymanr for suggestions and feedback.)
1. 🟠 Should fix: Tune-mode re-benchmark undoes the atQueryHint fallback override
When Who is affected: Callers with both adaptive mode ( 2. 🟠 Should fix: New re-selection tests never exercise an actual solution swap
The four new re-selection tests ( Who is affected: Maintainers relying on the test suite: the PR's main behaviour (launching a different solution when the hint changes) is untested, and default CI runs none of these cases. Verified by running it. 3. 🔵 Nit: Explicit algo silently swapped, breaking documented tuning contract
In adaptive mode ( Who is affected: Callers who enable both adaptive SM-count mode and a tuning mode, and who pass the top heuristic result back to Related past reviews: #5004 on 4. 🔵 Nit: Rank workspace uses max, contradicting comment about query's limit
The re-ranking key uses Who is affected: Readability only; no runtime behaviour differs for any caller. 5. 🔵 Nit: Unsynchronised memo and deep copy on hot launch path
Who is affected: Callers with adaptive SM count enabled ( Verified by running it. Related past reviews: #860 on 6. 🔵 Nit: Tag bytes in opaque algo break bytewise compare/cache/serialize
In adaptive mode only ( Who is affected: Applications that opt into 7. 🔵 Nit: Non-RAII restore of shared problem throws on heuristic failure
Who is affected: Only launches with an adaptive launch hint that differs from the tagged query hint, and only if the solution lookup throws, which this review did not show can happen. 8. 🔵 Nit: Workspace-fallback launch hint divergence undocumented
The CHANGELOG and Who is affected: Users of adaptive 9. 🔵 Nit: Docs claim only override-file results exempt; cache mode also exempt
The CHANGELOG and streamk doc say only Who is affected: Users who enable both adaptive Verified by running it. 10. 🔵 Nit: Hint override comment does not explain the atQueryHint rationale
The comment Who is affected: Maintainers reading the 11. 🔵 Nit: Memo eviction and multi-key miss paths untested
Who is affected: Maintainers relying on the adaptive SM-count gtests; the memo eviction and cross-key miss paths are only exercised when one descriptor sees more than one hint/workspace key. Verified by running it. Shared from a pilot of PR Pundit, an agentic code reviewer that builds and tests the change and reads the project's past reviews. Reviewed at |
…e-mode benchmarking A tune-mode call that benchmarks a new shape rewrites the Tensile problem afterwards, which reset the hint a re-selection had pinned to the query hint. The pin is now reported by reselectForAdaptiveHint and re-applied after that rewrite. Also: - test an actual solution swap under FORCE; - document the query-hint fallback, that tuning-cache results are never re-selected, and the adaptive exception to "an explicit algo launches as given"; - rank with algo->max_workspace_bytes, which hipblasLtMatmul already clamps the caller's workspace to. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
TheRock Submodule Bump ActivityNewest first
|
JIRA ID : N/A
Stack: #13692 (kernel probe) → #13892 (estimator) → #13263 (adaptive launches) → this PR (re-selection). Review in that order.
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
rocblaslt_auxiliary.cpp): the top heuristic result is tagged with the hint it was ranked under (rocblaslt_matmul_algo::data[4..7];tagAlgo/algoTaginrocblaslt_adaptive_sm_count.hpp).tensile_host.cpp,reselectForAdaptiveHint, called inrunContractionProblemright afterupdateTensileProblem): when the hint of the launch'sadaptive_sm::Launch(feat(hipblaslt): add adaptive sm_count_target mode #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 aLaunch(mode off, ext API, grouped GEMM) never re-select.Launchand is never re-selected: it runs the tagged solution as given.AdaptiveReselectionMemoonTensileDataGemm, 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.HIPBLASLT_TUNING_MODE=cache|tune) entries (they are not tagged).uniformSummationOrderSupported, the check heuristic selection uses)._LOG=1also logs re-selections; CHANGELOG entry.Test Plan
hipblaslt-test-hostunit: tag round trip and version check.hipblaslt-testGPU tests (FORCE-only cases skip otherwise):HIPBLASLT_ADAPTIVE_SM_COUNT=1,_FORCE=128and_FORCE=128 _FORCE_PROBE=1.Test Result
Submission Checklist
Risk level
Low. Only in adaptive mode, only for the tagged top heuristic result, and only to a solution that fits the caller's workspace.
Appendix
From query to launch
The caller's algo is never modified: a swapped solution runs through a local copy of it.
Launch-time decision
"Fits" means
requiredWorkspaceSizeat 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:
Example: a framework creates descriptor D for bf16 512×8192×8192, queries once at startup (hint 0, solution A tagged 0) and caches A.
reselect … hint=0->192 solution=A->BTag
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 tohipblasLtMatmul; otherwise the query's solution runs.Limits
🤖 Generated with Claude Code