feat(tensilelite): factored 2-D StreamK cluster mode for gfx1250 - #9612
Closed
jaopaulolc wants to merge 87 commits into
Closed
feat(tensilelite): factored 2-D StreamK cluster mode for gfx1250#9612jaopaulolc wants to merge 87 commits into
jaopaulolc wants to merge 87 commits into
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 ❌ Your project check has failed because the head coverage (76.92%) 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 @@
## develop #9612 +/- ##
===========================================
- Coverage 69.56% 68.58% -0.99%
===========================================
Files 2740 2741 +1
Lines 450875 457589 +6714
Branches 66376 69368 +2992
===========================================
+ Hits 313651 313806 +155
- Misses 117000 122628 +5628
- Partials 20224 21155 +931
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jaopaulolc
added
the
rocm:gemm algos
jaopaulolc
force-pushed
the
users/jolabega/streamk-factored-cluster-mode
branch
from
July 22, 2026 15:48
45dd3cd to
ca47b80
Compare
…ulticast
JIRA ID: AIHPBLAS-3929
The TDM multicast ("cluster load") mask machinery was duplicated inline across
KernelWriter, KernelWriterAssembly, and Subtile/SubtileGREmit, and Multicast was
implicitly coupled to ClusterDim != [1, 1]. That coupling made it impossible to
compose barrier-only clustering and cooperative-load clustering independently, or
to reuse the mask attach at a new load site.
- New ClusterLoad component (Components/ClusterLoad.py, registered via
Component.py + Components/__init__.py). ClusterLoadTDM centralizes the mask
value computation (computeMasks), the MulticastMask* SGPR declare/undeclare
(declareSgprs / undeclareSgprs), the topology decision (combined vs split
A/B masks), and the descriptor attach at each load site (applyToDescriptor).
It is a behavior-preserving extraction: every method emits byte-identical
assembly to the previous inline code, and computeMasks receives the exact
SGPR operands the caller already holds so register indices are unchanged.
Selection is capability-based (HasTDM + TDMInst == 3), like TensorDataMoverLoad.
- KernelWriter / KernelWriterAssembly / SubtileGREmit now route mask declare,
undeclare, compute, and descriptor attach through ClusterLoadTDM.find(...).
- Tri-state Multicast parameter (ValidParameters + GlobalParameters default -1):
-1 = auto (legacy: ClusterDim != [1,1] implies Multicast, minus the StreamK
cluster paths), 0 = force off, 1 = force on. Default -1 reproduces the historic
derivation exactly, so every existing YAML (which omits Multicast) derives
byte-identically.
- Solution.py derives Multicast/ClusterBarrier from the tri-state, keying the
legacy auto coupling and the ClusterBarrier gate on StreamK == 0.
- Common/Utilities.clusterEnabled() helper for the ClusterDim != [1,1] test.
- Unit: test_cluster_load_component.py (component find/declare/compute/attach),
test_multicast_tristate.py (-1/0/1 derivation, legacy equivalence).
Low. Behavior-preserving refactor; the tri-state defaults to -1 (legacy auto),
so emitted assembly is unchanged for all existing configs.
…icast) Add the gfx1250 StreamK DP cooperative B-multicast fast path (StreamKMulticast), auto-derived from StreamK=3 + ClusterDim on top of the ClusterLoad component + tri-state Multicast foundation. - Collapse: StreamK==3 && ClusterDim!=[1,1] auto-enables StreamKMulticast (derived-only internal state; no YAML opt-in), forcing Multicast on and enabling the cluster-scope barrier handshake (ClusterBarrier). - StreamK.py: multicast mask predicate (clusterMulticastValid), DP->SK boundary mask clear, and a prologue wave-0 s_barrier_signal -3 that pairs the InsertClusterBarrierPass first-load wait (balanced signal/wait counts). - KernelWriter.py: postMainLoopBarrierCheckAndReset preserves cluster-scope -3 barriers. - ClusterLoad.py: re-add the split-mask hooks (usesCombinedMask / maskSgprName) for StreamKMulticast. - _validateStreamKMulticast: SK3 + [C,1] pow2 + gfx1250/HasTDM/TDMInst=3 + XCC=0; Multicast==1 force-on hard-guard. - Host: streamKMulticast field threaded through SizeMapping / C++ structs; grid multiple-of-C guard + getSKGridImpl ceil(tiles/C)*C override; SK4/SK5 + ClusterDim reject. Mutual-exclusion with the barrier-only cluster reduction is deferred to the following stacked PR. JIRA ID : AIHPBLAS-3929
… the zero-iteration skip path The StreamKMulticast prologue emits a wave-0-elected cluster-scope arrive before the first cooperative-multicast load. Its only matching cluster-scope wait was the first-load wait, which sits after the last-iteration guard's long branch to PrefetchGlobalLastIterEnd. On the zero-full-iteration path (reachable when K is not a whole multiple of DepthU) that branch skips the first-load wait, leaving the prologue arrive unmatched and the cluster-scope barrier unbalanced on that edge. Add a matching cluster wait on the skip edge (guarded by StreamKMulticast) so the arrive is consumed on every control-flow path, preserving whole-cluster barrier symmetry. The wait branches over on scc0 (>=1 full iteration, where the first-load wait pairs the arrive) and executes an all-waves cluster wait on scc1 (zero iterations); it leaves scc intact so the subsequent long branch is unaffected. Update the three gfx1250 cluster char-test balance checks: the skip-edge wait and the first-load wait are mutually exclusive at runtime but both emitted statically, so the static cluster-wait count is now exactly one greater than the signal count (one prologue arrive consumed by exactly one of the two waits).
…multicast - Option A prologue prefetch handshake bracketed with a cluster-scope barrier round (StreamK.py / KernelWriter.py StreamKMulticast plumb). - Remove the PGR<=1 enumeration guard so PGR2 multicast solutions build. - InsertClusterBarrierPass: drain the cooperative broadcast tensor group with both a pre-round drain and a producer-side drain (after the cooperative tensor_load group), gated on StreamKMulticast && PGR>=2. - PrefetchGlobalRead: [1,2] in the mxf4/mxf8 cluster-multicast configs plus a new PGR2 characterization test (+ snapshot + designed yaml). GPU-validated on gfx1250. JIRA ID : N/A
StreamK GEMM tiles whose partial results are reduced across workgroups (WGs) synchronize that reduction with a cross-CU global-workspace flag spin-wait. On gfx1250 the WGs that reduce one tile can instead be co-located into a hardware workgroup cluster and synchronized with an intra-cluster split barrier, removing the spin-wait from the fast path. Adds an opt-in solution parameter StreamKClusterReduction. When enabled, a StreamK tile's fixup peers are placed in a single 1-D workgroup cluster (ClusterDim = [C, 1]); the HW WG-id remap (WorkGroup0 = cluster_x*C + wg_x) makes each cluster own the contiguous StreamK index range [c*C, c*C+C), i.e. exactly one tile's peers. - Kernel (Components/StreamK.py): clusterReduceSignal / clusterReduceWait / clusterReduceIntraCheck helpers emit the wave-0-elected s_barrier_signal / s_barrier_wait handshake in the fixup epilogue. A uniform intra-cluster predicate selects the barrier fast path or the retained global-flag path so a cluster never mixes the two. The WG-id reread is skipped under clustering so StreamKIdx stays the global index. - Host (src/ContractionSolution.cpp, ContractionSolution.hpp + serialization): getSKGridImpl reduction override rounds the SK grid to C*tiles (fixed even split, SKItersPerWG = itersPerTile/C) and keeps the launch grid a multiple of C; the two-tile cluster-reduction kernarg block is emitted only when sk.grid == C*tiles; the field is threaded through SizeMapping / Contractions. - Predicate (ContractionProblemPredicates.hpp + serialization + Contractions.py): ClusterReductionIterCheck rejects at selection time any problem where itersPerTile is not a multiple of C (the split barrier would otherwise over-signal); the global-flag reduction is retained as the runtime fallback. - Solution.py: _validateStreamKClusterReduction gates the opt-in (SK3, non-atomic, non-DP-only, [C,1] power-of-two cluster, gfx1250 with HasClusterBarrier, TDMInst != 0). Multicast and the mainloop ClusterBarrier are kept off for a StreamK reduction cluster: its peers iterate different K-splits and never march the mainloop in lockstep. Multicast and cluster reduction are mutually exclusive: the StreamKMulticast auto-derivation is gated on NOT StreamKClusterReduction, and both validators reject the simultaneous combination with a neutral message. This stacks on the reusable ClusterLoad component + tri-state Multicast and the StreamK cooperative multicast; the mutual exclusion is intended to be unified by a later factored combined-cluster mode. Validation: CPU asm-string unit tests (test_streamk_cluster_reduction.py), snapshot codegen characterization (test_streamk_cluster_reduction_gfx1250_char), the ClusterReductionIterCheck predicate characterization and C++ gtests, plus opt-in-off client GEMM YAMLs. The feature is off by default and strictly opt-in.
… the zero-iteration skip path The StreamKMulticast prologue emits a wave-0-elected cluster-scope arrive before the first cooperative-multicast load. Its only matching cluster-scope wait is the first-load wait, which sits after the last-iteration guard's long branch to PrefetchGlobalLastIterEnd. On the zero-full-iteration path (reachable when K is not a whole multiple of DepthU) that branch skips the first-load wait, leaving the prologue arrive unmatched and the cluster-scope barrier unbalanced on that edge. Add a matching cluster wait on the skip edge (guarded by StreamKMulticast) so the arrive is consumed on every control-flow path, preserving whole-cluster barrier symmetry. The wait branches over on scc0 (>=1 full iteration, where the first-load wait pairs the arrive) and executes an all-waves cluster wait on scc1 (zero iterations); it leaves scc intact so the subsequent long branch is unaffected. Update the gfx1250 cluster char-test balance checks: the skip-edge wait and the first-load wait are mutually exclusive at runtime but both emitted statically, so the static cluster-wait count is now exactly one greater than the signal count (one prologue arrive consumed by exactly one of the two waits).
…multicast - Option A prologue prefetch handshake bracketed with a cluster-scope barrier round (StreamK.py / KernelWriter.py StreamKMulticast plumb). - Remove the PGR<=1 enumeration guard so PGR2 multicast solutions build. - InsertClusterBarrierPass: drain the cooperative broadcast tensor group with both a pre-round drain and a producer-side drain (after the cooperative tensor_load group), gated on StreamKMulticast && PGR>=2. - PrefetchGlobalRead: [1,2] in the mxf4/mxf8 cluster-multicast configs plus a new PGR2 characterization test (+ snapshot + designed yaml). GPU-validated on gfx1250. JIRA ID : N/A
…ticast x K-split reduction) for gfx1250
Factor the 1-D gfx1250 StreamK HW cluster ClusterDim=[C,1] into two ORTHOGONAL
axes, C = Cs * Ck, selected by a new StreamKClusterKSplit (Ck) parameter:
* Cs = C // Ck spatial B-multicast peers -- process M-adjacent DISTINCT tiles
sharing the same B N-block over one K-slice (multicast along the s axis);
* Ck K-split reduction peers -- split one tile's K range and
reduce partials through the whole-cluster split barrier (reduction along k).
A single cluster now performs BOTH the spatial multicast and the K-split partial
reduction in one kernel. The within-cluster rank decodes "k fastest"
(k = StreamKIdx & (Ck-1), s = (StreamKIdx & (C-1)) >> log2(Ck)); the factored
B-multicast mask is maskB_base << k (Cs bits at stride Ck), with the self-only
fallback for a partial / M-unaligned cluster. The reduction reuses the existing
cluster-scope -3 barrier unchanged: it releases at the full C-membership arrival
count, which is benign over-synchronization across s-rows and keeps whole-kernel
signal/wait counts balanced (prologue arrive + mainloop lockstep + epilogue
reduce). Reduction owner/peer accounting falls out of skSplit = Ck.
Host: getSKGridImpl is unified to skGrid = roundUp(Ck*tiles, C); the two-tile
cluster-reduction kernarg block uses skSplit = Ck; ClusterDimCheck aligns on Cs
and ClusterReductionIterCheck balances Ck. New sizeMapping.streamKClusterKSplit
threads Ck to the host.
Mutual-exclusion relaxation: the multicast and cluster-reduction fast paths are
no longer mutually exclusive -- they are composable axes of one Cs x Ck
factoring. The collapse derives StreamKMulticast (Cs>1) and
StreamKClusterReduction (Ck>1) from (ClusterDim[0], StreamKClusterKSplit), and
the validators accept the combination; a new factoring validator enforces Ck |
C with Ck and Cs powers of two.
Degenerate collapse: Ck==1 (Cs==C) reproduces the shipped pure-multicast path
byte-for-byte; Ck==C (Cs==1) reproduces the shipped pure-reduction path
byte-for-byte (the legacy StreamKClusterReduction=1 opt-in is its Ck==C
degenerate).
Validation is CPU/snapshot only; real-HW / FFM validation of the combined axes
is deliberately deferred and gated behind multicast cooperative-load stability
at PrefetchGlobalRead>1. Adds factored CPU unit tests (decode, validation
matrix, predicate values, and degenerate byte-identity vs the pure paths) and a
gfx1250 codegen characterization test + designed config; regenerates only the
ValidParameters / SolutionClass roster goldens for the new parameter.
… the zero-iteration skip path The StreamKMulticast prologue emits a wave-0-elected cluster-scope arrive before the first cooperative-multicast load. Its only matching cluster-scope wait was the first-load wait, which sits after the last-iteration guard's long branch to PrefetchGlobalLastIterEnd. On the zero-full-iteration path (reachable when K is not a whole multiple of DepthU) that branch skips the first-load wait, leaving the prologue arrive unmatched and the cluster-scope barrier unbalanced on that edge. Add a matching cluster wait on the skip edge (guarded by StreamKMulticast) so the arrive is consumed on every control-flow path, preserving whole-cluster barrier symmetry. The wait branches over on scc0 (>=1 full iteration, where the first-load wait pairs the arrive) and executes an all-waves cluster wait on scc1 (zero iterations); it leaves scc intact so the subsequent long branch is unaffected. Update the two gfx1250 cluster char-test balance checks (multicast and coop load): the skip-edge wait and the first-load wait are mutually exclusive at runtime but both emitted statically, so the static cluster-wait count is now exactly one greater than the signal count (one prologue arrive consumed by exactly one of the two waits).
…multicast - Option A prologue prefetch handshake bracketed with a cluster-scope barrier round (StreamK.py / KernelWriter.py StreamKMulticast plumb). - Remove the PGR<=1 enumeration guard so PGR2 multicast solutions build. - InsertClusterBarrierPass: drain the cooperative broadcast tensor group with both a pre-round drain and a producer-side drain (after the cooperative tensor_load group), gated on StreamKMulticast && PGR>=2. - PrefetchGlobalRead: [1,2] in the mxf4/mxf8 cluster-multicast configs plus a new PGR2 characterization test (+ snapshot + designed yaml). GPU-validated on gfx1250. JIRA ID : N/A
jaopaulolc
force-pushed
the
users/jolabega/streamk-factored-cluster-mode
branch
from
July 22, 2026 17:28
ca47b80 to
61b867b
Compare
…ulticast Condense verbose prose that duplicates the design docs (ClusterLoad component docstrings, StreamK multicast prologue helpers, _validateStreamKMulticast) to concise summaries with a "See docs/design/..." pointer, and strip explanatory comments from the client/characterization test YAMLs (copyright/SPDX retained; parsed-YAML unchanged). No codegen change: char snapshots and targeted unit tests unchanged.
…eduction Condense verbose prose that duplicates the design docs (ClusterLoad docstrings, StreamK multicast/reduce cluster helpers, _validateStreamKMulticast / _validateStreamKClusterReduction) to concise summaries with a "See docs/design/..." pointer, and strip explanatory comments from the client/characterization test YAMLs (copyright/SPDX retained; parsed-YAML unchanged). No codegen change: char snapshots and targeted unit tests unchanged.
…cluster mode Condense verbose prose that duplicates the design docs (ClusterLoad docstrings, StreamK cluster helpers, the StreamKMulticast/Reduction/KSplit validators) to concise summaries with a "See docs/design/..." pointer, repoint stale citations of the untracked local plan doc to docs/design/streamk-wg-clusters.md, and strip explanatory comments from the client/characterization test YAMLs (copyright/SPDX retained; parsed-YAML unchanged). No codegen change: char snapshots and targeted unit tests (incl degenerate byte-identity) unchanged.
jaopaulolc
marked this pull request as ready for review
July 22, 2026 18:15
jaopaulolc
changed the title
…DPOnly configs Config-only additions closing gfx1250 StreamK coverage gaps (no codegen change): - Factored configs (sk_mxf4gemm_cluster_factored, sk_mxf8gemm_cluster_factored): add MacroTile MT128 (MI [16,16,128,1,1,4,4,2,2]), larger cluster shapes up to C=Cs*Ck=16 ([4,4],[2,8],[8,2]), K=4096, and (mxf4) an LDSTrInst [True,False] sweep. ProblemSizes chosen so every (ClusterDim, MacroTile) combo is cluster-aligned -> zero DID_NOT_SATISFY_ASSERTS (ClusterDimCheck ceil(M/MT0)%Cs==0 & ceil(N/MT1)%Ck==0, ClusterReductionIterCheck ceil(K/DepthU)%Ck==0). - Add sk_mxf8_2d_dual_multicast (F8 analogue of sk_mxf4_2d_dual_multicast) and sk_mxf8_force_dp_only (F8 analogue of sk_mxf4_force_dp_only). - Add PGR2 factored characterization: designed config streamk_factored_cluster_pgr2.yaml + char test + golden, mirroring the multicast PGR1/PGR2 pairing.
# Conflicts: # projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py # projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr
# Conflicts: # projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py # projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr
# Conflicts: # projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py # projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py # projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr
Decouple the cluster-scope barrier handshake from Multicast: a ClusterBarrier is now derived for ANY active cluster (ClusterDim != [1, 1] with TDM live and HasClusterBarrier), independent of whether the derived Multicast flag is on. Every cluster role -- multicast, partial reduction, factored and dual-2D -- needs its co-resident peers kept in lockstep, not just the B-multicast path (review request, jichangjichang). Also factor the StreamK=3 + ClusterDim StreamKMulticast derivation into a named _deriveStreamKMulticast helper so the Multicast / ClusterBarrier enable conditions stay readable. Behavior-preserving for the derivation (the StreamKMulticast value is unchanged). Update test_multicast_tristate::test_explicit_off to the new semantics: a force-off Multicast=0 on a [2, 2] cluster keeps ClusterBarrier on. On this branch the decouple is codegen-neutral: every clustered config in the characterization suite already derived ClusterBarrier via Multicast, so no golden moves.
Decouple the cluster-scope barrier handshake from Multicast: a ClusterBarrier is now derived for ANY active cluster (ClusterDim != [1, 1] with TDM live and HasClusterBarrier), independent of the derived Multicast flag (review request, jichangjichang, applied consistently across the stacked StreamK-cluster PRs). In particular a pure reduction cluster ([1, C], StreamKClusterReduction) derives Multicast=0 yet still needs the cluster barrier around its cooperative loads. This IS a codegen change on this branch: the reduction configs (ClusterDim=[1,2] and [1,4], StreamK=3, gfx1250) now emit the InsertClusterBarrierPass split barrier (s_barrier_signal/wait -3) in the cluster mainloop -> requires gfx1250 HW re-validation. Update test_multicast_tristate::test_explicit_off to the new semantics.
Decouple the cluster-scope barrier handshake from Multicast: a ClusterBarrier is now derived for ANY active cluster (ClusterDim != [1, 1] with TDM live and HasClusterBarrier), independent of the derived Multicast flag (review request, jichangjichang, applied consistently across the stacked StreamK-cluster PRs). A pure reduction cluster ([1, C]) and the reduction axis of a factored [Cs, Ck] cluster derive Multicast off yet still need the cluster barrier. This IS a codegen change on this branch: the pure-reduction configs (ClusterDim=[1,2] and [1,4], StreamK=3, gfx1250) now emit the InsertClusterBarrierPass split barrier (s_barrier_signal/wait -3) -> requires gfx1250 HW re-validation. The factored / dual-2D / forced-DP configs already derived ClusterBarrier via Multicast (no change). Update test_multicast_tristate::test_explicit_off to the new semantics.
The decoupled ClusterBarrier enabled the mainloop cluster split-barrier on ANY active cluster, including pure-reduction [1,C] (Cs=1). That path has no cooperative multicast tensor_load_to_lds to bracket, so the InsertClusterBarrier mainloop group emits an unmatched prologue s_barrier_wait -3 (Member N / Signal 0) that never completes -> FFM cluster deadlock (SIGABRT), and would hang real HW. The reduction is already synchronized by its own StreamK reduction -3 barriers. Gate the decouple on Cs = ClusterDim[0] > 1 so the mainloop barrier is emitted only where cooperative multicast loads exist ([C,1], [C,C], factored, dual-2D); pure-reduction [1,C] keeps ClusterBarrier False (pre-decouple, HW-green codegen). Spatial clusters (including Multicast forced off) are unaffected. Validated: FFM reduction [1,2] and [1,4] PASS (no deadlock); full CPU unit + characterization suite green; no golden movement.
The decoupled ClusterBarrier enabled the mainloop cluster split-barrier on ANY active cluster, including pure-reduction [1,C] (Cs=1). That path has no cooperative multicast tensor_load_to_lds to bracket, so the InsertClusterBarrier mainloop group emits an unmatched prologue s_barrier_wait -3 (Member N / Signal 0) that never completes -> FFM cluster deadlock (SIGABRT), and would hang real HW. The reduction is already synchronized by its own StreamK reduction -3 barriers. Gate the decouple on Cs = ClusterDim[0] > 1 so the mainloop barrier is emitted only where cooperative multicast loads exist ([C,1], [C,C], factored, dual-2D); pure-reduction [1,C] keeps ClusterBarrier False (pre-decouple, HW-green codegen). Spatial clusters (including Multicast forced off) are unaffected. Validated: FFM reduction [1,2] and [1,4] PASS (no deadlock); full CPU unit + characterization suite green; no golden movement.
The decoupled ClusterBarrier enabled the mainloop cluster split-barrier on ANY active cluster, including pure-reduction [1,C] (Cs=1). That path has no cooperative multicast tensor_load_to_lds to bracket, so the InsertClusterBarrier mainloop group emits an unmatched prologue s_barrier_wait -3 (Member N / Signal 0) that never completes -> FFM cluster deadlock (SIGABRT), and would hang real HW. The reduction is already synchronized by its own StreamK reduction -3 barriers. Gate the decouple on Cs = ClusterDim[0] > 1 so the mainloop barrier is emitted only where cooperative multicast loads exist ([C,1], [C,C], factored, dual-2D); pure-reduction [1,C] keeps ClusterBarrier False (pre-decouple, HW-green codegen). Spatial clusters (including Multicast forced off) are unaffected. Validated: FFM reduction [1,2] and [1,4] PASS (no deadlock); full CPU unit + characterization suite green; no golden movement.
PrefetchAcrossPersistent re-applies the TDM multicast mask on every persistent-loop iteration, but the mask SGPRs were freed in the prologue. Keeping the reuse valid pushed the StreamK [C,1] + StreamKForceDPOnly=0 kernel to 107 SGPRs (> the 106 budget), so KernelWriterAssembly replaced the body with an s_endpgm overflow stub and the output tensor D was left entirely unwritten -- exactly the PAP=1 + cluster + FDPO=0 failure. Keep the [C,1] broadcast mask (MulticastMaskB) live under PAP and free the self-only A mask (ClusterDim[1]==1 -> maskA==1, per-workgroup) whose re-application is a no-op, bringing the kernel back to 104 SGPRs. 2-D clusters (ClusterDim[1]>1, real A multicast) keep both masks live; PAP=0 and non-StreamK-multicast paths are byte-identical.
PrefetchAcrossPersistent re-applies the TDM multicast mask on every persistent-loop iteration, but the mask SGPRs were freed in the prologue. Keeping the reuse valid pushed the StreamK [C,1] + StreamKForceDPOnly=0 kernel to 107 SGPRs (> the 106 budget), so KernelWriterAssembly replaced the body with an s_endpgm overflow stub and the output tensor D was left entirely unwritten -- exactly the PAP=1 + cluster + FDPO=0 failure. Keep the [C,1] broadcast mask (MulticastMaskB) live under PAP and free the self-only A mask (ClusterDim[1]==1 -> maskA==1, per-workgroup) whose re-application is a no-op, bringing the kernel back to 104 SGPRs. 2-D clusters (ClusterDim[1]>1, real A multicast) keep both masks live; PAP=0 and non-StreamK-multicast paths are byte-identical.
PrefetchAcrossPersistent re-applies the TDM multicast mask on every persistent-loop iteration, but the mask SGPRs were freed in the prologue. Keeping the reuse valid pushed the StreamK [C,1] + StreamKForceDPOnly=0 kernel to 107 SGPRs (> the 106 budget), so KernelWriterAssembly replaced the body with an s_endpgm overflow stub and the output tensor D was left entirely unwritten -- exactly the PAP=1 + cluster + FDPO=0 failure. Keep the [C,1] broadcast mask (MulticastMaskB) live under PAP and free the self-only A mask (ClusterDim[1]==1 -> maskA==1, per-workgroup) whose re-application is a no-op, bringing the kernel back to 104 SGPRs. 2-D clusters (ClusterDim[1]>1, real A multicast) keep both masks live; PAP=0 and non-StreamK-multicast paths are byte-identical.
…op StreamKMulticast key
The StreamK=3 DP cooperative B-multicast fast path was tracked as a
derived-only internal state key (StreamKMulticast) plus a serialized C++
SizeMapping member. Both are redundant with ClusterDim, which is already
the source of truth and is serialized to the library.
Replace every read with a single derived helper
streamKMulticast(kernel) => kernel.get("StreamK",0)==3 and ClusterDim[0]>1
(Common/Utilities.py), and remove the derivation writes in Solution.py
(seed/derive/reset + the now-dead _deriveStreamKMulticast helper).
_validateStreamKMulticast now triggers off the helper.
C++: delete SizeMapping::streamKMulticast + its mapOptional (clean removal;
old libraries simply ignore the absent key) and rewrite the two grid guards
in ContractionSolution.cpp to the byte-exact equivalent
(streamK == 3 && clusterDim.x > 1).
Emitted assembly is byte-identical across the gfx1250 StreamK cluster
config matrix except the durable comment rename ("StreamKMulticast:" ->
"cluster B-multicast:", not snapshotted). Behavior-preserving: the only
golden movement is SolutionClass .ambr dropping StreamKMulticast from the
state-key list (num_keys 340 -> 339).
…op StreamKMulticast key
The StreamK=3 DP cooperative B-multicast fast path was tracked as a
derived-only internal state key (StreamKMulticast) plus a serialized C++
SizeMapping member. Both are redundant with ClusterDim, which is already
the source of truth and is serialized to the library.
Replace every read with a single derived helper
streamKMulticast(kernel) => kernel.get("StreamK",0)==3 and ClusterDim[0]>1
(Common/Utilities.py), and remove the derivation writes in Solution.py.
_validateStreamKMulticast and the reduction validator's factored-cluster
xor now trigger off the helper.
StreamKClusterReduction is deliberately kept exactly as-is (its own derived
state key, serialized C++ member, and the standalone reduction grid read).
C++: delete SizeMapping::streamKMulticast + its mapOptional (clean removal)
and rewrite the two grid guards to the byte-exact equivalent
(streamK == 3 && ... clusterDim), leaving the streamKClusterReduction terms
intact.
Emitted assembly is byte-identical across the gfx1250 StreamK cluster
config matrix (multicast + reduction) except the durable comment rename
("StreamKMulticast:" -> "cluster B-multicast:", not snapshotted). The only
golden movement is SolutionClass .ambr dropping StreamKMulticast from the
state-key list (num_keys 341 -> 340).
…ree self-only maskA on factored 2-D + PAP Part A -- drop the StreamKMulticast key (mirrors #9603/#9611): The StreamK=3 DP cooperative B-multicast fast path was tracked as a derived-only internal state key (StreamKMulticast) plus a serialized C++ SizeMapping member, both redundant with ClusterDim (already the source of truth and serialized). Replace every read with a single derived helper streamKMulticast(kernel) => kernel.get("StreamK",0)==3 and ClusterDim[0]>1 (Common/Utilities.py), remove the seed/derive/reset writes in Solution.py, and stop emitting the key to C++ (Contractions SizeMapping). The _validateStreamKMulticast validator and the derivation now trigger off the helper. StreamKClusterReduction is kept exactly as-is (its own derived key, serialized member, dual-2D caveat, and the standalone reduction grid read). C++: delete SizeMapping::streamKMulticast + its mapOptional (clean removal), and rewrite the two getSKGrid guards to the byte-exact equivalent (streamK == 3 && ...clusterDim), leaving the streamKClusterReduction terms intact. Part B -- factored self-only maskA guard refinement: ClusterLoad.papDropsSelfOnlyMaskA now mirrors computeMasks' aPeers==1 (papRefreshesMask and not streamKDual2DMulticast) instead of the conservative ClusterDim[1]==1. On a FACTORED 2-D cluster (Ck is a K-split reduction axis) A is per-workgroup, so the self-only maskA is freed under PAP -- factored 2-D + PAP + FDPO=0 now fits <=106 SGPRs (was 107 -> s_endpgm stub) and emits a real kernel. Dual-2D (FDPO=1 / StreamKDualMulticast) keeps both masks live (A is a real multicast). New designed + core PAP configs and _codegen char tests pin this (factored: maskA freed, sgpr=104; dual-2D: both masks live, sgpr=103; all err==0). Emitted assembly is byte-identical across the gfx1250 StreamK cluster matrix except the durable comment rename ("StreamKMulticast:" -> "cluster B-multicast:", not snapshotted) and the intended factored maskA drop. The only golden movement is SolutionClass .ambr dropping StreamKMulticast from the state-key list (num_keys 342 -> 341).
Port the full 2-D dual-multicast feature (both cluster axes broadcast A+B) onto the 1-D [C,1] B-multicast base so this branch owns all multicast modes: - Utilities: streamKDualMulticast / streamKForceDP2DMulticast / streamKDual2DMulticast detectors, derived from ClusterDim + StreamKDualMulticast/ForceDPOnly (no reference to the removed StreamKMulticast state key). - ValidParameters + GlobalParameters: register StreamKDualMulticast knob (default 0). - SolutionStructs/Solution: dual-2D derivation + validation branch; accept [Cs,Ck] both>1 as dual-multicast when the dual knob / ForceDPOnly is set; reject the factored (K-split) 2-D shape here. - Components/ClusterLoad: real A-multicast maskA when aPeers>1 (both-axes decode) shared machinery + cooperativeThreadPartition. - Components/StreamK: 2-D DP fold, DP->SK both-mask clear, dense 2-D mask emission and per-pass arrive for the standard-path dual-multicast. - Contractions + C++ (ContractionSolution.hpp / Serialization / getSKGridImpl): streamKDualMulticast serialized field + dual-multicast grid geometry (ForceDPOnly-2D and standard dual). - Configs/tests: sk_mxf4/8_2d_dual_multicast (FDPO=0), ForceDPOnly-2D configs (FDPO=1) incl. the 2-D+PAP home, designed char configs + _codegen char tests, and unit coverage. Existing 1-D goldens unchanged (additive-only).
…tion Port the full 2-D dual-multicast feature (both cluster axes broadcast A+B) onto the cluster-reduction branch so it stays a cumulative SUPERSET of the cluster-multicast branch (= 1-D [C,1] + 2-D dual-multicast + 1-D [1,C] reduction). Cherry-picked from the cluster-multicast port with the reduction machinery kept intact; the two 2-D shapes are made mutually exclusive: - Common: streamKForceDP2DMulticast / streamKDual2DMulticast detectors + StreamKDualMulticast knob (kept alongside streamKClusterFactors and the reduction derivation). - Solution derivation: StreamKClusterReduction is now derived only for a K-split cluster (Ck>1 AND NOT dual-2D); a 2-D dual cluster keeps reduction OFF so _validateStreamKMulticast accepts the [Cs,Ck] shape instead of rejecting it as factored. - SK cluster guard: [1,C] reduction and dual-2D ([Cs,Ck]+dual flag) both allowed; factored (both>1, no flag) still rejected. - StreamK.preLoop: dual-2D uses the M-fastest DP fold; the K-split reduction [1,C] keeps the WorkGroup0*Ck + WorkGroup1 fold (if/elif, mutually exclusive). - Contractions ClusterDimCheck: dual-2D keeps Ck as the N-tile divisor (value[4]=Ck); K-split reduction still pins value[4]=1. - C++ (getSKGridImpl / generateSingleCall): dual grid geometry (ForceDPOnly-2D tiles, standard-dual reshape) added ahead of the reduction grid override; streamKDualMulticast serialized field. - Configs/tests: dual-2D configs (FDPO=0/1), the 2-D+PAP home, designed char configs + _codegen char tests, and unit coverage. Reduction and 1-D multicast goldens unchanged (additive-only); the dual-2D nodes are new. No behavior change on the [1,C]/[C,1] paths.
# Conflicts: # projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py # shared/stinkytofu/include/stinkytofu/bindings/python/Module.hpp
…zes (gfx1250) Port #9690's clusterPadEarlyExit + computeMulticastMaskReduction onto the StreamK cluster-multicast path so non-multiple cluster launches work without ClusterDimCheck (now dropped). Covers 1-D [C,1] multicast and 2-D dual [Cs,Ck] (FDPO=1 ForceDP-2D and FDPO=0 StreamKDualMulticast); factored [Cs,Ck] (#9612) remains out of scope. - StreamK.py: streamKClusterPadEarlyExit() s_endpgm's padded boundary-cluster peers (WorkGroup0>=nWG0 or WorkGroup1>=nWG1*GSU) in the prologue BEFORE the first s_barrier_signal -3, so their WAVEDONE decrements the SQG barrier live-member count and the -3 cluster barrier still completes. - KernelWriterAssembly.py: gate computeMulticastMaskReduction to also fire on the 2-D dual multicast path so ld_bcst on a boundary cluster only waits on present peers. - ContractionSolution.cpp: round gridX up to Cs and gridY up to Ck for 2-D dual StreamK modes (legal HIP cluster launch); 1-D [C,1] unchanged. - Solution.py: correct stale ClusterDimCheck docstring (predicate dropped). - char tests: update the mask short-circuit marker to the reworded note. Divisible-size path is byte-identical (new logic guarded to boundary/partial).
…tes + unify 1D/2D multicast configs
The unified ClusterDim-driven configs (StreamKDualMulticast knob unset)
enable non-ForceDPOnly 2-D dual-operand cluster multicast purely via
ClusterDim=[Cs,Ck] both>1, but three host grid decisions were still gated on
the now-unset streamKDualMulticast knob, so the launched grid disagreed with
the kernel's ClusterDim-only 2-D fold. This produced ~50% wrong output values
on divisible grids and a client hang/SIGABRT on degenerate shapes.
Wrong values: gate the numWorkGroups reshape (generateSingleCall) and the
getSKGridImpl reshape on ClusterDim (clusterDim.x>1 && clusterDim.y>1) instead
of the knob, matching the kernel-side streamKDual2DMulticast detector, so the
host launches the [nWG0, gridY, batch] grid the 2-D fold expects (previously it
fell through to the 1-D [skGrid,1] grid, leaving most tiles unwritten).
Hang / strict-barrier SIGABRT ("wave ended with USER barrier signal_count
non-zero, member 4 signal 1" / reduction deadlock) on the degenerate case where
real N-tile rows < Ck (e.g. MT128x128 on 128x128, CD2_2 -> 1 tile, cluster
wants 4):
* getSKGridImpl: when (nWG1 / Ck) == 0, a full Ck A-multicast cluster row
cannot be formed from real tiles, so fall back to skGrid = tiles (one WG per
tile) instead of rounding up to a phantom Ck-row whose WG is pad-exited yet
still expected to contribute a K-partial.
* sk.grid final round-up to a multiple of Cs is 1-D-only: gate it on
clusterDim.y == 1. For 2-D dual the launch grid is padded to [Cs,Ck]
separately (generateSingleCall) and pad peers early-exit; rounding sk.grid
here re-inflated it above tiles, re-introducing the phantom K-split peer.
Host-only (ContractionSolution.cpp); no kernel codegen change. 1-D [C,1],
ForceDPOnly (1-D+2-D) and reduction paths are byte-identical. Validated on the
gfx1250 FFM simulator: FDPO=0/FDPO=1 mxf8+mxf4 cluster_multicast and isolated
2-D-dual sweeps (incl. batch and non-divisible sizes) are correct with no
hang/abort; CPU unit+char show no golden movement.
With ClusterDimCheck dropped, a boundary 1-D [C,1] B-multicast cluster (rounded-up grid) can contain no-work peers: padded tiles beyond the M extent (StreamKIdx >= totalTiles), or -- under the single-kernel tree K-split, where the rounded grid splits each tile's K across peers -- empty partials. Those peers used to fire the prologue cluster arrive (s_barrier_signal -3) and only exit at the later StreamK work-check, over-counting the split-barrier -> "signal for barrier (-3) that is already completed" (jitcu_barrier_state.h:180) abort at ClusterDim [8,1], 128x128, MT128x128 (K>DepthU). Fix: defer the 1-D multicast prologue arrive until AFTER the work-check (both the two-tile FDPO=0 path and the single-DP ForceDPOnly path) so only peers with real work signal -3; no-work peers s_endpgm at the work-check first, their WAVEDONE decrements the barrier's live-member count, and the surviving peers' first-load s_barrier_wait -3 completes exactly. This is correct for both reductions -- crucially it does NOT drop tree-K-split peers (which have StreamKIdx >= totalTiles yet real work), unlike a coord-based early-exit. Non-degenerate 1-D behavior is unchanged (every peer has work -> every peer arrives, just later in the prologue). The 2-D dual path keeps its existing streamKClusterPadEarlyExit.
… StreamKDualMulticast The StreamKDualMulticast valid parameter was added to the roster on this WIP branch without regenerating its characterization goldens (pre-existing drift). Regenerate only the affected nodes: the ValidParameters roster gains the additive 'StreamKDualMulticast' key ([0,1]) and the SolutionClass kernel names gain the _SKDM0 token (num_keys 340 -> 341). No behavioral change.
…barrier livelock on gfx1250 A boundary cluster whose real tile count is below the cluster width Cs rounds the 1-D StreamK launch grid up past `tiles`, K-splitting the lone tile across the Cs peers. When the tile's trailing summation chunk is a sub-DepthU partial (K % DepthU != 0) and the split is fine enough that the last peer receives only that partial (per-peer iters <= 1), that peer runs zero full DepthU iterations (numIterL == 0) and skips the per-iteration cluster split-barrier (-3) rounds its full-iter peers execute -- the survivors' `s_barrier_wait -3` never completes (deterministic livelock at MT256x256, ClusterDim [8,1], 256x256x1920). Host-only fix: decouple the cluster launch grid from the K-split degree. In solve(), when the degenerate partial-tail condition holds, keep sk.grid == tiles (one WG per tile owning the whole K-range incl. its own partial tail) instead of ceil(tiles/Cs)*Cs; generateSingleCall now rounds the 1-D launch grid up to Cs separately so the padded peers pad-exit via the normal StreamKIter work-check. Guarded to the tree-reduction 1-D pure-multicast boundary case, so non-degenerate 1-D, divisible-K, ForceDPOnly, 2-D dual and parallel-reduction paths stay byte-identical. No kernel/codegen change -> no golden movement. Validated on gfx1250 (FFM): the 4 shipped cluster_multicast configs pass 1128/1128 (0 FAILED / 0 incorrect / 0 hang / 0 abort), including the previously hanging MT256x256 CD8_1 256x256x1920.
…reduction Cascade the completed #9603 change-set (develop merge incl #9690 ClusterDimCheck drop, StreamK pad-exit + mask-reduction, refined ClusterBarrier gating, unified cluster_multicast configs, host-side grid fixes) onto the StreamK cluster REDUCTION feature (#9611), making #9611 a proper superset of the new #9603. Conflict resolution: - Predicates: ACCEPT ClusterDimCheck deletion (#9690); KEEP ClusterReductionIterCheck. - Solution.py: keep StreamKClusterReduction derivation ([1,C] Ck>1 and not dual2d); take #9603 refined ClusterBarrier gate (True for any active cluster, OFF for the 1-D reduction shape ClusterDim[0]==1 & [1]>1 & StreamK!=0 & FDPO==0); keep SK cluster guard that allows [C,1]/[1,C]/dual[Cs,Ck] and rejects factored. - ContractionSolution.cpp/.hpp + serialization: keep reduction kernarg/grid path guarded to streamKClusterReduction; unify multicast/reduction grid formula. - StreamK.py: keep dual-2D M-fastest vs reduction [1,C] K-split folds mutually exclusive; #9603 pad-exit/mask-reduction; #9603 prologue signal gating. - Configs: take #9603 unified cluster_multicast (1-D + 2-D dual); keep separate cluster_reduction configs. - Tests/goldens/docs: take #9603 multicast tests/snapshots/design-doc; keep the reduction-specific unit tests; snapshots to be regenerated + verified on gfx1250.
…ion key The merged cluster-reduction feature adds StreamKClusterReduction as a serialized SizeMapping state key, moving the solution-class char golden key count 341 -> 342. Regenerated ONLY the two affected nodes (test_solution_construction, test_mapping_interface); the sole diff is the new 'StreamKClusterReduction' entry.
Merge wip/streamk-cluster-reduction-9611 (5c97889) into the factored working branch (728bf88). Resolve conflicts keeping BOTH the factored cluster feature AND all #9611/#9603 infra + fixes: - Drop ClusterDimCheck (runtime pad-exit + mask reduction instead); keep ClusterReductionIterCheck (per-problem itersPerTile % Ck hard reject). - Refined ClusterBarrier gate: disabled only for pure reduction [1,C] (Cs==1); factored [Cs,Ck] with Cs>1 keeps the mainloop barrier. - Unified ClusterDim-driven grid: getSKGridImpl round-up roundUp(Ck*tiles, C) serves 1-D [C,1] multicast, [1,C] reduction, and [Cs,Ck] factored; dual-2D reshape gated on the StreamKDualMulticast knob to disambiguate factored. - generateSingleCall reduction branch (streamKClusterReduction && Ck>1) serves both pure reduction and factored; solve() 1-D degenerate-K-split liveness fix + dual-2D nwg1Ck<ck hang guard retained. - Factored derivation (StreamKClusterReduction = Ck>1 and not dual2D), decode/mask, and configs; StreamKDualMulticast kept solely for factored-vs-dual disambiguation. - Reduction validators accept factored [Cs,Ck] (no mutual-exclusion reject); shared _validateStreamK2DClusterShape (pow2, C in [2,16]). - Tests: keep factored-aware assertions (streamKMulticast helper), drop the obsolete mutual-exclusion tests. NOTE: SolutionClass golden (test_solution_class_char.ambr) len/num_keys/name carry interim values pending authoritative --snapshot-update regen.
After cascading local #9611 (reduction) into #9612 (factored), a knobless 2-D ClusterDim=[2,2] now derives as the FACTORED multicast+reduction path (Cs=2 B-broadcast + Ck=2 K-split) instead of dual-2D, so the #9611-era test assumptions about the unified cluster_multicast config no longer hold: - test_streamk_multicast / test_streamk_cluster_multicast_gfx1250_char: the factored [2,2] kernel carries no 1-D M-alignment predicate and no 2-D dual fold marker, so the negative "not 2-D marker" filter leaked it. Isolate the genuine 1-D [C,1] path positively (pin ClusterDim=[4,1] for emit / select on the "nWG0 aligned to C?" predicate). - test_streamk_dual_2d_multicast_gfx1250_char: dual-2D is now knob-gated (StreamKDualMulticast=1); repoint from the unified config to the dedicated streamk_dual_2d_multicast.yaml which sets the knob. - test_streamk_factored_cluster: ClusterDimCheck was dropped (runtime pad-exit + ClusterReductionIterCheck); assert the predicate is ABSENT rather than present. - SolutionClass golden: len/num_keys 341 -> 342 for the seeded derived-only StreamKClusterReduction key (targeted --snapshot-update, no name churn). CPU unit+char: all pass (761 snapshots). gfx1250 FFM validation of the factored, factored-pap, inherited multicast, force-DP-only and reduction configs: 0 FAILED / 0 abort / 0 hang across all shapes.
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
1 participant
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-3929
Motivation
The gfx1250 StreamK spatial B-multicast (#9603) and 1-D K-split cluster reduction (#9611) are two ways to use one hardware WG cluster. This PR adds the factored
[Cs,Ck]cluster mode that unifies them so one kernel can do both — B-multicast alongCsand K-reduction alongCk— on a genuine 2-D cluster (C = Cs * Ck).Stacked on #9611, and a cumulative superset of it. This branch contains everything in #9611 — the cluster-load component, 1-D
[C,1]B-multicast, all 2-D dual-operand multicast (ForceDPOnly-2D and standard-pathStreamKDualMulticast), the 1-D[1,C]cluster reduction, and the runtime pad-exit + mask-reduction alignment — plus the new factored mode.Technical Details
Incremental contribution over #9611:
[Cs,Ck]mode (multicast + cluster-reduction combined). A both-axes cluster[Cs,Ck]with no dual flag is interpreted as factored:StreamKMulticast = (Cs > 1)andStreamKClusterReduction = (Ck > 1)are both derived, so B is multicast alongCswhileCkpeers split and reduce K in the same kernel.StreamKDualMulticastdisambiguates factored-vs-dual. on a both-axes cluster,StreamKDualMulticast=1selects the dual-operand-multicast interpretation, while the default (0) selects the factored (multicast + reduction) interpretation — mutually exclusive by construction. (feat(tensilelite): StreamK cooperative cluster loads #9603/feat(tensilelite): add StreamK WG-cluster partial reduction for gfx1250 #9611 reject the non-dual both-axes shape; this PR accepts it.)maskB_base << k(Cs bits at stride Ck) with a self-only fallback for a partial / M-unaligned cluster; the reduction reuses the whole-cluster symmetric-3split barrier withskSplit = Ck.[Cs,Ck]passes both.Inherited unchanged from #9611/#9603: the cluster-load component, 1-D
[C,1]multicast, the 2-D dual-operand multicast, the 1-D[1,C]reduction, the symmetric cluster barrier, and the runtime pad-exit + mask-reduction alignment.Test Plan
--snapshot-update: inherited goldens byte-identical; new factored char nodes (incl. the PAP variant).[Cs,Ck]sweeps + inherited configs.Test Result
err == 0.Submission Checklist
Risk level
Medium.
[C,1]multicast,[1,C]reduction, and the 2-D dual multicast are byte-identical to #9611; the factored[Cs,Ck]combined mode is the new path (opt-in by cluster shape). Pending real gfx1250 hardware re-validation.