…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

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

…for zero-skip coverage

The mxf4/mxf8 cluster-reduction client YAMLs paired big tiles (MT64x64, MT128x128)
and high cluster factors (C up to 8) with low-M / low-K sizes, so 32 of 96
solution x size runs were rejected at selection time by ClusterDimCheck
(ceil(M/MT0) % C != 0) or ClusterReductionIterCheck (ceil(K/DepthU) % C != 0),
surfacing as DID_NOT_SATISFY_ASSERTS with no real coverage.

Redesign both configs so every (tile x C x size) combo runs:
- Drop MT128x128 (MT128 x C8 needs M a multiple of 1024 and is cluster-resource
  heavy); keep MT32x16 / MT32x32 / MT64x64.
- Sizes with M and N multiples of 512 and K a multiple of 2048 so
  itersPerTile = ceil(K/256) in {8,16} is divisible by C in {2,4,8} and
  ceil(M/MT0) is divisible by C for every kept tile:
  Exact [512,512,1,2048], [1024,512,1,2048], [512,512,1,4096], [512,1024,2,2048].
- Retain C in {2,4,8}, PrefetchGlobalRead [1,2], a multi-iter K (4096) and a
  batched size ([512,1024,2,2048]).

Result: 0 DID_NOT_SATISFY_ASSERTS (72 runs each, was 32/96), verified by
enumerating solutions via the char config harness and evaluating both host
predicates per solution x size, and end-to-end on the FFM client.

…r validator reject branches

Add unit coverage for previously-untested branches (codecov gaps): ClusterLoad
computeMasks sparse-metadata (Sparse==1/2) + undeclareSgprs metadata; and the
direct reject branches of _validateStreamKMulticast and
_validateStreamKClusterReduction that are unreachable through config derivation
(StreamK!=3, atomic, force-DP-only, XCC=3, non-[C,1]/non-pow2 cluster, non-gfx1250
ISA, missing HasTDM/HasClusterBarrier, TDMInst==0, multicast/reduction combo).

… to SCOPE_SYS

The gfx1250 cluster-reduction split-barrier handshake published peer partials
with a SCOPE_DEV global_wb and read them after a SCOPE_DEV global_inv. The split
barrier orders execution but carries no memory semantics, so on gfx1250's
partitioned L2 a SCOPE_DEV writeback was not guaranteed visible to a peer reading
on a different partition -> a rare (~0.1%, one boundary cluster) miscompute at the
largest grid (512,512,1,2048) on PGR1 (PGR2 masked it via prefetch timing).

Escalate only the cluster-reduction release (peer partial writeback) and the
owner's post-barrier acquire invalidate to SCOPE_SYS via a new optional `scope`
arg on releaseFence/acquireFence; the global-flag fallback and all non-cluster
StreamK keep SCOPE_DEV. InsertClusterBarrierPass.cpp (blob cf772a0) untouched.

Confirmed on real gfx1250: MXFP4 and MXFP8 cluster reduction pass (the PGR1
(512,512,1,2048) cases go FAILED->PASSED, no regression). Reduction char + unit
tests green (30 passed).

…rain

Reflow the producerDrainAnchors.push_back ternary to the repo clang-format style
(Google/IndentWidth 4/ColumnLimit 100). Whitespace only; no behavior change.

…sterDim-driven)

Derive the StreamK workgroup-cluster roles purely from the cluster shape
ClusterDim = [Cs, Ck] on StreamK=3 instead of a user opt-in:
  * StreamKMulticast        iff Cs = ClusterDim[0] > 1  ([C,1] = pure multicast)
  * StreamKClusterReduction iff Ck = ClusterDim[1] > 1  ([1,C] = pure reduction)

Remove the StreamKClusterReduction user/benchmark parameter (now a derived-only
internal state key, mirroring StreamKMulticast). Pure reduction migrates from the
1-D [C,1] cluster to a genuine 2-D [1,C] cluster: the launch grid becomes
[skGrid/Ck, Ck, 1] and the kernel folds the cluster Y rank into the linear index
(StreamKIdx = WorkGroup0*Ck + WorkGroup1), reusing the HW-validated 2-D mechanics.

The whole-cluster size C = Cs*Ck now drives the split-barrier span, the
ClusterReductionIterCheck predicate (Ck = ClusterDim[1]), the ClusterDimCheck
N-tile divisor (pinned to 1 for genuine 2-D clusters), and the host grid launch.

Pure multicast [C,1] is preserved byte-identically (Ck==1: no 2-D fold, 1-D
launch, C = ClusterDim[0]); the SCOPE_SYS reduction fence is unchanged.

Migrate reduction client + designed configs [C,1]+StreamKClusterReduction -> [1,C];
drop the now-invalid StreamKClusterReduction knob from multicast configs.
Regenerate reduction/ValidParameters/SolutionClass goldens (param removal +
[C,1]->[1,C] expression); multicast + StreamK codegen goldens stay byte-identical.

…rrier (no flags)

The gfx1250 StreamK [1,C] cluster reduction dropped exactly one peer's
partial for C >= 4 at the grid tail (device ~= (C-1)/C * reference, a
cluster-aligned column stripe confined to the last ~16 tiles / WG
960-1023). Root cause: the reduction cluster barrier handshake was
ASYMMETRIC -- each peer only issued s_barrier_signal -3 (arrive) and then
branched away / exited, while only the owner issued s_barrier_wait -3.
A peer-signal-only arrive is not a genuine cluster synchronization point
on gfx1250's partitioned L2: the owner could observe the barrier
satisfied and sum a peer slot before that cross-WGP peer's partial was
globally visible. Escalating the release/acquire fences to SCOPE_SYS did
not fix it because the defect was structural, not a missing drain.

Fix: mirror the HW-validated multicast handshake and make the reduction
barrier SYMMETRIC -- every cluster member (owner AND every peer) both
arrives (s_barrier_signal -3) AND waits (s_barrier_wait -3) on the same
cluster barrier, so no peer proceeds past the rendezvous until the whole
cluster has arrived and the owner's paired SCOPE_SYS acquire has a real
cross-cluster order to observe. Concretely: add the peer-side
clusterReduceWait after clusterReduceSignal in partialsWriteProcedure,
and remove the C <= 2 restriction in _streamKClusterBarrierFastEnabled so
the barrier path (not a per-peer global-flag fallback) carries
peer->owner synchronization for ALL C. The 64-bit slot-offset fix
(SMulHIU32 in computeWorkspaceSrd) and the SCOPE_SYS release/acquire on
the reduction path are retained. The per-tile global-flag handshake
remains ONLY for a cluster that straddles the SK grid boundary
(clusterReduceIntraCheck == SCC0); both owner and peers gate on that
identical uniform predicate, so a cluster never mixes barrier
participants with flag setters (deadlock invariant preserved).

Validation (gfx1250, physical GPU3), barrier-only kernels confirmed by
disassembly (s_barrier_signal/wait -3 present on owner and peer, peer
skips the flag store on the intra path, for both C=4 |3 and C=8 |7):
  256x512x4096 [1,4]:   1000/1000 pass, 0 fail, 0 hang
  512x512x4096 [1,4]:   1000/1000 pass, 0 fail, 0 hang
  512x512x4096 [1,8]:    500/500  pass, 0 fail, 0 hang
  1024x1024x4096 [1,4]:  500/500  pass, 0 fail, 0 hang
  1024x1024x4096 [1,8]:  500/500  pass, 0 fail, 0 hang
  hang check 512x1024 [1,8], 1024x1024 [1,8] (timeout 30): 0 hang
Full suites (fresh codegen from this source): sk_mxf8gemm_cluster_reduction
72/72 PASS, sk_mxf4gemm_cluster_reduction 72/72 PASS,
sk_mxf4gemm_cluster_multicast 296/296 PASS (multicast unregressed).
CPU unit tests: test_streamk_cluster_reduction 28/28, char emit test pass.

JIRA ID : N/A

…/factored

The ported symmetric-cluster-barrier fix folded a 64-bit workspace slot
offset (SMulHIU32 high word) into the SHARED computeWorkspaceSrd helper,
which the pure-multicast ([C,1]) and non-cluster StreamK paths also emit.
That perturbed the multicast kernel codegen (extra s_mul_hi_u32 + an SGPR
checkout that cascaded into register/schedule changes), so the multicast
asm was no longer byte-identical to the shipped baseline.

Gate the 64-bit high-word path on _streamKClusterReductionEnabled (the
same predicate every other hunk of the fix already uses): the
cluster-reduction / factored (derived Ck>1) path -- which drives the large
SK grids the 64-bit offset guards against, and is the only path this fix
targets -- keeps the HW-validated 64-bit codegen; the multicast and
non-cluster paths take the else branch and emit the original 32-bit offset
byte-for-byte. Verified: multicast STINKY_TOTAL_INST_BYTES unchanged and
the pre-vs-post multicast diff collapses to the emitter's nondeterministic
label/schedule floor (identical to a pre-vs-pre baseline), i.e. the
multicast instruction stream is byte-identical to the shipped path.

JIRA ID : N/A

Un-gate the 64-bit workspace slot-offset computation in the shared
computeWorkspaceSrd helper so the high-word (SMulHIU32) fold into SrdWS+1
is emitted for every StreamK path that addresses the partial-sum
workspace -- multicast ([C,1]), cluster reduction, factored, and
non-cluster StreamK -- not only the cluster-reduction/factored paths.

The overflow this guards against depends only on the per-slot tile
stride (MacroTile0*MacroTile1*bpe) and the StreamK slot count: the
partials workspace is partialTileSize == tileSize * skGrid regardless of
cluster mode, so the multicast path can silently wrap a 32-bit
slot*stride product on large problems exactly like the reduction path,
aliasing the wrong workspace slot. This intentionally changes the
multicast codegen (byte-identity with the prior 32-bit path is
deliberately dropped); the {basename, err} char goldens are unchanged
and still emit err==0. Real-HW validation pending (owner: user).

JIRA ID : N/A

…on diagnostic axis

- Extract StreamK._clusterElectArriveSignal, reused by clusterReduceSignal,
  streamKMulticastPrologueSignal + streamKMulticastProloguePrefetchHandshake
  (wait=True). Each call site keeps its own getNameInc label base / sgpr pool
  tag, so the emitted assembly is byte-identical to the inline copies.
- Remove dead ClusterLoadTDM.cooperativeThreadPartition + its test, and the
  unused clusterEnabled import in SubtileGREmit.py.
- Remove the stale post-derivation Multicast-coercion note in Solution.py.
- Prune test_streamk_multicast.py cases subsumed by other tests/goldens
  (TestEmit -> char golden; auto_enable/control_multicast -> tristate/baseline).
- Fix the reduction char-test docstring: pure K-reduction is ClusterDim=[1,C]
  (Ck=ClusterDim[1]) via streamk_cluster_reduction.yaml.
- Correct the ClusterReductionIterCheck host-side diagnostic (comment +
  debugEval message) from ClusterDim[0]/C to ClusterDim[1]/Ck to match the
  predicate (built from value[1]=Ck), and update the pinned Predicates_test.cpp
  string. Host-side diagnostic only; no kernel codegen impact.

No emitted-codegen change: streamk/cluster characterization goldens remain
byte-identical (verified without --snapshot-update).

Rewrite user-facing StreamK cluster validator reject messages to state the
supported configuration instead of git branch names, and remove branch-name
references from the surrounding comments and the ClusterDim shape-decoder
docstring. Trim bug-fix narrative from the cluster split-barrier gate docstring
down to the durable symmetric-barrier / SCOPE_SYS visibility / deadlock
invariant, and drop the "Held fix"/"red-pgr1" dangling references while keeping
the cross-L2 memory-ordering rationale. Convert before/after narration to
present tense, drop "byte-identical"/"HW-validated"/"lift" annotations (keeping
the durable fact each was attached to), and delete removed-test tombstones.
Rewrite both cluster design docs as concise as-shipped notes (reduction shape
[1,C], derived Ck, SCOPE_SYS symmetric barrier, boundary-clear present).

No emitted-codegen change: cluster/multicast/reduction char goldens are
byte-identical without --snapshot-update; CPU unit suite green.

# Conflicts:
#	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 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.

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.