Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
27 commits
Select commit Hold shift + click to select a range
4e3c6bd
feat(tensilelite): add reusable ClusterLoad component and tri-state M…
jaopaulolc Jul 15, 2026
8cac9fe
feat(tensilelite): add StreamK cooperative cluster loads (StreamKMult…
jaopaulolc Jul 20, 2026
ef53e58
fix(tensilelite): consume StreamKMulticast prologue cluster arrive on…
jaopaulolc Jul 21, 2026
049da7b
feat(gfx1250): enable + fix PrefetchGlobalRead=2 for StreamK cluster …
jaopaulolc Jul 22, 2026
e4ba1c5
chore(tensilelite): trim comment/YAML footprint for StreamK cluster m…
jaopaulolc Jul 22, 2026
391eaa2
test(tensilelite): add PGR2 coverage to StreamK force-DP-only cluster…
jaopaulolc Jul 22, 2026
74eba56
test(tensilelite): cover ClusterLoad sparse-metadata + StreamKMultica…
jaopaulolc Jul 22, 2026
116f3bd
style(stinkytofu): satisfy clang-format in cluster-barrier producer-d…
jaopaulolc Jul 22, 2026
7fdc899
Merge branch 'develop' into users/jolabega/streamk-cluster-multicast
jaopaulolc Jul 24, 2026
3df4641
fix(streamk): use 64-bit workspace slot offset on all StreamK paths
jaopaulolc Jul 25, 2026
3908c5c
test(streamk): add mxf8 ForceDPOnly cluster multicast config
jaopaulolc Jul 27, 2026
98925c1
refactor(tensilelite): share cluster elect-arrive helper; prune subsu…
jaopaulolc Jul 27, 2026
ee8372e
docs(streamk): scrub session artifacts from comments, docstrings, and…
jaopaulolc Jul 28, 2026
bbe45da
docs(streamk): fix cluster-load design note to match shipped ClusterL…
jaopaulolc Jul 28, 2026
9832571
Merge remote-tracking branch 'origin/develop' into HEAD
jaopaulolc Jul 28, 2026
03d3442
fix(tensilelite): enable ClusterBarrier for any active cluster
jaopaulolc Jul 28, 2026
99fb570
fix(tensilelite): gate ClusterBarrier on Cs>1 (exclude pure-reduction)
jaopaulolc Jul 28, 2026
9d88b96
fix(tensilelite): keep StreamK multicast B mask live across PAP refresh
jaopaulolc Jul 29, 2026
1410870
refactor(tensilelite): derive StreamK B-multicast from ClusterDim, dr…
jaopaulolc Jul 29, 2026
0ea20e5
feat(streamk): add 2-D dual-operand multicast to cluster-multicast PR
jaopaulolc Jul 29, 2026
4aae1d6
Merge remote-tracking branch 'origin/develop' into HEAD
jaopaulolc Jul 29, 2026
ee37ffc
feat(streamk): pad-exit + mask-reduction for non-divisible cluster si…
jaopaulolc Jul 29, 2026
98552f8
refactor(streamk): ClusterBarrier gating + ClusterDim-direct 1D/2D ga…
jaopaulolc Jul 29, 2026
43eb9a7
fix(streamk): correct FDPO=0 2-D-dual cluster-multicast grid on gfx1250
jaopaulolc Jul 30, 2026
a3e7f96
fix(streamk): balance 1-D [C,1] cluster-multicast -3 barrier on gfx1250
jaopaulolc Jul 30, 2026
c0147e3
test(streamk): regenerate ValidParameters/SolutionClass snapshots for…
jaopaulolc Jul 30, 2026
2f67f81
fix(streamk): cap 1-D [C,1] multicast degenerate K-split to avoid -3 …
jaopaulolc Jul 30, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
179 changes: 179 additions & 0 deletions docs/design/cluster-load-component-and-streamk-multicast.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,179 @@
<!--
Copyright Advanced Micro Devices, Inc., or its affiliates.
SPDX-License-Identifier: MIT
-->

# Design: `ClusterLoad` Component + StreamK Cooperative Multicast Loads (gfx1250)

**Target arch:** gfx1250 (ISA `(12,5,0)`, wave32, TDM: `MXLoadInst=TDM` → `TDMInst=3`).

This note describes, as shipped, three coordinated pieces:

1. The **`ClusterLoad` component**, which owns the multicast ("cluster load") mask
machinery (value compute, `MulticastMask*` SGPR declare/undeclare, combined-vs-split
topology decision, per-load-site descriptor attach), replacing copies that lived in
`KernelWriter` / `KernelWriterAssembly` / `SubtileGREmit`.
2. **`Multicast` as an explicit tri-state** solution parameter, decoupling multicast from
`ClusterDim` so barrier-only clustering and cooperative-load clustering compose
independently.
3. **StreamK cooperative cluster loads** for the data-parallel (DP) region: a `[C,1]`
cluster of consecutive workgroups multicasts B across M-adjacent tiles.

The barrier-only cluster split-barrier reduction is described separately in
`docs/design/streamk-wg-clusters.md`.

---

## 1. Background

### 1.1 What a "cluster load" is

A cluster load is a TDM `tensor_load_to_lds` whose descriptor `Group1[word0]` has a
multicast-mask bit field OR'd in (`TensorDataMover.setMulticastMask`, a single `SOrB32`).
The hardware broadcasts the loaded tile to every workgroup in the cluster whose bit is set.

### 1.2 Mask value and the three topologies

The mask value is computed from the workgroup's position within the cluster and `ClusterDim`:

- `maskA = OR over idx in range(ClusterDim[1]) of (1 << (idx*ClusterDim[0]))`, shifted left
by `wg_x`: selects every workgroup sharing the same `wg_x` column (same M block) across all
`ClusterDim[1]` rows.
- `maskB = (1 << ClusterDim[0]) - 1`, shifted left by `wg_y * ClusterDim[0]`: selects every
workgroup in the same `wg_y` row (same N block).

Three name topologies (the selection matrix `ClusterLoad` must preserve):

| Topology | Predicate | SGPR name(s) |
|---|---|---|
| Combined single-parity | `tdmA and tdmB and NumWaves>1 and not UseSubtileImpl` | `MulticastMask` (even wave = maskA, odd wave = maskB, chosen by `WaveIdx` parity) |
| Split A/B | otherwise (subtile, single-tensor, or `NumWaves==1`) | `MulticastMaskA`, `MulticastMaskB` |
| Metadata (sparse) | `enableTDMMetadata` | `MulticastMaskMetadata` (follows A for `Sparse==1`, B for `Sparse==2`) |

Declare and undeclare mirror the same predicate.

### 1.3 Data-reuse fact (drives the StreamK design)

Two DP workgroups share **A** iff same WG0 (M) and overlapping K; share **B** iff same WG1
(N) and overlapping K. DP tiles compute the whole tile over full K, so K always overlaps.

- K-split peers of one tile (the `[C,1]` reduction cluster) have the same WG0/WG1 but
disjoint K → zero reuse; multicasting a reduction cluster multicasts nothing.
- Real reuse is spatial over a common full-K range: adjacent DP tiles. Consecutive-WG
clustering (`[C,1]`) gives M-adjacent tiles → shared B. (Tiles `nWG0` apart are N-adjacent
→ shared A; not reachable by consecutive `[C,1]` clustering.)

**Conclusion: in the StreamK DP region a `[C,1]` cluster multicasts B (not A).** This maps
onto the mask math with `wg_y=0, C0=C, C1=1`: `maskB = (1<<C)-1` (all C workgroups share B),
`maskA = 1<<wg_x` (self only → A per-workgroup).

### 1.4 StreamK addressing (where multicast pays off)

- `skIndexToWG` maps a tile → `(WorkGroup0, WorkGroup1, WorkGroup2)` with `tileID =
WG2*(nWG0*nWG1) + WG1*nWG0 + WG0` — **WG0 (M) fastest**, so consecutive tiles are
M-adjacent (same WG1/N block, consecutive WG0).
- In the first DP round (`graWorkGroup`, SK3 TwoTileDPFirst) consecutive workgroups map to
consecutive tiles → M-adjacent → share the same B over full K.
- The gfx1250 kernel-side WG remap `WorkGroup0 = cluster_x*nwg_x + wg_x` makes cluster `c`
own a contiguous `WorkGroup0` range `[c*C, c*C + C)`.

---

## 2. The `ClusterLoad` component

`ClusterLoad` is a capability-selected `Component` (`asmCaps = {"HasTDM": True}`,
`kernel = {"TDMInst": 3}`), found exactly like `TensorDataMoverLoad`: `ClusterLoad.find`
returns the TDM impl on gfx1250 and `None` (fallback → no multicast) elsewhere. It owns mask
value + declare/undeclare + topology decision + descriptor attach; it does **not** own
descriptor-group SGPRs, LDS offsets, or the `tensor_load_to_lds` itself.

| Method | Responsibility |
|---|---|
| `usesCombinedMask(kernel)` | The single combined-vs-split predicate. |
| `maskSgprName(kernel, tc, *, subtile=False, waveSeparated=False)` | Central name resolver (combined `"MulticastMask"` vs split `f"MulticastMask{strip_MXS(tc)}"` vs metadata). |
| `declareSgprs` / `undeclareSgprs` | Allocate / free the `MulticastMask*` SGPRs. |
| `computeMasks(writer, kernel, *, sgprWgX, sgprWgY, sgprNWgX, sTmp)` | Compute the mask value(s); the caller passes the SGPR operands it already holds. |
| `applyToDescriptor(writer, kernel, group1, tc, *, subtile=False)` | Gate + name choice + the `SOrB32`; empty `Module` when `not (kernel["Multicast"] and clusterEnabled(ClusterDim))`. |

**SGPR-operand contract:** `computeMasks` does not allocate SGPRs; it receives the operands
the surrounding code already holds and emits into them, so callers own SGPR lifetime.

---

## 3. `Multicast` tri-state

`Multicast` is an explicit tri-state solution parameter (`ValidParameters.py`):

- `-1` = auto: `ClusterDim != [1,1]` implies multicast, except StreamK cluster paths, which
drive it through the derived `StreamKMulticast` collapse instead;
- `0` = force off;
- `1` = force on (independent of `ClusterDim`).

Default `-1` reproduces the historic `ClusterDim`-coupled derivation, so configs that omit
`Multicast` are unchanged. `Multicast=1` requires a matching `ClusterLoadTDM` (`TDMInst=3` on
gfx1250 with `HasTDM`) and is rejected otherwise.

---

## 4. StreamK cooperative cluster loads (DP region)

### 4.1 Derivation and validation

`StreamKMulticast` is a derived-only internal state key (no `ValidParameters.py` entry). It
is auto-enabled when `ClusterDim != [1,1]` on SK3; a StreamK cluster always carries a
cooperative role. `_validateStreamKMulticast` hard-rejects a derived config it cannot
satisfy:

- `StreamK == 3` (SK3 DP schedule + `skIndexToWG` assumptions);
- `Multicast != 0` (the mask SGPRs are gated on `Multicast` while the predicate /
boundary-clear emitters are gated on `StreamKMulticast`; `Multicast=0` would reference
undeclared masks);
- `StreamKAtomic == 0` (the atomic path skips the workspace/tile DP structure);
- `StreamKXCCMapping == 0` (WGM/XCC remap is bypassed under clustering; XCC=3 overflows the
SGPR budget);
- `ClusterDim = [C,1]`, `C` a power of two with `C ∈ [2,16]`;
- gfx1250 with `HasTDM`, `TDMInst == 3`, and `HasClusterBarrier`.

The `nWG0 % C` runtime requirement is enforced by the `ClusterDimCheck` predicate at
selection time (not a silent fallback).

### 4.2 Mask derivation and runtime validity

For StreamK the authoritative WG coordinates are the ones produced by `skIndexToWG`. For the
DP `[C,1]` cluster `wg_x = StreamKIdx & (C-1)`, `wg_y = 0`, so `computeMasks` yields the
B-broadcast mask `(1<<C)-1` with A self-only.

For the C workgroups to genuinely share B on every issued load, the C tiles they process
must be M-adjacent (same WG1/N block, consecutive WG0). An init-time predicate
`clusterMulticastValid` (from `NumWorkGroups0` / total tiles / cluster base; cheap AND/compare
with `C` a power of two) gates the mask: `applyToDescriptor` ORs the mask only when the
predicate holds, otherwise the workgroup loads B normally. Incorrect multicast would write
the wrong B tile into a peer's LDS, so the gate is conservative.

### 4.3 DP→SK boundary clear

The static spatial mask is valid only for the DP (full-tile) round. At the DP→SK boundary
`streamKMulticastBoundaryClear` drops the mask to self-only (a single `SAndB32` on
`Group1[word0]`) so the SK partial-tile loads are not multicast; the SK tail reduces via the
existing workspace / global-flag path. Cooperative DP loads therefore work inside a full SK3
kernel, not only under `StreamKForceDPOnly`.

### 4.4 Host plumbing

- `sizeMapping.streamKMulticast` mirrors `streamKClusterReduction`.
- `getSKGridImpl` rounds `skGrid` up to a multiple of `C`; tail workgroups (if
`totalTiles % C != 0`) are idle-but-present and the init-time predicate disables their
cluster's multicast.
- The launch `enableCluster` branch already keeps the grid cluster-friendly and sets
`rv.clusterDim`; no launch change beyond threading `streamKMulticast`.

---

## 5. Interaction notes

- gfx1250 StreamK already uses TDM (`MXLoadInst=TDM` → `TDMInst=3`), so `HasTDM` is the
natural cap; no new transport is introduced.
- `StreamKAtomic` is incompatible (the atomic path skips the workspace/tile structure the DP
schedule relies on) and is rejected.
- Consecutive `[C,1]` clustering shares B, not A; A stays per-workgroup. A-multicast needs an
N-strided / 2-D cluster.
Original file line number Diff line number Diff line change
Expand Up @@ -576,10 +576,14 @@
{"MIArchVgpr": [False]},
{"StreamK": [0]},
{"StreamKForceDPOnly": [0]},
{"StreamKDualMulticast": [0]},
{"StreamKAtomic": [0]},
{"StreamKWorkStealing": [0]},
{"StreamKXCCMapping": [0]},
{"StreamKFixupTreeReduction": [0]},
# NOTE: StreamKMulticast is a derived-only internal state key (like
# ClusterBarrier), auto-enabled by Solution.py for StreamK==3 + ClusterDim
# clusters; it is deliberately NOT a benchmark/default parameter here.
{"DebugStreamK": [0]},
{"DebugPersistentKernelLoopForever": [False]},
{"ActivationFused": [True]},
Expand Down Expand Up @@ -624,6 +628,9 @@
# [1, 1] disables clustering. Non-[1, 1] enables Multicast so workgroups within
# a cluster can share data loaded via TDM-multicast, reducing redundant global reads.
{"ClusterDim": [[1, 1]]},
# Multicast tri-state (see ValidParameters). Default -1 = legacy auto, so
# every existing YAML (which omits Multicast) derives identically to before.
{"Multicast": [-1]},
{"HalfPLR": [0]},
{"TDMIterateMode": [0]}
]
Expand Down
60 changes: 60 additions & 0 deletions projects/hipblaslt/tensilelite/Tensile/Common/Utilities.py
Original file line number Diff line number Diff line change
Expand Up @@ -365,6 +365,66 @@ def clusterEnabled(clusterDim):
"""True when a workgroup cluster is requested (ClusterDim [x, y] is not [1, 1])."""
return (clusterDim[0] * clusterDim[1]) != 1

def streamKMulticast(d):
"""True when the StreamK=3 DP cooperative B-multicast fast path is active.

Single source of truth derived from ClusterDim: on StreamK=3 a spatial
cluster (ClusterDim[0] = Cs > 1, i.e. Cs peers sharing B across M-adjacent
DP tiles) IS the cooperative B-multicast path. This replaces the former
derived/serialized ``StreamKMulticast`` state key -- ClusterDim is now the
only source of truth, so there is nothing extra to store or serialize.

``d`` may be a kernel or a solution ``state`` dict; both expose "StreamK"
and "ClusterDim". Uses ``.get`` for partial-state derivation call sites that
construct a dict without a StreamK / ClusterDim key.
See docs/design/cluster-load-component-and-streamk-multicast.md.
"""
return d.get("StreamK", 0) == 3 and d.get("ClusterDim", [1, 1])[0] > 1

def streamKForceDP2DMulticast(d):
"""True for the ForceDPOnly 2-D dual-operand multicast cluster.

This is a StreamK==3 ``StreamKForceDPOnly`` (dense data-parallel, no K-split
reduction) kernel given a GENUINE 2-D cluster ClusterDim = [Cs, Ck] with BOTH
axes > 1. Unlike the factored K-split cluster (where Ck is a reduction axis),
here the Ck (Y) axis maps to N-ADJACENT output tiles so the Y-peers reuse the
A operand (A-multicast), while the Cs (X) peers reuse B on M-adjacent tiles
exactly as in the shipped 1-D [C,1] ForceDPOnly multicast. Both operands are
multicast via the DENSE ClusterLoad 2-D masks.

Detected purely structurally (ForceDPOnly + ClusterDim[0]>1 + ClusterDim[1]>1).
``d`` may be a kernel or a solution ``state`` dict; both expose
"StreamKForceDPOnly" and "ClusterDim".
See docs/design/streamk-wg-clusters.md.
"""
return bool(d.get("StreamKForceDPOnly", 0)) \
and d["ClusterDim"][0] > 1 and d["ClusterDim"][1] > 1

def streamKDual2DMulticast(d):
"""True for a 2-D DUAL-operand multicast cluster (ClusterDim-derived).

A GENUINE 2-D cluster ClusterDim = [Cs, Ck] with BOTH axes > 1 on the StreamK
DP multicast path: the DP (full-tile) round does the 2-D dual multicast (Cs/X
peers share B on M-adjacent tiles, Ck/Y peers share A on N-adjacent tiles); the
standard (StreamKForceDPOnly=0) path's SK (partial-tile) round reduces 1-D via
the workspace as today, while ForceDPOnly=1 is a single dense DP round.

The 1-D-vs-2-D DISTINCTION is made PURELY on ClusterDim (Ck = ClusterDim[1] > 1),
NOT on the StreamKForceDPOnly / StreamKDualMulticast knobs: on the multicast PR
(#9603) there is no factored K-split cluster, so a [Cs, Ck] with both axes > 1
is unambiguously the dual-operand multicast. (The factored-vs-dual
disambiguation via the StreamKDualMulticast knob -- where a 2-D cluster CAN be
a K-split reduction with Ck a reduction axis -- is re-introduced by the factored
PR #9612; StreamKForceDPOnly stays a MODE axis distinguishing the dense DP-only
schedule from the two-tile DP+SK schedule, not a dimensionality gate.)

``d`` may be a kernel or a solution ``state`` dict; uses ``.get`` for
partial-state derivation call sites.
See docs/design/streamk-wg-clusters.md.
"""
clusterDim = d.get("ClusterDim", [1, 1])
return clusterDim[0] > 1 and clusterDim[1] > 1

def log2(x):
return int(log(x, 2) + 0.5)

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -901,6 +901,15 @@ def makeValidMatrixInstructions():
# 0: uses workspace to store partial tiles, accumulate in deterministic fix-up step
# 1: uses atomics to accumulate partial tiles
"StreamKAtomic": [0, 1],
# Opt-in for 2-D DUAL-operand multicast on the STANDARD two-tile StreamK
# path (StreamKForceDPOnly=0). On a genuine 2-D cluster ClusterDim=[Cs,Ck]
# (both > 1) the DP (full-tile) round does 2-D dual multicast (Cs/X peers
# share B on M-adjacent tiles, Ck/Y peers share A on N-adjacent tiles) and
# the SK (partial-tile) round reduces 1-D via the workspace as today. Ck is
# an N-tiling / A-multicast axis, NOT a K-split reduction axis. ForceDPOnly-2D
# does not need it (that shape is unambiguous). Limited to batch==1. See
# streamKDual2DMulticast and docs/design/streamk-wg-clusters.md.
"StreamKDualMulticast": [0, 1],
# Codegen-time toggle for single-hop next-neighbor work stealing in the
# dynamic-queue StreamK fetch (SK4 / SK5-dynamic). Queue count =
# archCaps['NumXCD'] (8 on gfx942/gfx950). When a workgroup's home queue
Expand All @@ -918,6 +927,14 @@ def makeValidMatrixInstructions():
# 0: use linear reduction
# 1: use tree reduction
"StreamKFixupTreeReduction": [0, 1],
# NOTE: the gfx1250 StreamK DP cooperative cluster-load / TDM B-multicast
# fast path is intentionally NOT a valid/benchmark parameter and is no longer
# a state key at all. It is derived on demand from ClusterDim via the
# streamKMulticast(state) helper (StreamK==3 and ClusterDim[0] > 1);
# _validateStreamKMulticast hard-rejects any cluster config that cannot
# satisfy its constraints. ClusterDim is the single source of truth and is
# what serializes to C++ (SizeMapping.clusterDim); there is no separate
# streamKMulticast key to store or fork on.
# Debug settings for stream-k kernels to disable parts of the kernel
# Bit 0: Don't generate fixup code
# Bit 1: Don't generate write to partials code
Expand Down Expand Up @@ -1199,6 +1216,14 @@ def makeValidMatrixInstructions():
# Cluster dimension. Clusters have up to 16 work-groups in a cluster, but each work-group in a
# cluster runs on a separate WGP.
"ClusterDim": validClusterDimensions,
# Multicast (cluster load) control. Decouples the TDM-multicast opt-in from
# ClusterDim so barrier-only clustering and cooperative-load clustering
# compose independently. Tri-state:
# -1 = auto (legacy: ClusterDim != [1,1] implies Multicast, except the
# StreamK cluster paths); reproduces historic behavior exactly.
# 0 = force multicast off.
# 1 = force multicast on (independent of the ClusterDim coupling).
"Multicast": [-1, 0, 1],
# Enable PLR 0.5 to save vgprs
# 0: Disabled
# 1: Use PLR 0.5 for A
Expand Down
5 changes: 5 additions & 0 deletions projects/hipblaslt/tensilelite/Tensile/Component.py
Original file line number Diff line number Diff line change
Expand Up @@ -312,6 +312,11 @@ class GL2Prefetch(Component):
GL2 Prefetch
"""

class ClusterLoad(Component):
"""
Cluster (multicast) TDM load: multicast-mask compute + descriptor attach.
"""

# Importing here allows auto-registry of components in the Components directory.
# Each file must be listed in __all__ in Components/__init__.py
# "noqa" prevents linter from complaining here.
Expand Down
Loading
Loading