Repository navigation
feat(tensilelite): StreamK launch summary diagnostic report - #10036
Merged
jaopaulolc merged 18 commits intoSep 1, 2026
Merged
Conversation
… self-reset (WGMXCC + SKXCC) The per-XCD StreamK work-queue atomic counter is designed to self-reset to zero every launch purely by atomic wraparound: each queue q must receive exactly tiles_q + W_q increments, where W_q = distribute(skGrid, q) assumes the queue index densely covers [0, skGrid). Deriving the queue index from the remapped StreamKIdx breaks that invariant, because neither the wgmXCC CU-count remap (WorkGroupMappingXCC == -1 auto path) nor the StreamKXCCMapping chiplet remap (SKXCC with WorkGroupMappingXCC > 1) is a % numQueues-count-preserving permutation when the grid does not block evenly. The skewed per-queue home-workgroup count leaves the counter drifting off zero across launches. Fix: derive the queue index from the raw pre-remap launch workgroup rank, which is a dense bijection onto [0, skGrid) (== physical XCD rank) and restores the counting invariant on both the SK4 and SK5 dynamic paths. The raw rank is snapshotted once in the prologue, before wgmXCC and the SKXCC preLoop XCCMapping remap rewrite WorkGroup0, into the persistent StreamKTileIdx SGPR -- an already-allocated slot that is provably dead in the [prologue, queue-read] window -- and read back masked % numQueues at the queue index computation. This adds zero additional persistent SGPRs, so it also fits tuned high-register SKXCC kernels that would overflow the SGPR file with a dedicated queue SGPR. Only once-per-workgroup setup instructions are added; no steady-state cost. Count-preserving regimes (fixed non-SKXCC WGMXCC == 1, WorkGroupIdFromTTM targets, single-queue arches) keep the cheap StreamKIdx % numQueues path.
…ILE_DB 0x200000) Add an observational diagnostic that reports the StreamK launch DECISIONS made for a (solution, problem, hardware) triple as an indented labeled block printed to stderr from ContractionSolution::solve(). It is gated solely by TENSILE_DB bit 0x200000 (sitting alongside the 0x100000 StreamK-mode-selection bit) and is off by default, so it changes no launch behaviour. The report is built from a StreamKDecisions snapshot populated by computeStreamKDecisions(), which is the single source of truth that solve() itself consumes to fill StreamKSettings (reduction, finalGrid, tiles). Sourcing the report from the same values solve() launches with means the summary cannot drift from the real launch. Reported fields: - mode: SK3/SK4/SK5 plus the effective dynamic/static resolution; - reduction: tree vs parallel (data-parallel); - grids: selected -> final with changedBy attribution (none/fixedGrid/treeBounds/workspaceDP), plus the pre-fallback grid when it differs; - tiles/skTiles/skSplit/totalItems and whether partials are present; - workspace: allocated flag plus required/ideal/given bytes in human units, with "unbounded" shown for SIZE_MAX; - dp-only: with the source distinguished (forceDPOnly param vs streamKDP debug vs runtime workspace-DP fallback); - work-queue fields: numQueues (NUM_XCD) and dynamicPartialsSlots, which print "NA (work-queues not used)" on non-dynamic paths; - the fallback flags that fired. Covered by gtests in StreamKLaunchSummary_test.cpp.
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ 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. |
jaopaulolc
added
the
rocm:gemm algos
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #10036 +/- ##
===========================================
+ Coverage 69.33% 69.36% +0.03%
===========================================
Files 2811 2809 -2
Lines 465641 464468 -1173
Branches 68698 68556 -142
===========================================
- Hits 322828 322149 -679
+ Misses 119254 118811 -443
+ Partials 23559 23508 -51
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Replace the general "single source of truth / cannot drift / observational" prose on StreamKDecisions with precise per-field provenance: each field now states whether it is available (read from the value the launch uses, or the same helper/param solve() uses) or recomputed (re-derived by mirroring the makeArgs/solve() sizing and packing logic).
AlexBrownAMD
approved these changes
Aug 19, 2026
Conflict resolution: - Tensile/Components/StreamK.py, Tensile/KernelWriter.py, Tensile/KernelWriterAssembly.py and Tensile/Tests/unit/test_streamk_queue_index.py are resolved to develop. The branch carried an earlier version of the count-preserving per-XCD StreamK work-queue counter fix from when this was a stacked PR; an equivalent fix has since landed independently on develop (#10008). That work is out of scope for this PR, so develop's landed version wins and these files are byte-identical to develop. - include/Tensile/ContractionSolution.hpp keeps both sides: develop's CustomKernel, StreamKStaticSplit / streamKStaticSplit() / streamKStaticSplitRowUniform(), UniformSummationOrderError, the four-element WGMParamsCache key and the checkUniformSummationOrder() declaration, plus this branch's StreamKDecisions struct and the computeStreamKDecisions() / printStreamKLaunchSummary() declarations. Post-merge mirror fixes in src/ContractionSolution.cpp, so the diagnostic report still describes the launch it observes (report-only; no launch behaviour changes): - getSKGridImpl() gained a StreamKForceDPOnly cluster-multicast clamp (#10189) that runs after the tree-fixup-bounds fallback and also resets the grid to the tile count. The summary had no way to attribute that clamp, so a launch taking it reported "changedBy = none". Added a clusterDPGridClamped out-param and StreamKDecisions field, reported under "fallbacks" and in the changedBy attribution (ordered latest-clamp-wins). - The static-split mirror now calls the shared streamKStaticSplit() helper develop factored out, instead of repeating its skTiles arithmetic inline. makeArgs() packs skTiles from that same helper, so the report can no longer drift from the kernel arguments.
The ForceDPOnly cluster-multicast clamp in getSKGridImpl() -- SK3 +
streamKForceDPOnly + clusterDim.x*clusterDim.y > 1 resetting the grid to
tiles -- had no unit coverage. Add four cases:
* the clamp resets the grid to tiles and is reflected in the snapshot,
with getSKGrid() cross-checked against the reported final grid;
* the three-way conjunction is a real guard (SK4, non-force-DP SK3, and
a z-only cluster each leave the grid alone);
* the printed summary names the clamp, credits it for the
selected-vs-final grid change, and surfaces preFallback;
* changedBy credits the clamp over an skFixedGrid override that the
clamp subsequently overwrote, so the report never names a grid the
launch did not use.
The clamp reads only sizeMapping, so all four run on the existing mock
AMDGPU with no analytical hardware and no gfx1250 device.
Also tighten a few existing assertions and comments:
* replace a totalItems assertion that restated the implementation
character-for-character with the concrete literals the fully
determined fixture produces (1792);
* check all four fallback flags in the workspace-DP attribution test
now that clusterDPGridClamped exists;
* turn an if(workspaceAllocated)-guarded workspace sizing check into
ASSERT_TRUE plus an unconditional check so it cannot silently vanish;
* pin d.tiles and d.skGridPreFallback to literals rather than only to
the same helpers the implementation calls;
* correct the ForceDPOnly comment, which claimed the param never resets
the grid to tiles -- true only because the fixture leaves clusterDim
at 1;
* restate the dynamic-path workspace guard in full (parallel OR
tiles%grid != 0) and note why the fourth combination is not tested;
* rewrite the comments that read as a diff against an earlier revision,
and rename NonStreamKProducesEmptySnapshot ->
NonStreamKProducesInertSnapshot and Sk3StaticPartialTiles ->
Sk3StaticPartialTilesReserveWorkspace.
The workspace / DP-fallback logic now lives in computeStreamKDecisions()
rather than solve(), but a dozen comments still described themselves as
mirroring solve(). There is nothing left in solve() to mirror: it copies
the reduction, grid and isDynamic predicate straight out of the snapshot
and does no further StreamK sizing.
Point those comments at the real drift peer instead. requiredWorkspaceSize()
is a second, independent implementation of the same reserve-or-not rule and
is what the caller allocates from, yet it already differs: it always uses
getSKReduction() and sizes parallel reduction with requiredWorkspaceSizeGsu()
rather than partialTileSize(). Note explicitly that
StreamKDecisions::requiredWorkspaceBytes is NOT that function's return value
despite the near-identical name.
Other corrections:
* skTiles/skSplit/totalItems were documented with a formula that is
false on the parallel and static paths -- the SK3 ABI has no
TotalItems arg, and the parallel path packs the split factor into the
skTiles slot. Document all three paths.
* "It has no side effects" in solve() was wrong: computeStreamKDecisions()
must precede the dynamic-queue guard because the guard consumes
isDynamic, and the helpers it calls can write TENSILE_DB diagnostics.
Record what is and is not reachable on a guard-rejected launch.
* The SIZE_MAX -> "unbounded" sentinel is ContractionProblem's default
workspace size (setWorkspaceSize() never called), not something the
bench client passes.
* The print-layout comment claimed a format was "preserved verbatim" for
existing log scrapers; the function is new, so describe the output
shape and which tokens are load-bearing instead.
* Restore "Only these three fields are refreshed" in
reloadDebugBitsForTest()'s doc -- it still refreshes exactly three.
* Document that effectiveDynamic is SK5-only (false for SK4), that
isDynamic is a launch input, why finalGrid and skGrid both exist, the
dpOnlySource ordering, and the (void)problem parameter retention.
jaopaulolc
marked this pull request as ready for review
August 27, 2026 21:17
jaopaulolc
enabled auto-merge (squash)
August 30, 2026 01:53
jaopaulolc
disabled auto-merge
August 30, 2026 01:54
Alex-Vasile
approved these changes
Aug 31, 2026
Alex-Vasile
left a comment
Alex-Vasile
left a comment
Contributor
There was a problem hiding this comment.
Approving CMakeLists changes.
… default Three consistency fixes around the StreamK launch summary: - skGridPreFallback documented the workspace-DP fallback as living in solve(); it is applied in computeStreamKDecisions(), which is what the workspaceDPFallbackFired doc a few lines below already says. - finalGrid's list of the fallbacks it accounts for omitted the ForceDPOnly cluster-multicast clamp. It now names all four clamps in application order using the same names the report's "changedBy" ladder prints, so the doc and the ladder can be compared directly. - The static-split path in computeStreamKDecisions() fell back to skFullTiles=0 when the AMDGPU downcast yielded null. AMDGPU::skFullTiles defaults to 1, and checkUniformSummationOrder() already uses 1 for the same fallback. The path is not reachable today (getSKGridImpl asserts and dereferences the same pointer first), so this only keeps the constant consistent with its two siblings. No behavioural change.
The StreamK launch-summary tests left four behaviours unasserted. Add five tests that close them: - Sk3ParallelReductionReservesPartialsWorkspace: the parallel (split-K DP) reduction, which no existing test reached. A 256x4096x4096 SK3 launch on the 256-CU analytical mock lands in select_reduction's parallel band, so the snapshot's packing (skSplit = grid/tiles) and its unconditional partials reservation are pinned. - Sk3ParallelReductionWorkspaceStarvedFallsBackToTree: the same shape with no workspace, demoting to tree with grid = tiles while the workspace-DP fallback is still credited even though the grid number did not move. - LaunchSummaryDebugBitIsOffByDefault: the summary is opt-in. Asserts the Debug predicate is clear with TENSILE_DB unset, set by 0x200000, not set by the neighbouring 0x100000 mode-selection bit, restored to the ambient value afterwards, and that building the snapshot writes nothing to stderr. - TreeBoundsFallbackWinsGridAttribution and FixedGridOverrideWinsGridAttribution: the two remaining changedBy attributions, each with a complementary run showing the grid is left alone when the clamp does not apply. Adds a ScopedTensileDb RAII guard that copies the prior TENSILE_DB value before setenv invalidates it and restores that value on scope exit, since the Debug singleton is process-wide. The four decision tests assert TENSILE_STREAMK_DATA_PARALLEL is clear, because it is latched at Debug construction and would otherwise change the answers silently. Also correct the StreamKDecisions::requiredWorkspaceBytes comment. It claimed the snapshot and requiredWorkspaceSize() must stay consistent about whether a workspace is needed; they do not. Parallel reduction with a grid equal to tiles gives a k-split of 1, requiredWorkspaceSizeGsu() short-circuits to 0 for a split of 1, and the snapshot still reserves partialTileSize(tiles). The comment now describes that case, why it is an over-reservation rather than an under-allocation, and why the allocate-then-launch loop settles at a self-consistent fixed point.
1 task done
jaopaulolc
added a commit
that referenced
this pull request
Sep 1, 2026
…eamk-per-tile-extra-iters Three files conflicted, all from develop-side StreamK work that landed in this window: #10036 (StreamK launch summary diagnostic), #10951 (multi-stream Stream-K Synchronizer fix) and #10925 (fused all-to-all epilogue). Debug.hpp / Debug.cpp: both sides claimed TENSILE_DB bit 0x200000. #10036 is merged and its bit is documented and exercised by StreamKLaunchSummary_test, so it keeps 0x200000; the uniform-summation-order no-solution report moves to 0x400000, the next free bit. Both diagnostics remain independently switchable. ContractionSolution.cpp, getSKGridImpl(): #10036 added out-params recording which post-selection fallback rewrote the grid, and captures the selected grid between the CU/override ladder and the tree-fixup bounds clamp -- the same point the uniform-summation-order F-star snap is inserted at. The snap now runs first and the capture follows it, because the snap is grid selection under uniform summation order rather than a fallback: it has no changedBy attribution, so capturing ahead of it would report selected != final with changedBy = none, the unattributed rewrite the out-params exist to prevent. #10951's flag-region clamp stays last, immediately before the return, so the launch summary's grid ordering is unchanged. ContractionSolution.cpp, solve(): #10036 extracted the reduction / grid / workspace-DP-fallback logic into computeStreamKDecisions(); this branch had extracted the same logic into resolveStreamKSettings(). resolveStreamKSettings() remains the source of truth for what is launched, since it additionally applies streamKReconcileReduction() so requiredWorkspaceSize() and the launch agree on the reduction. computeStreamKDecisions() now applies that same reconciliation at the same point, before the workspace-fit fallback, so the snapshot cannot report a fallback the launch did not take; solve() re-asserts the final reduction alongside the final grid before printing.
caio96
added a commit
that referenced
this pull request
Sep 2, 2026
…exes (#11515) #10951 gives Stream-K its own flag buffer — a private 2048-int block per (stream, problem) — and adds the first bound on `skGrid` against that block. The bound assumes all 2048 entries are addressable as flags, but `StreamK=4` and the SK4 sub-path of `StreamK=5` place the per-XCD work-queue counters at the *base* of that same region and start the ready flags after them. On gfx950 that prefix is 8 XCDs × 128 B = 256 ints, leaving 1792 indexable. A grid between 1793 and 2048 has its last workgroups write past the end of their own block; blocks are laid out as adjacent problem slots within a stream, so the overrun lands on the next problem's flags, or on the next stream's from the last slot of a block. The prefix is XCD-count × cache-line, so it varies by ISA and cannot be folded into the constant. Default grids stay well inside the bound (`skGrid` defaults to the CU count, 256 on gfx950), but `TENSILE_STREAMK_GRID_MULTIPLIER` scales it by the CU count with no cap of its own, so the bound is reachable through a supported tuning knob. Separately, the host reserved those same queue-region bytes in the **workspace** — a different buffer. The kernel takes them from `AddressFlags`, so the reservation was unread. ## Technical Details **Bound the grid by the entries the kernel can actually index.** `getSKGrid` subtracts the queue-counter prefix for the dynamic-queue paths and keeps the full region for the static ones, which index from offset 0 and would otherwise lose grid they are entitled to. Two fallbacks keep the full bound: an architecture whose XCD count cannot be determined, and a prefix that would swallow the region. **Stop reserving the queue region in the workspace.** `AddressWS` appears exactly once in codegen (`computeWorkspaceSrd`) and applies only `MacroTile0*MacroTile1*bpe * partialIdx` — no queue-region base offset — while every counter address is built from `AddressFlags` (`StreamK.py:642`, `:3640`, `:4436`). Removing it makes the reported workspace requirement match what the kernel writes. **Deduplication and comment fixes.** `streamKQueueRegionBytes()` and `streamKUsesDynamicQueue()` replace expressions open-coded at three sites each. The flag-region comments claimed indexing by CTA id, which holds only for the static paths. The grouped uniform-summation-order comment named a grid-divisibility check, but the rejection that actually fires is "resolved StreamK grid is 0". A user-facing error message also described sizing a workspace region this PR removes. Rebased over #10036, which moved this logic into `computeStreamKDecisions()`; the changes now land there and in `requiredWorkspaceSize()`. ## Test Plan New unit tests in `tensilelite/tests/CuCount_test.cpp` pin the bound for each path — SK4, SK3, and both SK5 sub-modes — reusing the existing `StreamK5AnalyticalEnv` scaffolding. Built for gfx950 and run on MI355X, plus a client-level Stream-K sweep as a regression guard on solution selection. ## Test Result `tensilelite-tests`: `StreamKFlagBound*` 4/4, `StreamK5WorkspaceRegressionTest*` 3/3, `CuCount*` 7/7; full binary 444 passed, 2 skipped, 0 failed. `hipblaslt-test --gtest_filter='*StreamK*'`: 33/33, no skips. Both skips are pre-existing and unrelated. A GPU probe confirmed the workspace change: a dynamic-queue launch succeeds at exactly `partialTileSize` and is rejected 1024 B below it, so the requirement dropped by exactly the queue region and the kernel is correct at the new minimum. JIRA ID : AIHPBLAS-4624
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AIHPBLAS-4513
Motivation
When a GEMM runs through the StreamK launch path there is currently no way to see WHICH decisions the launch-parameter logic took for a given problem/solution: the StreamK mode (SK3/SK4/SK5, and how SK5 resolved), the reduction strategy (tree vs data-parallel), the grid actually launched (and whether a fallback shrank it), the tile/split accounting, and the workspace/DP-only path.
These decisions involve several fallbacks — fixed-grid override, tree-fixup 24-bit bounds, the
StreamKForceDPOnlycluster-multicast clamp, and the workspace-insufficient DP fallback — that silently rewrite the grid and the reduction strategy. Two different sizes can land on the same kernel and still launch with different grids and different reductions, with nothing in the logs to say why. This change adds an opt-in diagnostic that surfaces those decisions.Technical Details
Adds an observational StreamK "launch summary" printed as a deeply-indented labeled block to stderr from
ContractionSolution::solve().Gating. Printed only when
TENSILE_DBbit0x200000is set (Debug::printStreamKLaunchSummary()), sitting alongside the existing0x100000StreamK-mode-selection bit. Off by default; the report is inert when the bit is unset.Field provenance. The report is built from a
StreamKDecisionssnapshot populated byContractionSolution::computeStreamKDecisions()and consumed bysolve(). A diagnostic that re-derives its own numbers is worse than useless if it drifts from the launch it claims to describe, so every field in the struct is annotated as one of exactly two kinds, and the annotation is part of the struct's documentation:solve()calls, so the field reflects the launch exactly. These are:streamKMode(sizeMapping.streamK),effectiveDynamic(streamK5EffectiveDynamic()),reductionandfinalGrid/skGrid(the exact valuessolve()wires intoStreamKSettings),tiles(problem.getNumTiles()),selectedGrid/skGridPreFallback/fixedGridUsed/treeBoundsFallbackFired/clusterDPGridClamped(getSKGridImpl()return value and out-params),forceDPOnly(sizeMapping.streamKForceDPOnly),streamKDP(Debug::useStreamKDataParrallel()),givenWorkspaceBytes(problem.workspaceSize()), andnumQueues(streamKBakedQueueCount()).computeStreamKDecisions()by mirroring themakeArgs()/solve()sizing and packing logic. These are:isDynamic(mode +effectiveDynamic),skTiles/skSplit/totalItems(themakeArgs()packing),partialsPresent(skTiles > 0),dpOnly(OR of the three DP triggers),workspaceAllocated/requiredWorkspaceBytes/idealWorkspaceBytes(the workspace sizing guard),dynamicPartialsSlots(skTiles * skSplit), andworkspaceDPFallbackFired.Recomputed fields are the ones that can drift, and they are called out as such so a future change to the packing logic has an obvious place to look. Where the mirrored arithmetic already exists as a shared helper it is called rather than copied: the static-split accounting uses
streamKStaticSplit(), the same helper the SK3 and SK5-static packers use, instead of a third hand-rolled copy of that arithmetic.requiredWorkspaceBytesin the snapshot is deliberately not the return value ofContractionSolution::requiredWorkspaceSize(), despite the similar name. That function is the separate, pre-existing, caller-facing implementation the allocator sizes from; it is untouched here. The snapshot instead mirrors the sizing guard insidesolve()and evaluates it on the final, post-fallback grid, so the report describes the launch rather than the allocation request. The two encode the same reserve-or-not rule and must stay consistent about whether a workspace is needed; both sites carry a comment saying so.Reported fields.
mode:SK3/SK4/SK5 plus the effective dynamic/static resolution for SK5 (e.g.SK5->dynamic(SK4)), andisDynamic.grid:selected -> finalwith achangedByattribution (none/workspaceDP/clusterDPMulticast/treeBounds/fixedGrid), plus the pre-fallback grid when it differs. The ladder is ordered latest-clamp-wins, so the attribution names the clamp that actually produced the final grid.tiles:tiles/skTiles/skSplit/totalItems/partials.workspace:the allocated flag plus required / ideal / given bytes in human units, withunboundedshown forSIZE_MAX.dp-only:with the source distinguished —forceDPOnly(param)(compile-time solution parameter) vsstreamKDP(debug)(debug override) vsworkspaceDP(runtime)(the runtime workspace-insufficient fallback).work-queue:numQueues(NUM_XCD)anddynamicPartialsSlots, which printNA (work-queues not used)on non-dynamic paths rather than misleading zeros.fallbacks:the flag for each fallback that can rewrite the grid.Cluster-multicast clamp.
getSKGridImpl()has a third post-selection grid clamp, which resetsskGrid = tilesforstreamK == 3withstreamKForceDPOnlyandclusterDim.x * clusterDim.y > 1. It runs after the tree-fixup-bounds clamp. Without explicit handling the report would showselected != finalwithchangedBy = none— misdiagnosing precisely the situation it exists to explain — so the snapshot carries aclusterDPGridClampedfield, thechangedByladder has an entry for it, and thefallbacks:section prints aclusterDPMulticastline.Ordering.
computeStreamKDecisions()runs beforesolve()'s unsupported-dynamic-queue guard, because the guard consumesisDynamicfrom the snapshot. This is not a behaviour change: the diagnostics and asserts that this could have reordered all live in the non-dynamicelsearm ofgetSKGridImpl(), which is unreachable whenever the guard fires;getSKReduction()is not called at all on the SK4 / SK5-dynamic path; and the0x100000note emitted bystreamK5EffectiveDynamic()already ran ahead of the guard.Observational only. The snapshot's
reduction,skGridandisDynamicare genuine launch inputs, but computing the snapshot and printing the summary change no launch behaviour.Example output
The report goes to stderr, so redirect it — piping stdout shows nothing:
A bare
hipblaslt-benchresolveslibhipblaslt.so.1against the installed ROCm library, not a local build. To exercise a build, point the loader and the library paths at it ($HIPBLASLT_BUILDbelow is the hipBLASLt build directory, and<arch>the target GPU architecture):One block is emitted per launch, so
-i 1 -j 0is worth passing: with the bench defaults (-i 10 -j 2) each of the first two scenarios below emits 12 identical blocks. The captures are from gfx950 (MI355X,NUM_XCD = 8). Kernel names are elided mid-string (Cijk_Ailk_Bjlk_..._MT64x64x64_..._WGMXCCn1) for width; the rest of each block is verbatim.The first two scenarios are a paired contrast: identical kernel, identical problem, only
--streamk_tile_schedulingdiffers, and SK5 resolves the other way. That flag selects the SK5 tile-scheduling sub-path through theHIPBLASLT_MATMUL_DESC_STREAMK_TILE_SCHEDULING_EXTattribute and acceptsoff|0,on|1,auto|2. Neither needs--algo_method all; both hit their path under default heuristic selection.Dynamic path. The per-XCD work-queue synchronizer is in use, so the work-queue fields report real counts:
Static / data-parallel path. The same kernel and shape with tile scheduling forced off resolves to SK3 with a parallel (DP) reduction. The grid goes 16 -> 256, a 4 MiB partials workspace is now reserved, and the work-queue fields are meaningless on this path so they print
NA (work-queues not used):A fallback rewrites the grid. With the workspace pinned to 0, a StreamK-dynamic tree-reduction solution cannot reserve its partials workspace, so the runtime workspace-insufficient DP fallback rewrites the grid:
selected = 768 -> final = 1024,changedBy = workspaceDP, anddp-only source = workspaceDP(runtime). Noteideal = 12.00 MiBagainstgiven = 0— the report shows what the launch wanted as well as what it got, which is the whole reason the fallback fired.This scenario does need
--algo_method allto enumerate solutions. It emitted 1670 blocks across all enumerated solutions, of which 461 showedchangedBy = workspaceDP; the block below is from the modal(selected, final) = (768, 1024)group and was byte-identical across two independent runs.Test Plan
tensilelite/tests/StreamKLaunchSummary_test.cpp(suiteStreamKLaunchSummaryTest, registered intensilelite/tests/CMakeLists.txt). They covercomputeStreamKDecisions()andprintStreamKLaunchSummary()across StreamK modes and fallback paths: SK5 resolving dynamic and static, SK3 static partial tiles reserving a workspace, theforceDPOnlyand workspace-DP DP-only sources,selectedvsfinalgrid attribution, the dynamic workspace-reservation cases (slots/divisibility matrix), the inert snapshot for non-StreamK solutions, the emitted field text, theNA (work-queues not used)non-dynamic behaviour, and four cases for the cluster-multicast clamp — that it clamps the grid totiles, that it requires SK3 +forceDPOnly+ an x*y cluster, that it is reported in the printed block, and that it winschangedByattribution overfixedGrid.tensilelite-testsbinary, to confirm no regression elsewhere.Test Result
Built and run for gfx950 (MI355X,
NUM_XCD = 8):StreamKLaunchSummaryTest: 19 of 19 pass.tensilelite-tests: 409 ran, 407 passed, 0 failed, 2 skipped. The two skips are pre-existing and unrelated to this change (RangeLibraryTest/RangeLibraryTest.SpecificSizes/0andFP4ZeroFrequency/MXDataGenFP4Test.ZeroFrequencyWithinBounds/3).Limits of the runtime evidence, stated plainly:
clusterDPMulticast = yes; it isnoin all 1670 blocks captured. The clamp requires astreamKForceDPOnlysolution (kernel tokenSKFDPO1), and the shipped gfx950 code-object library contains zero such kernels. The hardware runs therefore confirm the field is wired up and printing, but the unit tests are the only coverage of that clamp actually firing.treeBoundsFallback = yesorfixedGrid = yes; both arenothroughout. Those paths are covered by unit tests only.Submission Checklist
Risk level
Low. The report is diagnostic and observational only, and is gated off by default behind
TENSILE_DBbit0x200000. It introduces no change to StreamK launch behaviour.