Skip to content

feat(tensilelite): StreamK cooperative cluster loads#9603

Open
jaopaulolc wants to merge 10 commits into
developfrom
users/jolabega/streamk-cluster-multicast
Open

feat(tensilelite): StreamK cooperative cluster loads#9603
jaopaulolc wants to merge 10 commits into
developfrom
users/jolabega/streamk-cluster-multicast

Conversation

@jaopaulolc

@jaopaulolc jaopaulolc commented Jul 20, 2026

Copy link
Copy Markdown
Contributor

JIRA ID : AIHPBLAS-3929

Motivation

gfx1250 StreamK GEMMs that co-locate consecutive data-parallel (DP) workgroups in a 1-D WG cluster (ClusterDim = [C, 1]) process M-adjacent tiles that share the same B (N-block) over full K. Loading B once and TDM-multicasting it to the whole cluster (while A stays per-workgroup) removes redundant B traffic. This PR adds that cooperative-load fast path (StreamKMulticast) on top of the reusable ClusterLoad component + tri-state Multicast foundation from #9599.

Technical Details

  • Auto-derivation (collapse). StreamK==3 && ClusterDim!=[1,1] auto-enables StreamKMulticast in assignProblemIndependentDerivedParameters. It is a derived-only internal state key (like ClusterBarrier), never a YAML/benchmark parameter. The collapse forces Multicast on and enables the cluster-scope barrier handshake (ClusterBarrier). The bare index-only StreamK cluster state no longer exists.
  • Codegen (StreamK.py). New StreamKTwoTileDPFirst methods: streamKMulticastMaskPredicate (runtime clusterMulticastValid: nWG0 % C alignment + fully-populated cluster, else B loads normally) and streamKMulticastBoundaryClear (drops the B broadcast at the DP->SK boundary), plus the cluster-scope arrive/wait helpers that keep the InsertClusterBarrierPass signal/wait counts balanced: streamKMulticastPrologueSignal (wave-0 prologue s_barrier_signal -3), streamKMulticastProloguePrefetchHandshake (brackets the PGR>=2 double-buffer prologue prefetch load that the generic per-load bracketing's backward scan cannot reach), and streamKMulticastZeroIterClusterWait (consumes the prologue arrive on the zero-full-iteration skip path when K is not a whole multiple of DepthU).
  • PrefetchGlobalRead=2 (double-buffer) enabled. Alongside the Tensile-side prologue handshakes above, the native InsertClusterBarrierPass (shared/stinkytofu) emits a producer-side cooperative-load drain (s_wait_tensorcnt 0 immediately after the cooperative tensor_load_to_lds group, gated streamKMulticast && PGR>=2) so the peer-produced B broadcast is retired before the back edge and published coherently by the next loop-head barrier. This makes PGR2 numerically correct on the DP cooperative-multicast path (validated on gfx1250 B0). The multicast client configs sweep PrefetchGlobalRead: [1, 2].
  • KernelWriter.py. postMainLoopBarrierCheckAndReset preserves cluster-scope -3 barriers (only workgroup-scope barriers are rebuilt).
  • ClusterLoad.py. Re-adds the split-mask hooks (usesCombinedMask / maskSgprName) so StreamKMulticast binds MulticastMaskA (self) on A and MulticastMaskB (broadcast) on B.
  • Validation. _validateStreamKMulticast: SK3, ClusterDim=[C,1] (pow2 2..16), gfx1250 HasTDM/HasClusterBarrier/TDMInst=3, XCCMapping=0, not atomic; plus a Multicast==1 force-on hard-guard. SK4/SK5 + ClusterDim is rejected.
  • Host (ContractionSolution). streamKMulticast threaded through SizeMapping + C++ structs/serialization; grid multiple-of-C guard and a getSKGridImpl multicast override (ceil(tiles/C)*C, keeps every launched HW cluster full for the single DP round; trailing partial clusters disabled by the runtime predicate).

Test Plan

  • CPU unit: test_streamk_multicast.py, test_streamk_cluster_sk45_reject.py, test_multicast_tristate.py, test_cluster_load_component.py, test_multicast_legacy_coercion.py.
  • Snapshot char: full _codegen suite (non-multicast goldens byte-identical) + test_streamk_cluster_multicast_gfx1250_char.py, test_streamk_cluster_coop_load_gfx1250_char.py, test_streamk_cluster_multicast_pgr2_gfx1250_char.py.
  • gfx1250 C++ client GEMM roundtrip (PGR1 + PGR2): sk_mxf8gemm_cluster_multicast, sk_mxf4gemm_cluster_multicast, sk_mxf4_force_dp_only_cluster_multicast.

Test Result

  • Unit + char green; non-multicast _codegen goldens byte-identical. Multicast char kernels emit a balanced cluster barrier (s_barrier_signal -3 count == s_barrier_wait -3 count).
  • gfx1250 B0: MX-FP4 and MX-FP8 cooperative multicast pass at both PGR1 and PGR2 (PGR2 correctness fixed by the producer-side drain; no hangs, PGR1 unregressed).

Notes

Rebased onto latest develop. Comment and test-YAML footprint trimmed (verbose prose condensed to See docs/design/... pointers; the design doc docs/design/cluster-load-component-and-streamk-multicast.md is included on this PR). Cumulative over #9599.

Submission Checklist

Risk level

Low–moderate. StreamKMulticast is derived-only and gated to gfx1250 SK3 clusters; existing YAMLs derive byte-identically (non-multicast goldens unchanged). The PGR2 enablement adds a narrowly-gated (streamKMulticast && PGR>=2) producer-side drain in the shared cluster-barrier pass that is inert for every other path. Base is develop; the branch is cumulative over #9599.

Update — 64-bit workspace slot-offset fix (applies to the multicast path; HW validation pending)

Added the 64-bit workspace slot-offset computation to the shared computeWorkspaceSrd helper so every StreamK path that addresses the partial-sum workspace — including this PR's cluster multicast [C,1] path and non-cluster StreamK — emits a 64-bit offset instead of a 32-bit one. (This mirrors the now-universal fix landed on the stacked reduction/factored PRs #9611/#9612; on this branch the 64-bit hunk did not previously exist, so it is added here, un-guarded.)

Motivation — could the overflow happen on the multicast path? Yes.

computeWorkspaceSrd set SrdWS = AddressWS + offBytes * sPartialIdx with a 32-bit s_mul_i32, where offBytes = MacroTile0*MacroTile1*bpe (per-slot tile stride) and sPartialIdx ∈ [0, skGrid). The partials workspace is sized partialTileSize = offBytes * skGrid (host ContractionSolution::partialTileSize). The overflow depends only on the tile stride and the StreamK slot count, not on the cluster mode: the multicast [C,1] kernel emits and reads this exact workspace (partial-tile write via partialsWriteProcedure + owner fixup read), so for a large SK grid offBytes * sPartialIdx exceeds 2³², the 32-bit product silently wraps, and the peer-write / owner-read SRD aliases the wrong workspace slot. Concrete example: a 256×256 tile with fp32 partials is 256 KiB/slot, so skGrid ≳ 2³²/262144 = 16384 slots overflows a 32-bit offset (cluster C multiplies the effective slot count, so big-C multicast on large problems is the realistic trigger).

Technical Details

Compute the high word with s_mul_hi_u32 and fold it (plus the lo-add carry) into SrdWS+1 (s_mul_i32 low + s_mul_hi_u32 high + s_add_u32/s_addc_u32). Emitted for all StreamK paths; adds one SGPR + one s_mul_hi_u32 per workspace-SRD setup. This changes the multicast codegen (an extra instruction), which is intended.

Test Result (CPU, single-process; reused main worktree venv + editable rocisa .so)

  • Unit + char (streamk / cluster / multicast suites): 197 passed, 6 snapshots passed. The multicast (+ coop-load) {basename, err} char goldens are unchanged and still err == 0 (a {basename, err} digest does not move when only instructions are added), so no snapshot was regenerated.
  • Confirmed the emitted multicast [C,1] kernel now carries the 64-bit offset (s_mul_hi_u32 … // partials tile offset (high word) for 64-bit SRD folded into SrdWS+1).
  • Config zero-skip enumeration (gfx1250): streamk_cluster_multicast.yaml 4/4, streamk_cluster_multicast_pgr2.yaml 4/4, streamk.yaml 1/1 kernels → err == 0, 0 DID_NOT_SATISFY_ASSERTS, 64-bit SRD present in every kernel.

FFM/functional-sim is historically green regardless of HW correctness and is treated as a regression check only. Real gfx1250 hardware validation is still pending (owner: user).

@jaopaulolc
jaopaulolc requested a review from a team as a code owner July 20, 2026 18:36
@jaopaulolc jaopaulolc added the rocm:gemm algos Label to identify PRs of the GEMM Algorithms team. label Jul 20, 2026
@jaopaulolc
jaopaulolc changed the base branch from users/jolabega/clusterload-tristate-multicast to develop July 20, 2026 19:56
@therock-pr-bot

therock-pr-bot Bot commented Jul 20, 2026

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
🌿 Branch Name ✅ Pass
📝 PR Title/Description ✅ Pass
Forbidden Files ✅ Pass
🧪 Unit Test ✅ Pass
🚫 Draft PR 🔜 To Be Enabled
🚩 Feature Flag 🔜 To Be Enabled
📊 Code Coverage 🔜 To Be Enabled

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

@jaopaulolc jaopaulolc changed the title feat(tensilelite): StreamK cooperative cluster loads (StreamKMulticast) feat(tensilelite): StreamK cooperative cluster loads Jul 20, 2026
@codecov

codecov Bot commented Jul 20, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 85.13011% with 40 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...lt/tensilelite/Tensile/SolutionStructs/Solution.py 68.85% 11 Missing and 8 partials ⚠️
.../hipblaslt/tensilelite/src/ContractionSolution.cpp 0.00% 9 Missing and 1 partial ⚠️
...ipblaslt/tensilelite/Tensile/Components/StreamK.py 93.83% 2 Missing and 3 partials ⚠️
...lelite/Tensile/Components/Subtile/SubtileGREmit.py 0.00% 3 Missing ⚠️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py 84.62% 0 Missing and 2 partials ⚠️
...aslt/tensilelite/Tensile/Components/ClusterLoad.py 98.88% 0 Missing and 1 partial ⚠️

❌ Your project check has failed because the head coverage (76.84%) 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    #9603      +/-   ##
===========================================
+ Coverage    69.41%   69.61%   +0.20%     
===========================================
  Files         2770     2765       -5     
  Lines       451826   450822    -1004     
  Branches     66440    66350      -90     
===========================================
+ Hits        313602   313807     +205     
+ Misses      117205   116403     -802     
+ Partials     21019    20612     -407     
Flag Coverage Δ *Carryforward flag
TensileLite 34.39% <ø> (+0.04%) ⬆️ Carriedforward from 116f3bd
TensileLite-CPP 38.30% <0.00%> (-0.01%) ⬇️
TensileLite-Unit 65.21% <85.13%> (+0.08%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from 116f3bd
hipBLASLt 34.63% <ø> (ø)
hipCUB 82.68% <ø> (ø) Carriedforward from 116f3bd
hipDNN 86.34% <ø> (-<0.01%) ⬇️ Carriedforward from 116f3bd
hipFFT 48.90% <ø> (+1.99%) ⬆️ Carriedforward from 116f3bd
hipRAND 76.12% <ø> (ø) Carriedforward from 116f3bd
hipSOLVER 69.18% <ø> (ø) Carriedforward from 116f3bd
hipSPARSE 86.27% <ø> (ø) Carriedforward from 116f3bd
rocBLAS 47.95% <ø> (ø) Carriedforward from 116f3bd
rocFFT 48.51% <ø> (+1.43%) ⬆️ Carriedforward from 116f3bd
rocRAND 57.01% <ø> (ø) Carriedforward from 116f3bd
rocSOLVER 76.84% <ø> (-0.08%) ⬇️ Carriedforward from 116f3bd
rocSPARSE 72.49% <ø> (-0.09%) ⬇️ Carriedforward from 116f3bd
rocThrust 91.36% <ø> (ø) Carriedforward from 116f3bd

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...slt/tensilelite/Tensile/Common/GlobalParameters.py 98.02% <ø> (ø)
...aslt/tensilelite/Tensile/Common/ValidParameters.py 100.00% <ø> (ø)
...rojects/hipblaslt/tensilelite/Tensile/Component.py 95.56% <100.00%> (+0.03%) ⬆️
...pblaslt/tensilelite/Tensile/Components/__init__.py 100.00% <ø> (ø)
...ects/hipblaslt/tensilelite/Tensile/Contractions.py 87.68% <ø> (+1.47%) ⬆️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py 81.80% <100.00%> (+4.05%) ⬆️
...aslt/tensilelite/Tensile/SolutionStructs/Naming.py 96.97% <100.00%> (+0.02%) ⬆️
...aslt/tensilelite/Tensile/Components/ClusterLoad.py 98.88% <98.88%> (ø)
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py 80.53% <84.62%> (+1.61%) ⬆️
...lelite/Tensile/Components/Subtile/SubtileGREmit.py 70.65% <0.00%> (-0.12%) ⬇️
... and 3 more

... and 127 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@jaopaulolc
jaopaulolc force-pushed the users/jolabega/streamk-cluster-multicast branch from d5d9692 to e8b9c0c Compare July 20, 2026 21:51
@therock-pr-bot

Copy link
Copy Markdown

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

@therock-pr-bot

Copy link
Copy Markdown

Pre-commit check failed

pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

@jaopaulolc
jaopaulolc force-pushed the users/jolabega/streamk-cluster-multicast branch from 3db2ccf to 39772ee Compare July 22, 2026 15:48
…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
@jaopaulolc
jaopaulolc force-pushed the users/jolabega/streamk-cluster-multicast branch from 39772ee to 049da7b Compare July 22, 2026 17:28
…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.
@jaopaulolc
jaopaulolc marked this pull request as ready for review July 22, 2026 18:15
@jaopaulolc
jaopaulolc requested a review from a team as a code owner July 22, 2026 18:15
… multicast

Extend PrefetchGlobalRead to [1, 2] in sk_mxf4_force_dp_only_cluster_multicast.yaml
so the forced DP-only cooperative B-multicast path is exercised at PGR2 (the
B0-validated double-buffer broadcast + producer drain) as well as PGR1. No
validator guard blocked PGR2; solutions enumerate err==0.
…st reject branches

Add unit coverage for previously-untested branches surfaced by codecov:
- ClusterLoad.computeMasks sparse-metadata path (Sparse==1 follows-A, Sparse==2
  follows-B) and undeclareSgprs metadata SGPR free (ClusterLoad.py 90%->95%).
- _validateStreamKMulticast reject branches unreachable via config derivation
  (StreamK!=3, StreamKXCCMapping!=0, non-gfx1250 ISA, missing HasTDM, missing
  HasClusterBarrier), driven directly like test_accept_pgr2.
…rain

Reflow the producerDrainAnchors.push_back ternary to the repo clang-format style
(Google/IndentWidth 4/ColumnLimit 100). Whitespace only; no behavior change.
/// edge, so the publishing workgroup barrier at the next loop head correctly
/// orders it for the consuming waves. Matches PGR1, which drains each load in
/// its own iteration.
void insertProducerTensorDrainBefore(IRBase* anchor, AsmIRBuilder& irBuilder, GfxArchID archId) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@briannwu , could you help review joao's ClusterBarrier changes to see if that make sense to you?

The shared computeWorkspaceSrd helper addressed the StreamK partial-sum
workspace with a 32-bit SMulI32 slot*stride product. The per-slot stride
is MacroTile0*MacroTile1*bpe and the addressed slot index ranges over the
StreamK slot count (the partials workspace is partialTileSize ==
tileSize * skGrid, host ContractionSolution::partialTileSize), so for a
large SK grid the product exceeds 2^32, silently wraps, and the peer
write / owner read SRD aliases the wrong workspace slot.

The overflow depends only on the tile stride and the slot count, not on
the cluster mode, so the multicast [C,1] path -- which emits and reads
this exact workspace via partialsWriteProcedure / fixup -- can overflow
on large problems just like any other StreamK path. Compute the high
word with SMulHIU32 and fold it (plus the lo-add carry) into SrdWS+1,
emitted universally for every StreamK path (multicast, non-cluster, and
any future cluster mode). Adds one SGPR + one s_mul_hi_u32 per
workspace-SRD setup; the {basename, err} char goldens are unchanged and
still emit err==0. Real-HW validation pending (owner: user).

JIRA ID : N/A
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants