diff --git a/docs/design/cluster-load-component-and-streamk-multicast.md b/docs/design/cluster-load-component-and-streamk-multicast.md new file mode 100644 index 000000000000..89e1b8c05cd7 --- /dev/null +++ b/docs/design/cluster-load-component-and-streamk-multicast.md @@ -0,0 +1,561 @@ + + +# Design: `ClusterLoad` Component + StreamK Cooperative Multicast Loads (gfx1250) + +**Target arch:** gfx1250 (ISA `(12,5,0)`, wave32, TDM: `MXLoadInst=TDM` → `TDMInst=3`). +**Scope:** Two coordinated pieces of work. + +1. **`ClusterLoad` component + behavior-preserving refactor.** Extract the multicast + ("cluster load") mask machinery — value computation, SGPR declare/undeclare, and the + three descriptor apply sites (dense, wave-separated, subtile) — into a single reusable + `ClusterLoad` tensilelite `Component`, and refactor the existing subtile Multicast path + onto it. This is a **pure refactor**: zero change to emitted assembly for existing + configs, gated by codegen snapshots. +2. **Decouple `Multicast` from `ClusterDim`** into an independent opt-in, so barrier-only + clustering (the separate cluster split-barrier reduction feature) and cooperative-load + clustering compose independently instead of `ClusterDim != [1,1]` unconditionally forcing + `Multicast=True`. +3. **StreamK cooperative cluster loads.** Build TDM multicast on top of `ClusterLoad` for + the StreamK **data-parallel (DP) region**, resolving the tension between the spatial + `[C0,C1]` cluster cooperative loads want and the `[C,1]` same-tile K-peer cluster the + shipped reduction wants. + +The separate barrier-only cluster reduction (`docs/design/streamk-wg-clusters.md`) is +orthogonal and tracked independently. + +Source-line anchors reference the code as it exists now and were verified against the tree. + +--- + +## 1. Background (verified anchors) + +### 1.1 What a "cluster load" is + +A cluster load is a normal TDM `tensor_load_to_lds` whose descriptor `Group1[word0]` has a +multicast-mask bit field OR'd in. The HW then broadcasts (multicasts) the loaded tile to +every workgroup in the cluster whose bit is set. Attachment is a single `SOrB32`: + +```511:514:projects/hipblaslt/tensilelite/Tensile/Components/TensorDataMover.py + def setMulticastMask(self, group1: int | str, mask: str, writer: "KernelWriterAssembly") -> Module: + mod = Module() + mod.add(SOrB32(sgpr(f"{group1}"), sgpr(f"{group1}"), sgpr(f"{mask}"))) + return mod +``` + +### 1.2 Mask value computation and the three topologies + +The mask *value* is computed once in `defineAndResources` +(`KernelWriterAssembly.py:2646-2694`), from the WG's position within the cluster and +`ClusterDim`: + +- `maskA = OR over idx in range(ClusterDim[1]) of (1 << (idx*ClusterDim[0]))`, then shifted + left by `wg_x`. Bit `wg_y*ClusterDim[0] + wg_x` = the cluster-linear index of the WG. + So `maskA` selects every WG sharing the same `wg_x` column (same M block) across all + `ClusterDim[1]` rows. +- `maskB = (1 << ClusterDim[0]) - 1`, then shifted left by `wg_y * ClusterDim[0]`. Selects + every WG in the same `wg_y` row (same N block). + +Three name topologies (this exact selection matrix must be preserved by the refactor): + +| 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 (`KernelWriter.py:9163-9176`) and undeclare (`KernelWriter.py:2838-2848`) mirror the +same predicate. + +### 1.3 The three descriptor apply sites (the duplication to unify) + +All three OR the mask into `Group1` under the same gate +`kernel["Multicast"] and enableCluster`, where `enableCluster = prod(ClusterDim) != 1`: + +- **Dense** `initTDMDescriptor` — `KernelWriterAssembly.py:18901-18902`; local + `maskSgprName(tc)` returns `MulticastMask{A|B}` (split). +- **Wave-separated** `initTDMDescriptorWaveSeparatedImpl` — `KernelWriterAssembly.py:19059-19060`; + local `maskSgprName` returns the combined `"MulticastMask"`. +- **Subtile** `initTDMDescriptorSubtile` — `Components/Subtile/SubtileGREmit.py:1108-1113`; + uses `MulticastMask{tc}` (split). + +`initTDMDescriptorSubtile` (`SubtileGREmit.py:1061-1155`) is a near-clone of the writer +`initTDMDescriptor` (`KernelWriterAssembly.py:18843-18994`); the mask attachment is the +piece we lift into `ClusterLoad`. The rest of the descriptor build (LDS offset, tensor +dims/strides, padding) stays where it is — `ClusterLoad` does **not** own descriptor-group +allocation or LDS offsets. + +### 1.4 Cooperative-thread partition (shared math) + +`numCooperativeWGs = ClusterDim[1] for A / ClusterDim[0] for B`, duplicated in +`Components/GL2Prefetch.py:26` and `:78`. `ClusterLoad` centralizes this as +`cooperativeThreadPartition`. + +### 1.5 Component framework + +- Auto-registration metaclass: `Component.py:113-124` (`__init__` sets `implementations` + and `setattr` on each base). `matches()` partial-matches `asmCaps`/`archCaps`/`kernel` + (`Component.py:132-144`); `find()` requires exactly one match, raises on >1, returns + `None` on 0 → fallback (`Component.py:167-177`). +- Categories (abstract, no `__call__`) declared near `Component.py:305-313` + (`TensorDataMover`, `GL2Prefetch`); `from .Components import *` at `:318`. +- `Components/__init__.py:28-55` `__all__` lists each component module. +- `TensorDataMoverLoad` shows the concrete pattern: `asmCaps = {"HasTDM": True}`, + `kernel = {"TDMInst": 3}`, method-bag returning rocisa `Module`s + (`TensorDataMover.py:13-15`). +- New-file SPDX header (`# Copyright Advanced Micro Devices, Inc., or its affiliates.` / + `# SPDX-License-Identifier: MIT`). + +### 1.6 `Multicast` / `ClusterDim` coupling (current) + +`Solution.py:1046-1064` forces `Multicast=True` for any `ClusterDim != [1,1]`, except it is +already suppressed for Stream-K. `Multicast` is **not** a user parameter in +`ValidParameters.py` (it is a derived state var). + +### 1.7 StreamK addressing (where multicast pays off) + +- `skTileIndex` (`StreamK.py:439`, math `469-483`) magic-divides `StreamKIter` into a tile + index and `StreamKLocalStart`/`StreamKLocalEnd` (iter offsets within the tile). +- `skIndexToWG` (`StreamK.py:487-503`) maps tile → `(WorkGroup0, WorkGroup1, WorkGroup2)` + with the linearization `tileID = WG2*(nWG0*nWG1) + WG1*nWG0 + WG0` — **WG0 (M) is + fastest**, so consecutive tiles are M-adjacent (same WG1/N block, consecutive WG0). +- The K-direction StreamK offset (`StreamKLocalStart*DepthU*strideL`) is added in + `computeLoadSrdCommon` (`StreamK.py:548`, `555-562`) and `graAddressesCommon` + (`StreamK.py:629`, `636-645`); both are `tc`-generic (shared A/B). **The `WorkGroup0*MT0` + / `WorkGroup1*MT1` tile-base term lives in the base KernelWriter GRA/SRD code, keyed on + the post-`skIndexToWG` `WorkGroup0/1`.** A/B thus bind to WG0/M and WG1/N respectively; + the K offset is identical for A and B. +- `graWorkGroup` (SK3 TwoTileDPFirst, `StreamK.py:2846`; DP shift `2921-2925`): a WG `w` + processes tiles `w, w+skGrid, w+2*skGrid, …`. In the **first DP round** consecutive WGs + map to consecutive tiles → **M-adjacent → share the same B (N-block) over full K**. +- `StreamKIdx = WorkGroup0` (`StreamK.py:2658-2663`), preceded by the `ttmp9` raw-wgid + workaround (`StreamK.py:2645-2656`), which is **already skipped under any cluster + (`ClusterDim != [1,1]`)** so the cluster-remapped `WorkGroup0` survives. +- DP/SK region split in `graWorkGroup` at `StreamK.py:2907-2932`; the "started & + finished the whole tile → regular store, no reduction" branch at `StreamK.py:922-933`. +- gfx1250 kernel-side WG remap: `WorkGroup0 = cluster_x*nwg_x + wg_x` + (`KernelWriterAssembly.py:2628-2632`), so cluster `c` owns a contiguous `WorkGroup0` + range `[c*C, c*C + C)`. + +### 1.8 Data-reuse fact (drives the whole StreamK design) + +Two DP WGs 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 shipped `[C,1]` reduction cluster): same WG0/WG1 but + **disjoint K** → **zero reuse**. Multicasting the 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 +exactly onto the existing mask math with `wg_y=0, C0=C, C1=1`: `maskB = (1< bool` | The single predicate `tdmA and tdmB and NumWaves>1 and not UseSubtileImpl`. Kills the 3 duplicated copies (`KernelWriter.py:9170`, `:2842`, `KernelWriterAssembly.py:2668`). | +| `maskSgprName` | `(kernel, tc, *, subtile=False) -> str` | Central name resolver: combined `"MulticastMask"` when `usesCombinedMask` and not `subtile`; else `f"MulticastMask{strip_MXS(tc)}"`; `"MulticastMaskMetadata"` for metadata. Subsumes the two local `maskSgprName` closures + the subtile literal. | +| `declareSgprs` | `(writer, kernel) -> None` | Moves `KernelWriter.py:9163-9176` (uses `writer.defineSgpr`). | +| `undeclareSgprs` | `(writer, kernel) -> Module` | Moves `KernelWriter.py:2838-2848` (uses `writer.undefineSgpr`). | +| `computeMasks` | `(writer, kernel, *, sgprWgX, sgprWgY, sgprNWgX, sTmp) -> Module` | Moves `KernelWriterAssembly.py:2646-2694` verbatim, including the wave-parity branch and the metadata cases. Takes the exact SGPR operands the caller already holds so emitted asm is byte-identical. | +| `applyToDescriptor` | `(writer, kernel, group1, tc, *, subtile=False) -> Module` | Folds the gate + name choice + the `SOrB32`. Returns an **empty `Module`** when `not (kernel["Multicast"] and enableCluster)` — identical to today's skipped `if`. Internally calls `TensorDataMover.setMulticastMask`. Replaces the 3 apply sites. | +| `cooperativeThreadPartition` | `(kernel, tc) -> int` | `ClusterDim[1] if tc-ends-A else ClusterDim[0]`. Shared with `GL2Prefetch` (`:26`, `:78`). | + +Rationale for passing SGPR indices into `computeMasks`: the current code emits into +`sTmp+1..4` allocated by the surrounding `defineAndResources`. To guarantee **zero asm +diff**, the component must not re-allocate; it receives the same operands and emits the same +instructions in the same order. This is a mechanical lift, not a rewrite. + +### 2.3 Behavior-preserving refactor wiring + +- `KernelWriter.py:9163-9176` → `ClusterLoad.find(self).declareSgprs(self, kernel)`. +- `KernelWriter.py:2838-2848` → `... .undeclareSgprs(self, kernel)`. +- `KernelWriterAssembly.py:2646-2694` → `... .computeMasks(self, kernel, sgprWgX=..., sgprWgY=..., sgprNWgX=..., sTmp=...)` (same operands as today). +- `KernelWriterAssembly.py:18901-18902` (dense) → `... .applyToDescriptor(self, kernel, descSgprName(1), tc)`. +- `KernelWriterAssembly.py:19059-19060` (wave-separated) → same call (component resolves to combined name via `usesCombinedMask`). +- `SubtileGREmit.py:1108-1113` (subtile) → `... .applyToDescriptor(writer, kernel, descSgprName(1), tc, subtile=True)`. + +Guard: `ClusterLoad.find` returns `None` when `HasTDM`/`TDMInst` don't match; callers keep a +`comp = ClusterLoad.find(self)` + `if comp:` fallback exactly like `TensorDataMoverLoad.find` +usage, so non-TDM archs are untouched. + +**Regression gate (must stay green, zero asm change):** +`test_r4_xccremap_char.py` (drives `ClusterDim=[2,2]` with both a `MIWaveGroup` prod==1 → +split `MulticastMaskA/B` kernel and a prod==4 → combined `MulticastMask` kernel) and +`test_streamk_cluster_coop_load_gfx1250_char.py`. Add a byte-exact golden snapshot of the +emitted mask/apply region for a `[2,2]` combined config and a subtile split config before +the refactor; diff to zero after. + +--- + +## 3. Decoupling `Multicast` from `ClusterDim` + +### 3.1 Mechanism + +Add an explicit tri-state solution parameter (`ValidParameters.py`, near `ClusterDim`): + +```python +# -1 = auto (legacy: ClusterDim!=[1,1] implies Multicast, except StreamK cluster paths) +# 0 = force multicast off +# 1 = force multicast on (independent of ClusterDim coupling) +"Multicast": [-1, 0, 1], +``` + +Default `-1` preserves **all** existing emitted asm. Rewrite `Solution.py:1046-1064`: + +```python +state["ClusterBarrier"] = False +mc = state.get("Multicast", -1) +if mc == 1: + state["Multicast"] = True +elif mc == 0: + state["Multicast"] = False +elif state.get("StreamKMulticast", 0): + state["Multicast"] = True # derived cooperative-load path +else: # auto (legacy behavior) + state["Multicast"] = (state["ClusterDim"] != [1, 1] + and state["StreamK"] == 0) +# ClusterBarrier decision (legacy/cluster-subtile path + the derived StreamKMulticast path) +if state["ClusterDim"] != [1, 1] and state["Multicast"] and state["TDMInst"] != 0 \ + and (state["StreamK"] == 0 or state.get("StreamKMulticast", 0)) \ + and isaInfoMap[state["ISA"]].asmCaps.get("HasClusterBarrier", False): + state["ClusterBarrier"] = True +``` + +Net effect for existing configs: `Multicast` default `-1` → identical derivation to today +(subtile/dense clustered configs keep `Multicast=True`; non-multicast Stream-K keeps it +`False`). New capability: `Multicast=1`/`0` explicit override, and `StreamKMulticast` (§4) +turns on multicast for its own path without relying on the coupling. + +This is the one step in the sequence that *can* legitimately change asm — but only for +configs that opt in. Keep existing subtile/multicast YAMLs on `Multicast=-1` so their +snapshots don't move. + +--- + +## 4. StreamK cooperative cluster loads (DP region) + +### 4.1 The spatial vs reduction tension, resolved + +- Cooperative loads need a **spatial** cluster over **distinct** tiles (`[C,1]` consecutive + WGs → M-adjacent → shared B). +- The shipped reduction needs a **`[C,1]` same-tile K-peer** cluster (host ties + `skGrid = C*tiles`, `skIndexToWG` collapses the C peers onto one tile). +- A WG's HW cluster membership is fixed at launch, so a single kernel cannot have the + cluster mean "spatial DP" and "K-peers" at once. + +**Resolution: mutual exclusion.** For v1, `StreamKMulticast` and the barrier-only cluster +reduction are mutually exclusive. When `StreamKMulticast` is on, `ClusterDim=[C,1]` is free +to denote the spatial cluster; the reduction is off, so no `[C,1]` is claimed for K-peers. **No +separate `SpatialClusterDim` is needed in v1** — the same `ClusterDim=[C,1]` is reused, its +meaning selected by which StreamK cluster mode is enabled. (A future 2-D `[C0,C1]` mode for +A-multicast or combined reduction+coop-load would introduce a distinct spatial dim; deferred.) + +### 4.2 MVP target: DP-only, B-multicast + +The multicast mask is computed **once** at kernel init and baked into the descriptor +`Group1[word0]`; it is applied to every mainloop TDM load. For correctness, the sharing +relationship it encodes must hold for **every** load the WG issues. Therefore v1 requires +that the WG never enters a partial (K-split) tile with a stale spatial mask: + +**v1 MVP = cooperative B-multicast for the DP region with `StreamKForceDPOnly=1`** (no SK +partial tiles → no reduction → the static spatial mask is valid for all issued loads), and +single DP round per cluster (see §4.4). + +Recommended smallest correct MVP: +- **`StreamKMulticast=1`** requires `StreamKForceDPOnly=1` and no barrier-only cluster reduction. +- Reuse `ClusterDim=[C,1]`, `C` a power of two in `2..16`. +- Multicast B across the cluster; A loaded per-WG. + +Deferred (see §7): mask-toggle at the DP→SK boundary to allow cooperative DP loads inside a +full SK3 kernel (SK tail reduces via the existing global-flag path); combined +reduction+coop-load; A-multicast / 2-D spatial cluster; multi-round with per-round validity. + +### 4.3 Mask derivation from post-`skIndexToWG` coords + +The existing `computeMasks` derives `wg_x`/`wg_y` from the raw HW cluster registers +(`ttmp*`). For StreamK the authoritative WG coordinates are the ones produced by +`skIndexToWG` (`StreamK.py:503`), not the raw HW position, because StreamK re-derives +`(WG0,WG1)` from `StreamKIdx`. For the DP `[C,1]` MVP the two coincide +(`wg_x = StreamKIdx & (C-1)`, since the reduction ttmp9 workaround is skipped and +`WorkGroup0 = cluster_x*C + wg_x`), so the MVP feeds `computeMasks` with +`sgprWgX = StreamKIdx & (C-1)`, `wg_y = 0`, `C0=C`, `C1=1`, yielding the B-broadcast mask +`(1< **Status (as shipped):** `StreamKMulticast` is **not** a public/benchmark parameter. +> It is a derived-only internal state key (like `ClusterBarrier`), enabled automatically +> when `ClusterDim != [1,1]` on SK3 — see §3.1 / the +> collapse in `assignProblemIndependentDerivedParameters`. It has no `ValidParameters.py` +> entry; `_validateStreamKMulticast` is the internal guard that hard-rejects a derived +> config it cannot satisfy. The `[0,1]` "param" framing below reflects the original plan. + +### 5.1 Parameters + +| Param | Values | Meaning | +|---|---|---| +| `Multicast` | `[-1,0,1]` | Decoupled multicast control (§3). Default `-1` = legacy auto. | + +`StreamKMulticast` is internal-derived (see status note above), not a user-settable param. +The separate barrier-only cluster reduction feature is orthogonal and unchanged. + +### 5.2 Validation rules (`Solution.py`, reject-with-reason pattern) + +The derived `StreamKMulticast` on-state is rejected (`_validateStreamKMulticast`) unless all hold: +- `StreamK == 3` (SK3; DP schedule + `skIndexToWG` assumptions). +- `StreamKForceDPOnly == 1` (v1: no partial tiles reach the static mask). +- not the barrier-only cluster reduction (mutually exclusive — cannot claim `[C,1]` two + ways; the collapse gives the reduction precedence, and the reject itself is added + alongside that feature). +- `ClusterDim == [C,1]`, `C` power of two, `2 <= C <= 16` (`ClusterDim[1] == 1`; StreamK + grid is effectively 1-D along x). +- gfx1250 `asmCaps["HasTDM"]` and `TDMInst != 0` (MX TDM loads; multicast is a TDM feature). +- `StreamKXCCMapping == 0` (XCC=3 overflows SGPRs; WGM/XCC remap is bypassed under + clustering anyway). + +Also set `Multicast=True` for this path via the decoupled derivation (§3.1) so +`ClusterLoad.applyToDescriptor` fires. + +### 5.3 Interaction notes + +- gfx1250 StreamK already uses TDM (`MXLoadInst=TDM` → `TDMInst=3`), so `HasTDM` is the + natural cap and no new transport is introduced. +- `StreamKAtomic` is incompatible (atomic path skips the workspace/tile structure the DP + schedule relies on); reject. + +--- + +## 6. Ordered implementation task breakdown + +Sequencing keeps the tree green: **pure extraction first (zero asm), then the decouple, +then the feature.** `[P]` marks tasks that can proceed in parallel within a group. + +### Group 0 — Snapshot baseline (do first, no source change) +- **T0.1** Add byte-exact golden snapshots for the mask compute+apply region: a `[2,2]` + combined-mask config, a `[2,2]` split-mask config, and a subtile split config. Anchor: + extend `test_r4_xccremap_char.py` and the subtile char config under + `Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/`. + +### Group 1 — `ClusterLoad` extraction (pure refactor, snapshot-gated) +- **T1.1** Declare `class ClusterLoad(Component)` at `Component.py:305-313`; add + `"ClusterLoad"` to `Components/__init__.py:__all__`. +- **T1.2** Create `Components/ClusterLoad.py` with `ClusterLoadTDM` (SPDX header; + `asmCaps={"HasTDM":True}`, `kernel={"TDMInst":3}`); implement `usesCombinedMask`, + `maskSgprName`, `cooperativeThreadPartition`. +- **T1.3** Move mask compute into `computeMasks`; wire `KernelWriterAssembly.py:2646-2694` + to call it with the existing SGPR operands. +- **T1.4 [P]** Move declare/undeclare into `declareSgprs`/`undeclareSgprs`; wire + `KernelWriter.py:9163-9176` and `:2838-2848`. +- **T1.5 [P]** Implement `applyToDescriptor`; wire the 3 apply sites + (`KernelWriterAssembly.py:18901-18902`, `:19059-19060`, `SubtileGREmit.py:1108-1113`). +- **T1.6** Point `GL2Prefetch.py:26,78` at `cooperativeThreadPartition` (optional but DRY). +- **Gate:** T0.1 snapshots + `test_r4_xccremap_char.py` + + `test_streamk_cluster_coop_load_gfx1250_char.py` all diff to zero. + +### Group 2 — Decouple `Multicast` (may change asm only for opt-in configs) +- **T2.1** Add `"Multicast": [-1,0,1]` to `ValidParameters.py`. +- **T2.2** Rewrite `Solution.py:1046-1064` per §3.1; keep existing YAMLs on `-1`. +- **Gate:** existing multicast/subtile/reduction snapshots unchanged; add a snapshot for an + explicit `Multicast=1`/`0` override. + +### Group 3 — StreamK cooperative DP multicast (feature) +- **T3.1** Derive `StreamKMulticast` internally (no `ValidParameters.py` entry — it is a + derived-only state key, auto-enabled on SK3 + `ClusterDim!=[1,1]` not reduction) + + `Solution.py` validation (§5.2) + decoupled `Multicast` enablement. +- **T3.2 [P]** Host: `sizeMapping.streamKMulticast` (`ContractionSolution.hpp:146`, + `Contractions.py:630,725`); `getSKGridImpl` grid = `ceil(tiles/C)*C` + (`ContractionSolution.cpp:4087-4090`/`:3214-3218` region); thread through launch (no + launch change — `:1783-1794` already cluster-aware). +- **T3.3** Kernel: at StreamK preLoop compute `wg_x = StreamKIdx & (C-1)` and the + `clusterMulticastValid` predicate (§4.4) from `NumWorkGroups0`/`totalTiles`; feed + `ClusterLoad.computeMasks` with `wg_y=0, C0=C, C1=1` and gate the mask OR on the predicate + (via `applyToDescriptor`). Anchors: `StreamK.py:2645-2663` (StreamKIdx/ttmp9), + `skIndexToWG` `:487-503`. +- **T3.4** Ensure A stays per-WG (maskA = self) and B multicast (maskB = all-C); verify the + `[C,1]` path selects split `MulticastMaskA/B` names. + +Dependencies: Group 1 → Group 2 → Group 3 (T3.1). T3.2 (host) is parallel to T3.3/T3.4 +(kernel) once T3.1 lands the param. + +--- + +## 7. Correctness, edge cases, risks + +| Risk | Why | De-risk | +|---|---|---| +| **Wrong-B multicast → silent wrong results** | Static mask applied to a load where the C WGs don't actually share B (SK partial tile, non-adjacent, partial cluster). | v1 requires `StreamKForceDPOnly` (no partial tiles) + single round + init-time `clusterMulticastValid` predicate that ORs the mask only when provably valid; else load normally. | +| **Idle WG in a partial last cluster** | `totalTiles % C != 0` leaves tail WGs with no tile; a multicast expecting them may hang/mis-deliver. | Predicate disables multicast for a not-fully-populated cluster. **OPEN Q**: exact HW behavior when a masked target WG is not at a matching TDM load (see §8). | +| **DP→SK boundary with a live spatial mask** | In a full SK3 kernel the mask would still be applied to SK-tile loads. | v1 forbids the SK region (`StreamKForceDPOnly`). Deferred: clear the mask bit (single `SAndB32` on `Group1[word0]`) at the boundary (`StreamK.py:2907-2932`) since DP is first and SK last. | +| **Refactor changes asm** | Component lift could reorder/realloc SGPRs. | `computeMasks` takes the exact existing SGPR operands; snapshot gate (Group 0) diffs to zero. | +| **`Multicast` decouple regresses existing configs** | Changing the coupling could flip `Multicast` for subtile/dense clustered kernels. | Default `-1` = identity derivation; existing YAMLs untouched; snapshot-gated. | +| **Mask from raw HW coords instead of skIndexToWG** | StreamK re-derives `(WG0,WG1)`. | For DP `[C,1]` the two coincide (`wg_x=StreamKIdx&(C-1)`); documented general rule for 2-D follow-ups. | +| **SGPR pressure** | gfx1250 StreamK SGPRs tight; XCC=3 overflows. | Reuse existing `MulticastMask*` SGPRs; predicate uses `allocTmpSgpr` scopes and bitwise ops (C power of two); forbid `StreamKXCCMapping=3`. | +| **A-multicast expectation** | Consecutive `[C,1]` cluster shares B, not A. | Documented; A stays per-WG in v1; A-multicast needs an N-strided/2-D cluster (deferred). | + +--- + +## 8. Test plan (behavior → test) — as shipped + +Mirrors existing patterns: CPU asm-string unit, syrupy snapshot characterization, GPU +roundtrip (run under the arch's simulator/hardware), and C++ client end-to-end. + +| Behavior | Test type | Where | +|---|---|---| +| **Refactor is byte-exact**: `computeMasks`/`applyToDescriptor`/declare/undeclare emit identical asm for `[2,2]` combined, `[2,2]` split, and subtile split configs | snapshot char (regression gate) | `test_r4_xccremap_char.py`, `test_streamk_cluster_coop_load_gfx1250_char.py` | +| `ClusterLoad.find` returns TDM impl on gfx1250, `None` (fallback) on non-TDM; `usesCombinedMask`/`maskSgprName` matrix (combined vs split vs metadata) | CPU unit | `Tests/unit/test_cluster_load_component.py` | +| `Multicast` tri-state: `-1` reproduces legacy derivation exactly; `0` forces off; `1` forces on independent of `ClusterDim` (derived value is an int on all paths) | CPU unit | `Tests/unit/test_multicast_tristate.py`, `Tests/unit/test_multicast_legacy_coercion.py` | +| `StreamKMulticast` validation matrix: accepted only for SK3 + `ClusterDim=[C,1]` (pow2 2..16) + gfx1250 `HasTDM`/`TDMInst=3` + XCC=0 + not `StreamKAtomic`; auto-enabled config with `Multicast=0` rejected; rejected otherwise | CPU unit | `Tests/unit/test_streamk_multicast.py` | +| Cluster config emits real gfx1250 asm (`err==0`); DP loads carry the B-broadcast mask (`MulticastMaskB` OR into `Group1`); A load carries self-only mask; predicate gate + DP→SK boundary clear present | snapshot char | `test_streamk_cluster_multicast_gfx1250_char.py` (designed config `_designed/gfx1250/streamk_cluster_multicast.yaml`, `__snapshots__/*.ambr`) | +| DP GEMM with `StreamKMulticast`, `C∈{2,4}` reduces correctly and does not hang; matches non-multicast reference | GPU roundtrip | `Tests/unit/test_streamk_cluster_reduction_gpu.py` and cluster char roundtrips (`@requires_gpu_gfx1250`; watchdog on hang) | +| Real multi-WG cluster StreamK DP GEMM end-to-end | C++ client | `Tests/common/streamk/gfx1250/core/sk_mxf8gemm_cluster_multicast.yaml` (+ mxf4 sibling) | + +The gfx1250 GPU marker reuses `Tests/unit/gpu_test_helpers.py` +(`HAS_GFX1250`/`requires_gpu_gfx1250`, `TENSILE_GPU_TARGET=gfx1250`). + +--- + +## 9. Open questions (need a human decision at the design gate) + +1. **HW multicast semantics when a masked target WG is not executing a matching TDM load** + (idle tail WG, or a WG that finished its tiles): does the loader hang, drop the + broadcast, or corrupt? This determines whether the init-time predicate is sufficient or + whether the grid must be constrained to always-full clusters. (Mirrors the reduction + feature's deadlock concern.) +2. **Is `StreamKForceDPOnly`-scoped v1 valuable enough**, or should we prioritize the + DP→SK boundary mask-clear so cooperative loads apply inside a normal SK3 kernel (the + larger real-world win)? The boundary-clear is a modest add (§7) but widens the test + matrix. +3. **`Multicast` as tri-state `[-1,0,1]` vs a separate boolean** (`MulticastEnable`) — the + tri-state keeps one knob and a clean legacy default; a separate boolean is more explicit + but adds a second interacting param. Recommend tri-state; confirm. +4. **Grid rounding policy** for `StreamKMulticast`: `ceil(tiles/C)*C` (idle tail WGs) vs + requiring `totalTiles % C == 0` at selection time (fewer sizes, no idle WGs). Depends on + Q1. + +--- + +## 10. Summary of decisions + +- **`ClusterLoad` component**: new capability-selected category (`asmCaps HasTDM`, + `kernel TDMInst=3`) owning mask value (`computeMasks`), declare/undeclare, topology + decision (`usesCombinedMask`/`maskSgprName`), descriptor attach (`applyToDescriptor`, + gate folded in), and `cooperativeThreadPartition`. It does **not** own descriptor groups + or LDS offsets. The three apply sites and the subtile clone unify onto it; refactor is + byte-exact and snapshot-gated. +- **Decouple**: `Multicast` becomes a tri-state param (`-1` auto = legacy identity, `0` + off, `1` on); `Solution.py:1046-1064` derivation rewritten so `ClusterDim!=[1,1]` no + longer unconditionally forces multicast. +- **StreamK target = DP region, B-multicast** on a `[C,1]` consecutive-WG cluster + (M-adjacent tiles share the N-block). Spatial-vs-reduction tension resolved by **mutual + exclusion** (`StreamKMulticast` xor the barrier-only cluster reduction), reusing + `ClusterDim=[C,1]` with mode-selected meaning; no separate spatial dim in v1. +- **MVP** = `StreamKMulticast=1` requiring `StreamKForceDPOnly=1`, single DP round, init-time + `clusterMulticastValid` predicate gating the mask. Deferred: DP→SK boundary mask-clear, + multi-round, A-multicast / 2-D cluster, combined reduction+coop-load. +- **Enablement**: `Multicast` (`[-1,0,1]`) is the one public knob; `StreamKMulticast` is + internal-derived (auto-enabled on SK3 + `ClusterDim!=[1,1]`, not reduction), not a param. +- **Sequencing**: extraction (green) → decouple → feature. diff --git a/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py b/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py index e701a5695a6a..1c045d9c8ecf 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py +++ b/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py @@ -571,6 +571,9 @@ {"StreamKAtomic": [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]}, @@ -612,6 +615,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]} ] diff --git a/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py b/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py index 123dc69a803c..2b953f96f5f1 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py +++ b/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py @@ -837,6 +837,16 @@ def makeValidMatrixInstructions(): # 0: use linear reduction # 1: use tree reduction "StreamKFixupTreeReduction": [0, 1], + # NOTE: StreamKMulticast (the gfx1250 StreamK DP cooperative cluster-load / + # TDM B-multicast fast path) is intentionally NOT a valid/benchmark + # parameter. It is a DERIVED-ONLY internal state key (see the ClusterBarrier + # precedent): Solution.assignProblemIndependentDerivedParameters auto-enables + # it for StreamK==3 + ClusterDim != [1,1] (the "bare StreamK cluster" + # collapse), and _validateStreamKMulticast then hard-rejects any cluster + # config that cannot satisfy its constraints. It must not be user/YAML- + # settable, so it has no entry here (checkParametersAreValid would otherwise + # accept it as a fork/constant param). It still serializes to C++ via + # Contractions.SizeMapping (streamKMulticast, read with d.get()). # 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 @@ -1117,6 +1127,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 diff --git a/projects/hipblaslt/tensilelite/Tensile/Component.py b/projects/hipblaslt/tensilelite/Tensile/Component.py index 9639eced4bef..46a27b893e88 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Component.py +++ b/projects/hipblaslt/tensilelite/Tensile/Component.py @@ -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. diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/ClusterLoad.py b/projects/hipblaslt/tensilelite/Tensile/Components/ClusterLoad.py new file mode 100644 index 000000000000..6b4548554a59 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Components/ClusterLoad.py @@ -0,0 +1,170 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +"""Cluster (multicast) TDM load component. + +Centralizes the multicast ("cluster load") mask machinery (value compute, +``MulticastMask*`` SGPR declare/undeclare, combined-vs-split topology decision, +and per-load-site descriptor attach) that was previously duplicated across +``KernelWriter``/``KernelWriterAssembly``/``SubtileGREmit``. Behavior-preserving: +every method emits byte-identical assembly, receiving the SGPR operands the +caller already holds rather than re-allocating. Capability-selected +(``HasTDM`` + ``TDMInst == 3``), like ``TensorDataMoverLoad``. +""" + +from ..Component import ClusterLoad +from ..Common import clusterEnabled +from typing import Mapping +from rocisa.code import Module, Label +from rocisa.container import sgpr +from rocisa.instruction import SLShiftLeftB32, SMulI32, SBitcmp1B32, SCBranchSCC1, SBranch + + +class ClusterLoadTDM(ClusterLoad): + asmCaps = {"HasTDM": True} + kernel = {"TDMInst": 3} + + def __call__(self, writer: "KernelWriterAssembly", kernel: Mapping): + # Abstract-satisfying no-op, mirrors TensorDataMoverLoad.__call__. + pass + + # -- topology decision --------------------------------------------------- + + def usesCombinedMask(self, kernel: Mapping) -> bool: + """True when the single-parity combined ``MulticastMask`` applies. + + Single source of truth for the combined-vs-split decision. Subtile and + StreamK DP multicast both need the split A/B masks: subtile issues A and + B on every wave (no wave-parity split), and StreamK broadcasts only B + across the [C,1] cluster while A stays per-workgroup -- the combined + parity mask would be wrong for both. + """ + if kernel.get("StreamKMulticast", 0): + return False + tdmA: bool = kernel["enableTDMA"] + tdmB: bool = kernel["enableTDMB"] + return tdmA and tdmB and kernel["NumWaves"] > 1 and not kernel.get("UseSubtileImpl") + + def maskSgprName(self, kernel: Mapping, tc: str, *, subtile: bool = False, + waveSeparated: bool = False) -> str: + """Resolve the multicast-mask SGPR name. + + Wave-separated (non-subtile) uses the combined ``"MulticastMask"``; + dense/subtile and StreamK multicast use the split ``f"MulticastMask{tc}"`` + (any ``MXS`` prefix stripped). StreamK forces the split name so B never + resolves to the never-declared combined SGPR. + """ + if kernel.get("StreamKMulticast", 0): + string = tc.removeprefix("MXS") if tc.startswith("MXS") else tc + return f"MulticastMask{string}" + if waveSeparated and not subtile: + return "MulticastMask" + string = tc.removeprefix("MXS") if tc.startswith("MXS") else tc + return f"MulticastMask{string}" + + # -- SGPR declare / undeclare ------------------------------------------- + + def declareSgprs(self, writer: "KernelWriter", kernel: Mapping) -> None: + """Allocate the ``MulticastMask*`` SGPRs (lift of KernelWriter).""" + if not kernel["Multicast"]: + return + tdmM: bool = kernel["enableTDMMetadata"] + if self.usesCombinedMask(kernel): + writer.defineSgpr("MulticastMask", 1) + else: + writer.defineSgpr("MulticastMaskA", 1) + writer.defineSgpr("MulticastMaskB", 1) + if tdmM: + writer.defineSgpr("MulticastMaskMetadata", 1) + + def undeclareSgprs(self, writer: "KernelWriter", kernel: Mapping) -> Module: + """Free the ``MulticastMask*`` SGPRs (lift of KernelWriter).""" + mod = Module() + if not (kernel["Multicast"] and kernel["TDMInst"] != 0): + return mod + tdmM: bool = kernel["enableTDMMetadata"] + if self.usesCombinedMask(kernel): + mod.add(writer.undefineSgpr("MulticastMask")) + else: + mod.add(writer.undefineSgpr("MulticastMaskA")) + mod.add(writer.undefineSgpr("MulticastMaskB")) + if tdmM: + mod.add(writer.undefineSgpr("MulticastMaskMetadata")) + return mod + + # -- mask value computation --------------------------------------------- + + def computeMasks(self, writer: "KernelWriterAssembly", kernel: Mapping, *, + sgprWgX: int, sgprWgY: int, sgprNWgX: int, sTmp: int) -> Module: + """Compute the multicast mask value(s) into the ``MulticastMask*`` SGPRs. + + Verbatim lift of the ``defineAndResources`` mask compute; the caller + passes the operands it already holds (``sgprWgX``/``sgprWgY``/``sgprNWgX`` + and ``sTmp`` whose ``+4`` slot is scratch) so the output is byte-identical. + """ + mod = Module() + if not kernel["Multicast"]: + return mod + mod.addComment0("Calculate multicast mask") + + maskA = 1 + for idx in range(kernel["ClusterDim"][1]): + maskA |= (1 << (idx * kernel["ClusterDim"][0])) + + maskB = (1 << kernel["ClusterDim"][0]) - 1 + + if kernel["enableTDMMetadata"]: + if kernel["ProblemType"]["Sparse"] == 1: + mod.add(SLShiftLeftB32(dst=sgpr("MulticastMaskMetadata"), shiftHex=sgpr(sgprWgX), src=hex(maskA),\ + comment="Setting metadata mask (follows sparse A)")) + elif kernel["ProblemType"]["Sparse"] == 2: + mod.add(SMulI32(dst=sgpr(sTmp+4), src0=sgpr(sgprWgY), src1=sgpr(sgprNWgX),\ + comment="Shift factor: wg_y * nwg_x (metadata)")) + mod.add(SLShiftLeftB32(dst=sgpr("MulticastMaskMetadata"), shiftHex=sgpr(sTmp+4), src=hex(maskB),\ + comment="Setting metadata mask (follows sparse B)")) + + if self.usesCombinedMask(kernel): + setMulticastMaskLblOdd = Label(f"setMulticastMask_OddWave", "") + setMulticastMaskLblEven = Label(f"setMulticastMask_EvenWave", "") + setMulticastMaskLblEnd = Label(f"setMulticastMaskEnd", "") + + mod.add(SBitcmp1B32(sgpr("WaveIdx"), 0, "Check parity of wId")) + mod.add(SCBranchSCC1(setMulticastMaskLblOdd.getLabelName(), "Jump if wId is odd")) + + mod.add(setMulticastMaskLblEven) + mod.add(SLShiftLeftB32(dst=sgpr("MulticastMask"), shiftHex=sgpr(sgprWgX), src=hex(maskA),\ + comment="Setting maskA for even wave")) + mod.add(SBranch(setMulticastMaskLblEnd.getLabelName())) + mod.add(setMulticastMaskLblOdd) + mod.add(SMulI32(dst=sgpr(sgprWgY), src0=sgpr(sgprWgY), src1=sgpr(sgprNWgX),\ + comment="Shift factor: wg_y * nwg_x")) + mod.add(SLShiftLeftB32(dst=sgpr("MulticastMask"), shiftHex=sgpr(sgprWgY), src=hex(maskB),\ + comment="Setting maskB for odd wave")) + mod.add(setMulticastMaskLblEnd) + + else: + mod.add(SLShiftLeftB32(dst=sgpr("MulticastMaskA"), shiftHex=sgpr(sgprWgX), src=hex(maskA),\ + comment="Setting maskA")) + + mod.add(SMulI32(dst=sgpr(sgprWgY), src0=sgpr(sgprWgY), src1=sgpr(sgprNWgX),\ + comment="Shift factor: wg_y * nwg_x")) + mod.add(SLShiftLeftB32(dst=sgpr("MulticastMaskB"), shiftHex=sgpr(sgprWgY), src=hex(maskB),\ + comment="Setting maskB")) + return mod + + # -- descriptor attach --------------------------------------------------- + + def applyToDescriptor(self, writer: "KernelWriterAssembly", kernel: Mapping, + group1: int | str, tc: str, *, subtile: bool = False, + waveSeparated: bool = False) -> Module: + """OR the multicast mask into descriptor ``Group1[word0]``. + + Folds the ``Multicast and enableCluster`` gate, mask-name choice, and the + ``SOrB32`` attach; returns an empty ``Module`` when the gate is not met. + """ + from .TensorDataMover import TensorDataMoverLoad + mod = Module() + if kernel["Multicast"] and clusterEnabled(kernel["ClusterDim"]): + mask = self.maskSgprName(kernel, tc, subtile=subtile, waveSeparated=waveSeparated) + tdm = TensorDataMoverLoad.find(writer) + mod.add(tdm.setMulticastMask(group1, mask, writer)) + return mod diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py b/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py index 7b5c93d7a346..69b62254517b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py @@ -27,7 +27,7 @@ from rocisa.instruction import GlobalInv, GlobalWb, SAddCU32, SAddU32, SAndB32, SBarrier, \ SBranch, SCBranchSCC0, SCBranchSCC1, SCMovB32, SCSelectB32, SCmpEQU32, SCmpEQU64, \ SCmpGtU32, SCmpLeU32, SCmpLtU32, SLShiftLeftB32, SLShiftLeftB64, SLShiftRightB32, VLShiftLeftB32, SLoadB32, \ - SMaxI32, SMinU32, SMovB32, SMovB64, SMulI32, SNop, SOrB32, SSleep, SStoreB32, SSubU32, \ + SMaxI32, SMinU32, SMovB32, SMovB64, SMulHIU32, SMulI32, SNop, SOrB32, SSleep, SStoreB32, SSubU32, \ SWaitCnt, SWaitXCnt, VAddF32, VAddF64, VAddPKF16, VAddU32, VSubU32, VLShiftRightB32, VMovB32, \ VReadfirstlaneB32, VCmpXEqU32, VCvtBF16toFP32, GlobalAtomicIncU32Saddr, BufferLoadB32, BufferStoreB32, \ SAtomicInc, DSLoadB32, DSStoreB32, SLongBranch, SLongBranchPositive @@ -1121,9 +1121,31 @@ def computeWorkspaceSrd(self, writer, kernel, sPartialIdx, tmpSgpr = None): assert kernel["BufferStore"] module.addSpaceLine() - module.add(SMulI32(dst=sgpr(tmpSgpr), src0=hex(kernel["MacroTile0"]*kernel["MacroTile1"]*writer.states.bpeCinternal), src1=sPartialIdx, comment="Offset to correct partials tile")) + # 64-bit slot byte offset. The per-tile workspace stride + # MacroTile0*MacroTile1*bpe times the StreamK partial index can exceed + # 2^32 for large SK grids (grows with C: C>=4/C>=8 on big tiles), so a + # 32-bit SMulI32 product silently wraps and the peer write / owner read + # SRD then aliases the wrong workspace slot. Compute the high word with + # SMulHIU32 and fold it (plus the lo-add carry) into SrdWS+1 instead of + # adding only the carry. + # + # Applied universally to every StreamK path that addresses the partials + # workspace (multicast [C,1], cluster reduction, factored, and + # non-cluster StreamK). The overflow depends only on the per-slot tile + # stride and the StreamK slot count -- the partials workspace is + # partialTileSize == tileSize * skGrid regardless of cluster mode -- so + # the multicast path can overflow on large problems just like the + # reduction/factored paths. Emitting the 64-bit offset everywhere costs + # one extra SGPR and one s_mul_hi_u32; it intentionally changes the + # multicast codegen (byte-identity with the prior 32-bit path is + # deliberately dropped now that the fix is universal). + offBytes = hex(kernel["MacroTile0"]*kernel["MacroTile1"]*writer.states.bpeCinternal) + tmpHi = writer.sgprPool.checkOut(1, "SKSlotOffsetHi") + module.add(SMulI32(dst=sgpr(tmpSgpr), src0=offBytes, src1=sPartialIdx, comment="Offset to correct partials tile (low word)")) + module.add(SMulHIU32(dst=sgpr(tmpHi), src0=offBytes, src1=sPartialIdx, comment="partials tile offset (high word) for 64-bit SRD")) module.add(SAddU32(dst=sgpr("SrdWS+0"), src0=sgpr("SrdWS+0"), src1=sgpr(tmpSgpr), comment="add lo to SRD")) - module.add(SAddCU32(dst=sgpr("SrdWS+1"), src0=sgpr("SrdWS+1"), src1=0, comment="add hi to SRD")) + module.add(SAddCU32(dst=sgpr("SrdWS+1"), src0=sgpr("SrdWS+1"), src1=sgpr(tmpHi), comment="add hi (offset high word + lo carry) to SRD")) + writer.sgprPool.checkIn(tmpHi) if tmpLocal is not None: writer.sgprPool.checkIn(tmpLocal) @@ -2484,6 +2506,182 @@ class StreamKTwoTileDPFirst(StreamK): requiresWorkspaceReductionStorePath = True supportsSubtileImpl = True + def streamKMulticastMaskPredicate(self, writer, kernel): + """Gate the DP B-multicast on the runtime clusterMulticastValid predicate. + + The B-broadcast mask value (MulticastMaskB = (1< share the same B), and + 2. the cluster is fully populated: clusterBase + C <= totalTiles, where + clusterBase = StreamKIdx & ~(C-1) and totalTiles = nWG0*nWG1 (so no + peer is an idle tail workgroup with no matching load). + + When either fails we rewrite MulticastMaskB to the self-only MulticastMaskA + so B is loaded per-workgroup (a normal load) -- correct for all sizes and, + critically, never leaving a masked target without a matching load (the + conservative full-cluster path chosen because the masked-idle-target HW + semantics could not be verified on silicon). Inert unless StreamKMulticast. + """ + module = Module("StreamK multicast mask predicate") + if not kernel.get("StreamKMulticast", 0): + return module + c = kernel["ClusterDim"][0] + module.addComment0("StreamKMulticast: gate B-broadcast on clusterMulticastValid") + skConstsInVgprs = writer.isStreamKConstantsToVgprEnabled(kernel) + mcInvalid = Label(writer.labels.getNameInc("SKMC_Invalid"), "") + mcEnd = Label(writer.labels.getNameInc("SKMC_End"), "") + with writer.allocTmpSgpr(3, tag="SKMulticastPredicate") as tRes: + t0 = tRes.idx + t1 = tRes.idx + 1 + t2 = tRes.idx + 2 + # cond1: nWG0 % C == 0 (C is a power of two) + module.add(SAndB32(dst=sgpr(t0), src0=sgpr("NumWorkGroups0"), src1=hex(c - 1), + comment="nWG0 %% C (C power of two)")) + module.add(SCmpEQU32(src0=sgpr(t0), src1=0, comment="nWG0 aligned to C?")) + module.add(SCBranchSCC0(labelName=mcInvalid.getLabelName(), + comment="unaligned M -> B not shared, load normally")) + # cond2: clusterBase + C <= totalTiles + sIdx = writer.acquireStreamKConstSgpr(kernel, "StreamKIdx") + if skConstsInVgprs: + module.add(VReadfirstlaneB32(dst=sgpr(sIdx), src=vgpr(writer.states.skConstVgprs["StreamKIdx"]))) + module.add(SAndB32(dst=sgpr(t2), src0=sgpr(sIdx), src1=hex((~(c - 1)) & 0xFFFFFFFF), + comment="clusterBase = StreamKIdx & ~(C-1)")) + writer.releaseStreamKConstSgpr(sIdx) + module.add(SAddU32(dst=sgpr(t2), src0=sgpr(t2), src1=hex(c), comment="clusterBase + C")) + module.add(SMulI32(dst=sgpr(t1), src0=sgpr("NumWorkGroups0"), src1=sgpr("NumWorkGroups1"), + comment="totalTiles = nWG0 * nWG1")) + module.add(SCmpLeU32(src0=sgpr(t2), src1=sgpr(t1), + comment="cluster fully populated? clusterBase+C <= totalTiles")) + module.add(SCBranchSCC1(labelName=mcEnd.getLabelName(), + comment="valid cluster -> keep B broadcast mask")) + module.add(mcInvalid) + module.add(SMovB32(dst=sgpr("MulticastMaskB"), src=sgpr("MulticastMaskA"), + comment="invalid cluster -> B loaded normally (self-only mask)")) + module.add(mcEnd) + return module + + def streamKMulticastBoundaryClear(self, writer, kernel): + """Clear the B-broadcast mask at the DP->SK boundary. + + The B multicast mask lives in the persistent ``MulticastMaskB`` SGPR and + is OR'd fresh into the B TDM descriptor by ``initTDMDescriptor`` + (``ClusterLoad.applyToDescriptor``) on *every* persistent-loop iteration. + Manipulating the descriptor word directly would be undone by the next + rebuild, so the correct lever is the mask SGPR itself. + + When a workgroup transitions out of the single DP round into SK + partial-tile work, its cluster peers no longer co-issue the identical + full-K B load, so the spatial broadcast is no longer valid. We rewrite + ``MulticastMaskB`` to the self-only ``MulticastMaskA`` bit; from here on + every descriptor rebuild ORs only the self bit -> normal per-workgroup B + load for the remaining SK iterations. Because DP is a single round first + and SK is last (no return to DP), this one-shot rewrite is sufficient. + Inert unless StreamKMulticast. + """ + module = Module("StreamK multicast DP->SK boundary clear") + if not kernel.get("StreamKMulticast", 0): + return module + module.addComment0("StreamKMulticast: clear B-broadcast mask at DP->SK boundary") + module.add(SMovB32(dst=sgpr("MulticastMaskB"), src=sgpr("MulticastMaskA"), + comment="DP->SK: drop B broadcast -> self-only (normal B load)")) + return module + + def _clusterElectArriveSignal(self, writer, module, *, labelBase, electTag, wait=False): + """Emit the wave-0-elected cluster split-barrier arrive shared by the + StreamKMulticast prologue signal and prologue prefetch handshake. + + Wave election reuses the Serial/readfirstlane idiom the StreamK flag path + already uses (rather than sgpr("WaveIdx"), which may be undefined in the + epilogue): one wave per workgroup arrives (``s_barrier_signal -3``) while + the remaining waves branch over it via ``labelBase``. When ``wait`` is set + an all-waves cluster wait (``s_barrier_wait -3``) follows the arrive. + ``labelBase``/``electTag`` are supplied per call site so the emitted label + and pool tag -- and hence the assembly -- stay byte-identical to the + pre-extraction inline copies. Instructions are appended to ``module``. + """ + skipSignal = Label(label=writer.labels.getNameInc(labelBase), comment="") + elect = writer.sgprPool.checkOut(1, electTag) + module.add(VReadfirstlaneB32(dst=sgpr(elect), src=vgpr("Serial"), comment="wave 0 signals the cluster")) + module.add(SCmpEQU32(src0=sgpr(elect), src1=0, comment="Check for wave 0")) + module.add(SCBranchSCC0(labelName=skipSignal.getLabelName(), comment="only wave 0 signals the cluster")) + module.add(SBarrier(True, False, True, comment="cluster_barrier signal (arrive)")) + module.add(skipSignal) + if wait: + module.add(SBarrier(True, True, True, comment="cluster_barrier wait")) + writer.sgprPool.checkIn(elect) + return module + + def streamKMulticastPrologueSignal(self, writer, kernel): + """Elect wave 0 to arrive at the cluster split barrier once per workgroup. + + Supplies the prologue ``s_barrier_signal -3`` that the gfx1250 + cluster-barrier pass's first-load wait expects but that is otherwise + never anchored on the StreamKMulticast path (GlobalSplitU == 0). One + wave per workgroup arrives (others branch over it), uniformly across + peers, keeping cluster-scope signal/wait counts balanced. Inert unless + StreamKMulticast. See docs/design/cluster-load-component-and-streamk-multicast.md. + """ + module = Module("StreamK multicast prologue signal") + if not kernel.get("StreamKMulticast", 0): + return module + assert writer.states.asmCaps.get("HasClusterBarrier", False), \ + "StreamKMulticast requires the HasClusterBarrier asm capability" + module.addComment0("StreamKMulticast: elect wave 0 to signal the cluster barrier (pairs first-load wait)") + self._clusterElectArriveSignal( + writer, module, labelBase="SKMC_SkipSignal", electTag="SKMulticastElect") + return module + + def streamKMulticastProloguePrefetchHandshake(self, writer, kernel): + """Bracket the PGR>=2 prologue double-buffer prefetch multicast load with + a self-contained cluster-scope arrive/wait handshake. + + That "LDS1" prefetch load sits inside the single-iteration guard branch, + which the generic per-load bracketing's backward anchor scan stops at, so + it needs its own handshake (one wave arrives, all waves wait). The guard + branches on LoopCounterL, uniform across co-located peers, so peers run it + in lockstep. Inert unless StreamKMulticast. + See docs/design/cluster-load-component-and-streamk-multicast.md. + """ + module = Module("StreamK multicast prologue prefetch cluster handshake") + if not kernel.get("StreamKMulticast", 0): + return module + assert writer.states.asmCaps.get("HasClusterBarrier", False), \ + "StreamKMulticast requires the HasClusterBarrier asm capability" + module.addComment0("StreamKMulticast: bracket prologue double-buffer prefetch load with cluster handshake") + self._clusterElectArriveSignal( + writer, module, labelBase="SKMC_SkipPrefetchSignal", electTag="SKMulticastPrefetchElect", wait=True) + return module + + def streamKMulticastZeroIterClusterWait(self, writer, kernel): + """Consume the prologue cluster arrive on the zero-iteration skip path. + + The prologue arrive's only matching wait is the pass's first-load wait, + which sits after the last-iteration guard; on the zero-full-iteration + path (K not a whole multiple of DepthU) that guard skips the wait, + leaving the arrive unbalanced. Emit the matching all-waves wait on the + skip edge (scc1 == numIterL == 0; branch over it on scc0) so every peer + does exactly one arrive + one wait on every path. The wait leaves scc + intact for the following long branch. Inert unless StreamKMulticast. + See docs/design/cluster-load-component-and-streamk-multicast.md. + """ + module = Module("StreamK multicast zero-iteration cluster wait") + if not kernel.get("StreamKMulticast", 0): + return module + assert writer.states.asmCaps.get("HasClusterBarrier", False), \ + "StreamKMulticast requires the HasClusterBarrier asm capability" + module.addComment0("StreamKMulticast: zero-iteration skip path consumes the prologue cluster arrive (pairs prologue arrive)") + skipWait = Label(label=writer.labels.getNameInc("SKMC_SkipZeroIterClusterWait"), comment="") + module.add(SCBranchSCC0(labelName=skipWait.getLabelName(), + comment=">=1 full iteration: the first-load cluster wait pairs the arrive")) + module.add(SBarrier(True, True, True, comment="cluster_barrier wait")) + module.add(skipWait) + return module + def preLoop(self, writer, kernel): module = Module("StreamK TwoTileDPFirst openLoop") skConstsInVgprs = writer.isStreamKConstantsToVgprEnabled(kernel) @@ -2506,6 +2704,17 @@ def preLoop(self, writer, kernel): module.add(SMovB32(dst=sgpr("StreamKIdx"), src=sgpr("WorkGroup0"), comment="Save original StreamK index")) + # StreamKMulticast: with StreamKIdx now known, gate the DP B-broadcast + # mask on the runtime clusterMulticastValid predicate (full + M-aligned + # cluster). No-op unless StreamKMulticast is enabled. + module.add(self.streamKMulticastMaskPredicate(writer, kernel)) + + # StreamKMulticast: arrive once per workgroup at the cluster split + # barrier here in the prologue, before the first tensor_load_to_lds, so + # it pairs the cluster-barrier pass's first-load wait. No-op unless + # StreamKMulticast is enabled. + module.add(self.streamKMulticastPrologueSignal(writer, kernel)) + if kernel["StreamKForceDPOnly"]: sIdx = writer.acquireStreamKConstSgpr(kernel, "StreamKIdx") sIpt = writer.acquireStreamKConstSgpr(kernel, "ItersPerTile") @@ -2811,6 +3020,14 @@ def graWorkGroup(self, writer, kernel, tPA, tPB): module.add(SCmpLtU32(src0=sgpr("StreamKIter"), src1=sgpr(sTotalIters), comment="Make sure there's work to do")) module.add(writer.longBranchScc0(Label("KernelEnd", ""), posNeg=1)) + # DP->SK boundary: this WG just transitioned out of the DP region, so the + # spatial B-broadcast mask baked into the B descriptor is no longer valid + # (SK partial-tile peers do not co-issue the identical full-K B load). + # Clear it to a normal self-only B load for the remaining SK iterations. + # Reached only on the DP->SK switch path (the DP-continue / SK-continue + # cases branch to skUpdateDone above). No-op unless StreamKMulticast. + module.add(self.streamKMulticastBoundaryClear(writer, kernel)) + # If in SK, next iteration is sTmp+2 # Increment StreamK iteration module.add(skUpdateDone) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGREmit.py b/projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGREmit.py index f6c13ee53e84..9643a1d0183b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGREmit.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGREmit.py @@ -42,7 +42,7 @@ from math import ceil, log, log2, prod from rocisa.code import Label -from ...Common import INDEX_CHARS, clusterEnabled +from ...Common import INDEX_CHARS from ...Common.DataType import DataType @@ -1109,8 +1109,10 @@ def sizeRefName(idx): # OR the per-tensor broadcast mask into the descriptor for TDM multicast. # Subtile loads both A and B on every wave, so it uses split masks # (MulticastMask{tc}), not the non-subtile single parity mask. - if kernel["Multicast"] and clusterEnabled(kernel["ClusterDim"]): - mod.add(comp.setMulticastMask(descSgprName(1), f"MulticastMask{tc}", writer)) + from ...Components.ClusterLoad import ClusterLoadTDM + clusterComp = ClusterLoadTDM.find(writer) + if clusterComp: + mod.add(clusterComp.applyToDescriptor(writer, kernel, descSgprName(1), tc, subtile=True)) with writer.allocTmpSgpr(1) as tmpSgprRes: waveOffsetSgprIdx = tmpSgprRes.idx diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/__init__.py b/projects/hipblaslt/tensilelite/Tensile/Components/__init__.py index 43b0cbbffe48..9ebd19ba097a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/__init__.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/__init__.py @@ -52,4 +52,5 @@ def use(): pass "LSU", "TensorDataMover", "GL2Prefetch", + "ClusterLoad", ] diff --git a/projects/hipblaslt/tensilelite/Tensile/Contractions.py b/projects/hipblaslt/tensilelite/Tensile/Contractions.py index 645a5676c8a5..2bbf6ba5995f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Contractions.py +++ b/projects/hipblaslt/tensilelite/Tensile/Contractions.py @@ -627,6 +627,7 @@ class SizeMapping: 'streamK', 'streamKForceDPOnly', 'streamKAtomic', + 'streamKMulticast', 'prefetchAcrossPersistent', 'sourceKernel', 'globalAccumulation', @@ -721,6 +722,7 @@ def convertSFCWGMListToHex(value): streamK = d['StreamK'] if 'StreamK' in d else 0, streamKForceDPOnly = d.get('StreamKForceDPOnly', 0), streamKAtomic = d['StreamKAtomic'] if 'StreamKAtomic' in d else 0, + streamKMulticast = d.get('StreamKMulticast', 0), prefetchAcrossPersistent = d.get('PrefetchAcrossPersistent', 0), magicDivAlg = d.get('MagicDivAlg', 1), sourceKernel = d['KernelLanguage'] == 'Source', diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 01aa1c7226df..10ae8c7c1a71 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -48,6 +48,7 @@ from .Component import Component, LraTileProperties from .Components.Signature import UserArgumentsInfo from .Components.CustomSchedule import customMainLoopSchedule +from .Components.ClusterLoad import ClusterLoadTDM from .Components.StreamK import streamKVariantClass from .Components.Subtile.Kernel import * from .SolutionStructs import Solution, isPackedIndex @@ -2835,17 +2836,9 @@ def setupNewTile(self, kernel, tensorParametersA, tensorParametersB, isOptNLL=Fa module.add(self.graUnrollOffsets(kernel, tensorParametersB)) # Free sgpr that will not be used - if kernel["Multicast"] and kernel["TDMInst"] != 0: - tdmA: bool = kernel["enableTDMA"] - tdmB: bool = kernel["enableTDMB"] - tdmM: bool = kernel["enableTDMMetadata"] - if tdmA and tdmB and kernel["NumWaves"] > 1 and not kernel.get("UseSubtileImpl"): - module.add(self.undefineSgpr("MulticastMask")) - else: - module.add(self.undefineSgpr("MulticastMaskA")) - module.add(self.undefineSgpr("MulticastMaskB")) - if tdmM: - module.add(self.undefineSgpr("MulticastMaskMetadata")) + clusterComp = ClusterLoadTDM.find(self) + if clusterComp: + module.add(clusterComp.undeclareSgprs(self, kernel)) # tile edges if kernel["EdgeType"] == "ShiftPtr" and not tdmA and not tdmB: @@ -5473,6 +5466,15 @@ def kernelBody( self, kernel, tensorParametersA, tensorParametersB ): if kernel["PrefetchGlobalRead"] >= 2: for idxPgr in range(1, kernel["PrefetchGlobalRead"]): module.add(self.openPrefetchGlobalRead2orMore(kernel, idxPgr)) + # StreamKMulticast: the cooperative-multicast loads emitted below for + # this prefetch stage sit inside the single-iteration guard branch, + # past the generic per-load cluster-barrier bracketing boundary. + # Bracket them with a self-contained cluster-scope handshake so every + # multicast load stays synchronized and signal/wait counts stay + # balanced. Gated on StreamKMulticast (only ever set on the StreamK=3 + # component), so the emitted code is unchanged for every other path. + if kernel.get("StreamKMulticast", 0): + module.add(skComponent.streamKMulticastProloguePrefetchHandshake(self, kernel)) # For UnrollLoopSwapGlobalReadOrder, we also need to swap ds write A/B order. # In scheduling, we always schedule lwa first then lwb second, # Putting lwb in lwa's code object can easily change the order. @@ -6680,6 +6682,12 @@ def _restoreNtabState(): # Cluster-barrier handshake insertion in Gfx1250Backend # (kernel-scope at every OptLevel when set). "ClusterBarrier": bool(kernel.get("ClusterBarrier", False)), + # StreamKMulticast gates the per-iteration cooperative-broadcast + # drain in InsertClusterBarrierPass Rule 4 (mainloop) mode c: with + # PGR>=2 an `s_wait_tensorcnt 0` is emitted before the cluster-scope + # `s_barrier_signal -3` arrive so the broadcast retires before peers + # re-enter the next round. Defaults off; no-op for every other kernel. + "StreamKMulticast": bool(kernel.get("StreamKMulticast", 0)), # PrefetchGlobalRead (PGR) passed to InsertClusterBarrierPass. # Gates Rule 3 (`LCL <= PGR` skip) and Rule 4 (`LCL == PGR+1` # skip in fresh-gate mode; inherits upstream `LCL == PGR` cmp @@ -9218,20 +9226,9 @@ def vgprAllocationImplSubtile(): if kernel["enableTDMA"] or kernel["enableTDMB"]: self.defineSgpr("WaveIdx", 1) - if kernel["Multicast"]: - tdmA: bool = kernel["enableTDMA"] - tdmB: bool = kernel["enableTDMB"] - tdmM: bool = kernel["enableTDMMetadata"] - # Subtile issues both A and B loads on every wave (no wave-parity load - # split), so the single parity MulticastMask is wrong there -- it would - # OR one tensor's mask into both descriptors. Use the split A/B masks. - if tdmA and tdmB and kernel["NumWaves"] > 1 and not kernel.get("UseSubtileImpl"): - self.defineSgpr("MulticastMask", 1) - else: - self.defineSgpr("MulticastMaskA", 1) - self.defineSgpr("MulticastMaskB", 1) - if tdmM: - self.defineSgpr("MulticastMaskMetadata", 1) + clusterComp = ClusterLoadTDM.find(self) + if clusterComp: + clusterComp.declareSgprs(self, kernel) # SGPR above are user SGPR which are set by GPU hardware when the kernel is launched self.states.firstInitSgpr = self.sgprPool.size() @@ -10587,7 +10584,13 @@ def postMainLoopBarrierCheckAndReset(self, kernel, rootModule): if isinstance(item, Module): modulesToScan.append(item) keptItems.append(item) - elif isinstance(item, SBarrier): + elif isinstance(item, SBarrier) and "-3" not in str(item).split("//", 1)[0]: + # Pass-2 rebuilds only workgroup-scope barriers from token-state + # transitions, so only those are cleared here. Cluster-scope split + # barriers (s_barrier_signal/wait -3), e.g. the StreamKMulticast + # prologue arrive, are placed deliberately by other components and + # carry no LDS token, so preserve them rather than dropping a half of + # a cluster handshake. removedCount += 1 continue else: diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index 9b761f6369b9..57810e3884dc 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -75,6 +75,7 @@ from .Component import Component, TensorDataMover, GL2Prefetch from .Components.TensorDataMover import TensorDataMoverLoad from .Components.GL2Prefetch import GL2PrefetchLoad +from .Components.ClusterLoad import ClusterLoadTDM from .Components.GlobalWriteBatch import GlobalWriteBatchWriter from .KernelWriterModules import * from .AsmMemoryHelpers import dsStore, dsLoad, _vgprOffset @@ -2640,55 +2641,15 @@ def defineAndResources(self, kernel, tPA, tPB, tPM): comment="WorkGroup2 = (cluster_z * nwg_z) + wg_z")) moduleRegInit.add(label_calculate_workgroup_done) - if kernel["Multicast"]: - moduleRegInit.addComment0("Calculate multicast mask") - sgprWgX = sTmp+1 - sgprWgY = sTmp+2 - sgprNWgX = sTmp+3 - - maskA = 1 - for idx in range(kernel["ClusterDim"][1]): - maskA |= (1 << (idx * kernel["ClusterDim"][0])) - - maskB = (1 << kernel["ClusterDim"][0]) - 1 - - if kernel["enableTDMMetadata"]: - if kernel["ProblemType"]["Sparse"] == 1: - moduleRegInit.add(SLShiftLeftB32(dst=sgpr("MulticastMaskMetadata"), shiftHex=sgpr(sgprWgX), src=hex(maskA),\ - comment="Setting metadata mask (follows sparse A)")) - elif kernel["ProblemType"]["Sparse"] == 2: - moduleRegInit.add(SMulI32(dst=sgpr(sTmp+4), src0=sgpr(sgprWgY), src1=sgpr(sgprNWgX),\ - comment="Shift factor: wg_y * nwg_x (metadata)")) - moduleRegInit.add(SLShiftLeftB32(dst=sgpr("MulticastMaskMetadata"), shiftHex=sgpr(sTmp+4), src=hex(maskB),\ - comment="Setting metadata mask (follows sparse B)")) - - if tdmA and tdmB and kernel["NumWaves"] > 1 and not kernel.get("UseSubtileImpl"): - setMulticastMaskLblOdd = Label(f"setMulticastMask_OddWave", "") - setMulticastMaskLblEven = Label(f"setMulticastMask_EvenWave", "") - setMulticastMaskLblEnd = Label(f"setMulticastMaskEnd", "") - - moduleRegInit.add(SBitcmp1B32(sgpr("WaveIdx"), 0, "Check parity of wId")) - moduleRegInit.add(SCBranchSCC1(setMulticastMaskLblOdd.getLabelName(), "Jump if wId is odd")) - - moduleRegInit.add(setMulticastMaskLblEven) - moduleRegInit.add(SLShiftLeftB32(dst=sgpr("MulticastMask"), shiftHex=sgpr(sgprWgX), src=hex(maskA),\ - comment="Setting maskA for even wave")) - moduleRegInit.add(SBranch(setMulticastMaskLblEnd.getLabelName())) - moduleRegInit.add(setMulticastMaskLblOdd) - moduleRegInit.add(SMulI32(dst=sgpr(sgprWgY), src0=sgpr(sgprWgY), src1=sgpr(sgprNWgX),\ - comment="Shift factor: wg_y * nwg_x")) - moduleRegInit.add(SLShiftLeftB32(dst=sgpr("MulticastMask"), shiftHex=sgpr(sgprWgY), src=hex(maskB),\ - comment="Setting maskB for odd wave")) - moduleRegInit.add(setMulticastMaskLblEnd) - - else: - moduleRegInit.add(SLShiftLeftB32(dst=sgpr("MulticastMaskA"), shiftHex=sgpr(sgprWgX), src=hex(maskA),\ - comment="Setting maskA")) - - moduleRegInit.add(SMulI32(dst=sgpr(sgprWgY), src0=sgpr(sgprWgY), src1=sgpr(sgprNWgX),\ - comment="Shift factor: wg_y * nwg_x")) - moduleRegInit.add(SLShiftLeftB32(dst=sgpr("MulticastMaskB"), shiftHex=sgpr(sgprWgY), src=hex(maskB),\ - comment="Setting maskB")) + # Guard the compute site like the apply sites: find() returns None + # unless TDMInst==3 + HasTDM match, so Multicast with TDMInst in {1,2} + # or a non-TDM arch would otherwise None-deref here. + clusterComp = ClusterLoadTDM.find(self) + if kernel["Multicast"] and clusterComp: + # Same SGPR operands allocated above (wg_x=sTmp+1, wg_y=sTmp+2, + # nwg_x=sTmp+3, scratch=sTmp+4) are passed through. + moduleRegInit.add(clusterComp.computeMasks( + self, kernel, sgprWgX=sTmp+1, sgprWgY=sTmp+2, sgprNWgX=sTmp+3, sTmp=sTmp)) # SrdD can be used as temp sgprs for a bit if self.states.doShadowInit and kernel["BufferStore"]: self.addSgprVarToPool("SrdD") @@ -9733,6 +9694,15 @@ def openSumAtLeastUnroll(self, kernel, prefetch, isOptNLL, isNGLL=False, NLLinde if self.isPrefetchAcrossPersistentEnabled(kernel): module.add(SCMovB32(dst=sgpr("SkPrefetchPrimed"), src=0, comment="discard primed PAP group when current slice skips NLL")) + # StreamKMulticast: the long branch below skips the pass's first-load + # cluster wait on the zero-iteration path; emit the matching + # cluster-scope wait on that skip edge so the prologue cluster arrive + # is consumed on every control-flow path (whole-cluster barrier + # symmetry). scc (from checkLastIter) is preserved for the branch + # below. No-op unless StreamKMulticast. + if kernel.get("StreamKMulticast", 0): + skComponent = Component.StreamK.find(self) + module.add(skComponent.streamKMulticastZeroIterClusterWait(self, kernel)) # use positive offset only long jump with self.allocTmpSgpr(3, tag="openSumAtLeastUnroll_tmpSgprInfo") as tmpSgprInfo: module.add(self.longBranchScc1(lastIterEnd, posNeg=1, tmpSgprInfo=tmpSgprInfo)) @@ -18906,7 +18876,6 @@ def initTDMDescriptor(self, kernel: Mapping, tP: Mapping) -> Module: unrolledMajor = not tlu ti: int = tP["idx"] tileChar: str = tP["tileChar"] - enableCluster = clusterEnabled(kernel["ClusterDim"]) mod = Module(f"Init TDM Descriptor {tc}") def descSgprName(idx: int) -> str: @@ -18920,13 +18889,6 @@ def sizeRefName(idx: int) -> str: idxChar= INDEX_CHARS[idx] return f"Size{idxChar}" - def maskSgprName(tc: str) -> str: - if tc.startswith("MXS"): - string = tc.removeprefix("MXS") - else: - string = tc - return f"MulticastMask{string}" - dtype: DataType = kernel["ProblemType"][f"DataType{tc}"] mt: int = kernel[f"MacroTile{ti}"] du: int = kernel["DepthU"] @@ -18957,8 +18919,9 @@ def maskSgprName(tc: str) -> str: mod.add(comp.initOperands(descSgprName(0), descSgprName(1), group2Name, None)) mod.add(comp.setDataType(dtype, descSgprName(1), isMetadata)) mod.add(comp.setGlobalAddr(descSgprName(0), f"Address{tc}")) - if kernel["Multicast"] and enableCluster: - mod.add(comp.setMulticastMask(descSgprName(1), maskSgprName(tc), self)) + clusterComp = ClusterLoadTDM.find(self) + if clusterComp: + mod.add(clusterComp.applyToDescriptor(self, kernel, descSgprName(1), tc)) with self.allocTmpSgpr(2, tag="initTDMDescriptor_tmpSgprRes") as tmpSgprRes: waveOffsetSgprIdx: int = tmpSgprRes.idx @@ -19059,7 +19022,6 @@ def initTDMDescriptorWaveSeparatedImpl(self, kernel, tP, waveIdxSgpr: int | str unrolledMajor = not tlu ti: int = tP["idx"] tileChar: str = tP["tileChar"] - enableCluster = clusterEnabled(kernel["ClusterDim"]) mod = Module(f"Init TDM Descriptor {tc}") def descSgprName(idx: int) -> str: @@ -19073,9 +19035,6 @@ def sizeRefName(idx: int) -> str: idxChar= INDEX_CHARS[idx] return f"Size{idxChar}" - def maskSgprName(tc: str) -> str: - return f"MulticastMask" - dtype: DataType = kernel["ProblemType"][f"DataType{tc}"] mt: int = kernel[f"MacroTile{ti}"] du: int = kernel["DepthU"] @@ -19115,8 +19074,9 @@ def maskSgprName(tc: str) -> str: mod.add(comp.initOperands(descSgprName(0), descSgprName(1), group2Name, None)) mod.add(comp.setDataType(dtype, descSgprName(1), tc == "Metadata")) mod.add(comp.setGlobalAddr(descSgprName(0), f"Address{tc}")) - if kernel["Multicast"] and enableCluster: - mod.add(comp.setMulticastMask(descSgprName(1), maskSgprName(tc), self)) + clusterComp = ClusterLoadTDM.find(self) + if clusterComp: + mod.add(clusterComp.applyToDescriptor(self, kernel, descSgprName(1), tc, waveSeparated=True)) with self.allocTmpSgpr(2, tag="initTDMDescriptorWaveSeparatedImpl_tmpSgprRes") as tmpSgprRes: waveOffsetSgprIdx: int = tmpSgprRes.idx diff --git a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Naming.py b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Naming.py index 650e15ca75c9..a598d07357c5 100644 --- a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Naming.py +++ b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Naming.py @@ -213,6 +213,11 @@ def _getName(state, requiredParameters: frozenset, splitGSU: bool, ignoreInterna if state.get("LDSSegmentInterleave") == 1: requiredParametersTemp.add("LDSSegmentInterleave") + # Multicast is a tri-state derivation knob (-1 auto / 0 off / 1 on) that does + # not alter the emitted assembly, so it is intentionally excluded from the + # kernel name to keep names byte-identical to develop (no _M token). + requiredParametersTemp.discard("Multicast") + for key in sorted(requiredParametersTemp): if key not in state or key == "CustomKernelName": continue diff --git a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py index d06d896dbfd2..1e0c79da0720 100644 --- a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py +++ b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py @@ -231,6 +231,97 @@ def _validateStreamKForceDPOnly(state, printRejectionReason): return True +def _validateStreamKMulticast(state, printRejectionReason, isaInfoMap): + """Validate the gfx1250 StreamK DP cooperative cluster-load fast path. + + StreamKMulticast co-locates C consecutive StreamK DP workgroups in a 1-D + cluster (ClusterDim = [C, 1]); those M-adjacent tiles share the same B over + full K, so B is TDM-multicast to the cluster while A stays per-workgroup. + Solution-level requirements are rejected here at build time; the runtime + nWG0 % C "multiple-of-cluster-size" requirement is enforced by the + ClusterDimCheck predicate at selection time (not a silent fallback). + Auto-derived for StreamK=3 + ClusterDim != [1, 1] (the bare index-only cluster + state collapsed into this path), so the rejects below reject an unusable + cluster rather than an explicit opt-in. + See docs/design/cluster-load-component-and-streamk-multicast.md. + """ + if not state.get("StreamKMulticast", 0): + return True + + # StreamKMulticast is auto-enabled by ClusterDim on SK3, but the multicast mask + # SGPRs are gated on the (derived) Multicast flag while the mask predicate / + # boundary-clear emitters are gated on StreamKMulticast. An effective + # Multicast=off (force-off 0) therefore leaves MulticastMaskA/B undeclared while + # still being referenced -> broken codegen. Reject rather than emit it. + if not state.get("Multicast", 0): + reject(state, printRejectionReason, + "StreamKMulticast (auto-enabled by ClusterDim on SK3) is incompatible " + "with Multicast=0; use Multicast=-1 or 1, or remove ClusterDim.") + return False + + # SK3 (StreamKTwoTileDPFirst) only: the DP schedule + skIndexToWG addressing + # the mask derivation relies on are SK3-specific. + if state["StreamK"] != 3: + reject(state, printRejectionReason, + "StreamKMulticast requires StreamK=3 (two-tile DP-first)") + return False + + # The atomic path skips the workspace/tile DP structure the cooperative loads + # rely on. + if state["StreamKAtomic"]: + reject(state, printRejectionReason, + "StreamKMulticast is not supported with StreamKAtomic") + return False + + # StreamKXCCMapping remap is bypassed under clustering and XCC=3 overflows the + # SGPR budget alongside the cluster coords; require the default (no remap). + if state["StreamKXCCMapping"] != 0: + reject(state, printRejectionReason, + "StreamKMulticast requires StreamKXCCMapping=0 (WGM/XCC remap is bypassed under clustering)") + return False + + # 1-D cluster [C, 1] with C a power of two in [2, 16]. ClusterDim[1] == 1 + # because the StreamK grid is effectively 1-D along x and consecutive-WG + # clustering is what produces M-adjacent (shared-B) DP tiles. + clusterDim = state["ClusterDim"] + c = clusterDim[0] + if clusterDim[1] != 1: + reject(state, printRejectionReason, + "StreamKMulticast requires ClusterDim = [C, 1] (got %s)" % clusterDim) + return False + if c < 2 or c > 16 or (c & (c - 1)) != 0: + reject(state, printRejectionReason, + "StreamKMulticast requires ClusterDim[0] a power of two in [2, 16] (got %d)" % c) + return False + + # gfx1250 with TDM multicast loads (multicast is a TDM feature). + isa = tuple(state["ISA"]) + if isa != (12, 5, 0): + reject(state, printRejectionReason, + "StreamKMulticast requires gfx1250 ISA (12, 5, 0)") + return False + if not isaInfoMap[isa].asmCaps.get("HasTDM", False): + reject(state, printRejectionReason, + "StreamKMulticast requires asmCap HasTDM") + return False + # The cluster-scope barrier handshake that keeps the C multicast peers in + # lockstep around each tensor_load_to_lds needs the HasClusterBarrier asm cap. + if not isaInfoMap[isa].asmCaps.get("HasClusterBarrier", False): + reject(state, printRejectionReason, + "StreamKMulticast requires asmCap HasClusterBarrier (cluster-scope " + "barrier handshake around the multicast loads)") + return False + # ClusterLoadTDM (the component that emits/applies the multicast masks) matches + # only TDMInst == 3, so TDMInst in {1, 2} would produce no multicast component + # and silently drop the masks. Require TDMInst == 3. + if state["TDMInst"] != 3: + reject(state, printRejectionReason, + "StreamKMulticast requires TDMInst == 3 (TDM multicast loads on A and B)") + return False + + return True + + # _getExpectedTypes / _expectedParamTypes / _skipTypeCheck were moved into # Tensile/Common/ValidParameters.py to keep the registry and its derived # type map co-located (and to keep the Common -> Solution import direction). @@ -353,6 +444,21 @@ def validateParameterTypes(state, srcFile=""): return records +def coerceLegacyMulticastType(state) -> None: + """Normalize a legacy bool ``Multicast`` to the tri-state int, in place. + + Shipped library-logic artifacts serialize the derived value as a bool, but + ``Multicast`` is now an int ([-1, 0, 1]); bool and int are different msgpack + wire types, so a bool trips the create-library type gate (would ``std::bad_cast`` + at C++ deserialization). Coerce ``false -> 0`` / ``true -> 1`` (never -1 + "auto", to avoid switching a tuned kernel into legacy auto-coupling). Scoped to + the serialized-state loading path (the config/derivation path stays bool). + """ + mc = state.get("Multicast") + if type(mc) is bool: + state["Multicast"] = int(mc) + + def raiseIfTypeMismatches() -> None: """Raise when collected type mismatches exist. @@ -574,7 +680,17 @@ def __init__( self["AssignedProblemIndependentDerivedParameters"] = False if "AssignedDerivedParameters" not in self._state: self["AssignedDerivedParameters"] = False - + + # ``raiseProblemTypeOnTypeMismatch=False`` marks the paths that load a + # serialized / pre-derived solution state (library-logic artifacts, + # solution files, OriginalSolution). Those legacy artifacts carry a bool + # ``Multicast``; normalize it to the tri-state int so both the strict type + # gate and the emitted msgpack see an int. The config/derivation path + # (raiseProblemTypeOnTypeMismatch=True) is left untouched. + loadingSerializedState = not raiseProblemTypeOnTypeMismatch + if loadingSerializedState: + coerceLegacyMulticastType(self._state) + # Validate parameter types against the validParameters registry. # Catches bool-vs-int mismatches (YAML false vs 0) that would cause # std::bad_cast at C++ msgpack deserialization time. The mismatch @@ -1022,16 +1138,74 @@ def assignProblemIndependentDerivedParameters(state, printRejectionReason: bool, if state["DirectToVgprMXSA"] or state["DirectToVgprMXSB"]: reject(state, printRejectionReason, "UseSubtileImpl=1 PrefetchAcrossPersistent not supported with DirectToVgpr MX scale tensors") - state["Multicast"] = False state["ClusterBarrier"] = False - # Multicast uses a mask fixed to the physical cluster position, but Stream-K remaps - # each WG's tile per iteration, so the broadcast would target the wrong partner. - # Keep the cluster WG-id decode (gated on ClusterDim) but leave multicast off for Stream-K. - if state["ClusterDim"] != [1, 1] and state["StreamK"] == 0: - state["Multicast"] = True - # ClusterBarrier emits SCmp/branch on sgpr("WaveIdx"), which is only allocated when TDM is enabled. - if state["TDMInst"] != 0 and isaInfoMap[state["ISA"]].asmCaps.get("HasClusterBarrier", False): - state["ClusterBarrier"] = True + # StreamKMulticast is a DERIVED-ONLY internal state key (not a valid/benchmark + # parameter -- see ValidParameters.py). Seed it off here (mirroring the + # ClusterBarrier default above) so it is always present on state; the collapse + # below is the ONLY place it is turned on. + state["StreamKMulticast"] = 0 + # Collapse the bare StreamK cluster state: on StreamK=3 a non-[1,1] ClusterDim + # AUTO-ENABLES the DP cooperative B-multicast path. A StreamK cluster with no + # cooperative loads no longer exists -- "ClusterDim without cluster loads" is + # not a supported state. StreamKMulticast is derived-only (no user/YAML + # opt-in): this collapse is its sole enable site. + # + # If the resulting cooperative-load config cannot be satisfied (e.g. TDMInst!=3 + # or non-gfx1250), _validateStreamKMulticast (called later in + # assignDerivedParameters) rejects it at build time rather than silently + # degrading to a no-op cluster -- an unusable cluster is a hard reject. + # + # Scope note: the TDM B-multicast fast path is StreamK=3-only, so there is + # nothing to auto-enable for the dynamic (SK4) / hybrid (SK5) queue modes. + # SK4/SK5 with a non-[1,1] ClusterDim is rejected later in + # assignDerivedParameters (no cluster-load impl for those modes). + # ClusterDim is tested first so this (like the legacy Multicast auto branch + # below) never dereferences StreamK for the common non-clustered state, which + # some partial-state derivation call sites construct without a StreamK key. + if state["ClusterDim"] != [1, 1] and state.get("StreamK", 0) == 3: + state["StreamKMulticast"] = 1 + # Multicast tri-state (see ValidParameters): -1 auto (legacy), 0 off, 1 on. + # Default -1 reproduces the historic ClusterDim-coupled derivation, so YAML + # that omits Multicast is byte-identical. + mc = state.get("Multicast", -1) + if mc == 1: + # Force-on requires a matching ClusterLoadTDM (TDMInst==3 on gfx1250 with + # HasTDM); without it the multicast masks are never emitted or applied, so + # reject rather than silently generate a degenerate kernel. + isa = tuple(state["ISA"]) + if state["TDMInst"] != 3 or isa != (12, 5, 0) \ + or not isaInfoMap[isa].asmCaps.get("HasTDM", False): + reject(state, printRejectionReason, + "Multicast=1 requires TDMInst=3 on gfx1250 (HasTDM); " + "no cluster-load multicast component matches otherwise") + state["Multicast"] = 1 + elif mc == 0: + state["Multicast"] = 0 + elif state.get("StreamKMulticast", 0): + # StreamKMulticast (auto-derived above from StreamK=3 + ClusterDim) drives + # TDM B-multicast through the ClusterLoad component (its [C,1] cluster is + # spatial DP peers, not the legacy subtile coupling). The C co-resident + # peers must stay in lockstep around each multicast tensor_load_to_lds, so + # the cluster-scope barrier handshake is required (enabled below). + state["Multicast"] = 1 + else: # -1 auto (legacy) + # A legacy broadcast targets a fixed physical cluster position, which + # Stream-K tile remapping would send to the wrong partner, so auto-multicast + # is off for ALL Stream-K here; the derived StreamKMulticast is the one + # exception (branch above). Non-Stream-K clustered paths are unchanged. + state["Multicast"] = int(state["ClusterDim"] != [1, 1] + and state["StreamK"] == 0) + # The cluster-scope barrier handshake (s_barrier_signal/wait -3, inserted by + # StinkyTofu's InsertClusterBarrierPass around each multicast tensor_load_to_lds) + # keeps the C cluster peers in lockstep on the multicast loads. It applies to + # both clustered multicast paths: the legacy/subtile clustered multicast + # (StreamK == 0) and the StreamK DP cooperative multicast (StreamKMulticast). + # A forced-off Multicast keeps it off. + if state["ClusterDim"] != [1, 1] and state["Multicast"] \ + and state["TDMInst"] != 0 \ + and (state["StreamK"] == 0 or state.get("StreamKMulticast", 0)) \ + and isaInfoMap[state["ISA"]].asmCaps.get("HasClusterBarrier", False): + state["ClusterBarrier"] = True # done state["AssignedProblemIndependentDerivedParameters"] = True @@ -1670,11 +1844,16 @@ def assignDerivedParameters( state["GlobalSplitUAlgorithm"] = "MultipleBuffer" # Set default Algorithm state["AdaptiveGemmGSUA"] = 0 # Disable AdaptiveGemmGSUA for Stream-K if state["ClusterDim"] != [1, 1]: - # Only SK3 (two-tile DP-first) is cluster-aware; SK4 (dynamic per-XCD - # work queues) and SK5 (hybrid) have no cluster WG-id decode support. + # WG-cluster support is StreamK==3-only. The [C,1] cluster feature is the + # SK3 DP cooperative B-multicast (StreamKMulticast, auto-derived above); + # there is no cluster-load implementation for the dynamic (SK4) or hybrid + # (SK5) work-queue modes, so a ClusterDim there would only decode a + # cluster WG-id that no feature consumes. Reject outright rather than emit + # an unusable cluster kernel. if state["StreamK"] in (4, 5): reject(state, printRejectionReason, - "Stream-K modes 4 and 5 do not support ClusterDim != [1, 1]") + "StreamK dynamic/hybrid (SK4/SK5) do not support ClusterDim " + "(cluster support is SK3-only)") # Stream-K launches a 1-D grid in X, so StreamKIdx = WorkGroup0 = cluster_x*nwg_x # + wg_x must stay a unique linear index. A Y-extent > 1 collides WorkGroup0 across # WGs that differ only in Y, so restrict clustering to the X dimension. @@ -1705,6 +1884,7 @@ def assignDerivedParameters( if not state["BufferStore"]: reject(state, printRejectionReason, "Stream-K requires BufferStore") _validateStreamKForceDPOnly(state, printRejectionReason) + _validateStreamKMulticast(state, printRejectionReason, isaInfoMap) if state["StreamKAtomic"] == 1: if state["StreamK"] == 4: reject(state, printRejectionReason, "Atomic Stream-K is not supported with dynamic work queue mode") @@ -1771,6 +1951,7 @@ def assignDerivedParameters( state["StreamKAtomic"] = 0 state["StreamKXCCMapping"] = 0 state["StreamKFixupTreeReduction"] = 0 + state["StreamKMulticast"] = 0 state["DebugStreamK"] = 0 state["PrefetchAcrossPersistent"] = 0 state["DebugPersistentKernelLoopForever"] = False diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster_multicast.yaml similarity index 86% rename from projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster.yaml rename to projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster_multicast.yaml index 38eca11ac228..2cd34939272e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster_multicast.yaml @@ -1,5 +1,5 @@ TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] # not supported by arch + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] GlobalParameters: SyncsPerBenchmark: 0 @@ -19,8 +19,8 @@ GlobalParameters: ForceGenerateKernel: True BenchmarkProblems: - - # gfx1250 MXFP4, StreamK force DP-only persistent mode, 1-D ClusterDim range [2,1]/[4,1]/[8,1] - - # ProblemType, TN + - + - OperationType: GEMM DataType: F4 DestDataType: s @@ -36,8 +36,7 @@ BenchmarkProblems: MXBlockA: 32 MXBlockB: 32 - - # Persistent DP-only StreamK, 1-D ClusterDim range [2,1]/[4,1]/[8,1]; exercises the - # enableCluster guard in StreamKTwoTileDPFirst.preLoop. + - InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] @@ -59,9 +58,9 @@ BenchmarkProblems: - LdsBlockSizePerPadMXSB: [-1] - LdsPadMetadata: [0] - LDSTrInst: [False] - - PrefetchGlobalRead: [1] + - PrefetchGlobalRead: [1, 2] - PrefetchLocalRead: [1] - - ScheduleIterAlg: [3, 4] + - ScheduleIterAlg: [4] - SourceSwap: [False] - StoreRemapVectorWidth: [0] - GlobalSplitU: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4gemm_tdm_cluster.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4gemm_cluster_multicast.yaml similarity index 69% rename from projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4gemm_tdm_cluster.yaml rename to projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4gemm_cluster_multicast.yaml index f1b701f42855..0b0cf1d71b4a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4gemm_tdm_cluster.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4gemm_cluster_multicast.yaml @@ -1,7 +1,10 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] # not supported by arch + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] GlobalParameters: SyncsPerBenchmark: 0 + Architecture: gfx1250 NumElementsToValidate: -1 BoundsCheck: 0 KernelTime: False @@ -18,15 +21,8 @@ GlobalParameters: ForceGenerateKernel: True BenchmarkProblems: - ####################################### - # TN - StreamK=3, full tree-reduction (no StreamKForceDPOnly), 1-D ClusterDim - # range [2,1]/[4,1]/[8,1]. Same base sweep as sk_mxf4gemm_tdm.yaml with a 1-D - # ClusterDim range added: exercises the enableCluster guard in - # StreamKTwoTileDPFirst.preLoop on the general/split tile-splitting + fixup path - # (not just the DP-only grid-stride path) across cluster widths 2/4/8. - ####################################### - - - # ProblemType, TN + - OperationType: GEMM DataType: F4 DestDataType: s @@ -41,19 +37,19 @@ BenchmarkProblems: UseBias: 0 MXBlockA: 32 MXBlockB: 32 - - - # BenchmarkProblemSizeGroup - Standard + - InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] ForkParameters: - UseSgprForGRO: [0] - - ForceDisableShadowInit: [true] + - ForceDisableShadowInit: [True] - MatrixInstruction: - - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 2, 2, 2, 2] - - [16, 16, 128, 1, 1, 4, 4, 2, 2] - - [16, 16, 128, 1, 1, 8, 8, 2, 2] + - [16, 16, 128, 1, 1, 1, 1, 2, 2] + - [16, 16, 128, 1, 1, 2, 2, 2, 2] + - [16, 16, 128, 1, 1, 4, 4, 2, 2] + - [16, 16, 128, 1, 1, 8, 8, 2, 2] + - [16, 16, 128, 1, 1, 1, 1, 2, 1] - DepthU: [256] - ClusterLocalRead: [0] - TransposeLDS: [-1] @@ -67,34 +63,42 @@ BenchmarkProblems: - LdsBlockSizePerPadMXSB: [-1] - LdsPadMetadata: [0] - LDSTrInst: [True, False] - - TDMInst: [3] + - 1LDSBuffer: [0] + - DirectToVgprSparseMetadata: [False] - PrefetchGlobalRead: [1, 2] - PrefetchLocalRead: [1] - - ScheduleIterAlg: [0, 4] - - SourceSwap: [false] + - ScheduleIterAlg: [4] + - SourceSwap: [False] - StoreRemapVectorWidth: [0] - GlobalSplitU: [0] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - ExpandPointerSwap: [false] + - ExpandPointerSwap: [False] - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [0] - - DirectToVgprSparseMetadata: [false] - StoreVectorWidth: [-1] - StreamK: [3] + - StreamKAtomic: [0] + - StreamKXCCMapping: [0] + - ClusterDim: [[2, 1], [4, 1], [8, 1]] - StaggerU: [0] - DirectToVgprA: [False] - DirectToVgprB: [False] - WavefrontSize: [32] - WorkGroupMapping: [1] - - ClusterDim: [[2, 1], [4, 1], [8, 1]] + - TDMInst: [3] + - MXLoadInst: ["TDM"] + - MXScaleFormat: ["InMemorySwizzle"] BenchmarkForkParameters: JoinParameters: - BenchmarkJoinParameters: BenchmarkFinalParameters: - ProblemSizes: + - Exact: [512, 512, 1, 256] + - Exact: [256, 256, 1, 256] + - Exact: [128, 128, 1, 512] + - Exact: [256, 256, 2, 256] + - Exact: [256, 256, 1, 1920] - Exact: [2048, 256, 1, 1024] - Exact: [2048, 512, 1, 1056] - ActivationArgs: diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only_cluster_multicast.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only_cluster_multicast.yaml new file mode 100644 index 000000000000..529ccade9f6f --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only_cluster_multicast.yaml @@ -0,0 +1,94 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +TestParameters: + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] + +GlobalParameters: + SyncsPerBenchmark: 0 + NumElementsToValidate: -1 + BoundsCheck: 0 + KernelTime: False + CpuThreads: 32 + DataInitTypeA: 3 + DataInitTypeB: 3 + DataInitTypeD: 2 + DataInitTypeMXSA: 3 + DataInitTypeMXSB: 3 + DataInitTypeAlpha: 1 + DataInitTypeBeta: 1 + PrintLevel: 2 + PrintSolutionRejectionReason: True + ForceGenerateKernel: True + +BenchmarkProblems: + - + - + OperationType: GEMM + DataType: F8 + DestDataType: s + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + MXBlockA: 32 + MXBlockB: 32 + + - + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - MatrixInstruction: + - [16, 16, 128, 1, 1, 1, 1, 2, 2] + - DepthU: [256] + - ClusterLocalRead: [0] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMXSA: [-1] + - LdsPadMXSB: [-1] + - LdsBlockSizePerPadMXSA: [-1] + - LdsBlockSizePerPadMXSB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - PrefetchGlobalRead: [1, 2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [4] + - SourceSwap: [False] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [0] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - ExpandPointerSwap: [False] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - 1LDSBuffer: [0] + - DirectToVgprSparseMetadata: [false] + - StoreVectorWidth: [-1] + - StreamK: [3] + - StreamKForceDPOnly: [1] + - StaggerU: [0] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - WorkGroupMapping: [1] + - TDMInst: [3] + - MXLoadInst: ["TDM"] + - MXScaleFormat: ["InMemorySwizzle"] + - ClusterDim: [[2, 1], [4, 1], [8, 1]] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [256, 256, 1, 1056] + - Exact: [512, 512, 1, 1056] + - ActivationArgs: + - [Enum: none] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8gemm_cluster_multicast.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8gemm_cluster_multicast.yaml new file mode 100644 index 000000000000..1e5ff6abdd58 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8gemm_cluster_multicast.yaml @@ -0,0 +1,98 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +TestParameters: + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] +GlobalParameters: + SyncsPerBenchmark: 0 + Architecture: gfx1250 + NumElementsToValidate: -1 + BoundsCheck: 0 + KernelTime: False + CpuThreads: 32 + DataInitTypeA: 3 + DataInitTypeB: 3 + DataInitTypeD: 2 + DataInitTypeMXSA: 3 + DataInitTypeMXSB: 3 + DataInitTypeAlpha: 1 + DataInitTypeBeta: 1 + PrintLevel: 2 + PrintSolutionRejectionReason: True + ForceGenerateKernel: True + +BenchmarkProblems: + - + - + OperationType: GEMM + DataType: F8 + DestDataType: s + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + MXBlockA: 32 + MXBlockB: 32 + - + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - MatrixInstruction: + - [16, 16, 128, 1, 1, 1, 1, 2, 2] + - [16, 16, 128, 1, 1, 2, 2, 2, 2] + - [16, 16, 128, 1, 1, 4, 4, 2, 2] + - [16, 16, 128, 1, 1, 1, 1, 2, 1] + - DepthU: [256] + - ClusterLocalRead: [0] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMXSA: [-1] + - LdsPadMXSB: [-1] + - LdsBlockSizePerPadMXSA: [-1] + - LdsBlockSizePerPadMXSB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - PrefetchGlobalRead: [1, 2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [4] + - SourceSwap: [False] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [0] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - ExpandPointerSwap: [False] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - StoreVectorWidth: [-1] + - StreamK: [3] + - StreamKAtomic: [0] + - StreamKXCCMapping: [0] + - ClusterDim: [[2, 1], [4, 1], [8, 1]] + - StaggerU: [0] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - WorkGroupMapping: [1] + - TDMInst: [3] + - MXLoadInst: ["TDM"] + - MXScaleFormat: ["InMemorySwizzle"] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 256] + - Exact: [256, 256, 1, 256] + - Exact: [128, 128, 1, 512] + - Exact: [256, 256, 2, 256] + - Exact: [256, 256, 1, 1920] + - ActivationArgs: + - [Enum: none] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionArms/test_solution_arms_char.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionArms/test_solution_arms_char.py index 1ffd26f21e79..1ffb278e91b5 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionArms/test_solution_arms_char.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionArms/test_solution_arms_char.py @@ -47,6 +47,11 @@ def _reset(state): """Clear derivation-done flags so static methods re-run from scratch.""" state["AssignedProblemIndependentDerivedParameters"] = False state["AssignedDerivedParameters"] = False + # Multicast is a tri-state derived param (-1 auto / 0 off / 1 on). The vendored + # fixture states are already fully derived, so Multicast carries a concrete + # 0/1; restore the -1 auto sentinel a fresh (pre-derivation) config would have, + # so re-derivation exercises the legacy ClusterDim-coupled auto path. + state["Multicast"] = -1 return state @@ -378,7 +383,7 @@ def test_piap_cluster_dim_non_default_sets_multicast(isa_info_map, hss_state): state["ClusterBarrier"] = False state = _reset(state) _piap(state, isa_info_map) - assert state.get("Multicast") is True + assert state.get("Multicast") == 1 def test_piap_cluster_dim_default_multicast_false(isa_info_map, hss_state): @@ -388,7 +393,7 @@ def test_piap_cluster_dim_default_multicast_false(isa_info_map, hss_state): state["ClusterBarrier"] = False state = _reset(state) _piap(state, isa_info_map) - assert state.get("Multicast") is False + assert state.get("Multicast") == 0 def test_piap_cluster_barrier_no_cluster_dim_cleared(isa_info_map, hss_state): @@ -408,7 +413,7 @@ def test_piap_cluster_barrier_no_cluster_dim_cleared(isa_info_map, hss_state): _piap(state, isa_info_map) assert state.get("Valid") is not False assert state.get("ClusterBarrier") is False - assert state.get("Multicast") is False + assert state.get("Multicast") == 0 def test_piap_cluster_barrier_no_tdm_cleared(isa_info_map, hss_state): @@ -427,7 +432,7 @@ def test_piap_cluster_barrier_no_tdm_cleared(isa_info_map, hss_state): _piap(state, isa_info_map) assert state.get("Valid") is not False assert state.get("ClusterBarrier") is False - assert state.get("Multicast") is True + assert state.get("Multicast") == 1 # --------------------------------------------------------------------------- diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr index 71229dcd381a..81019888e28a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr @@ -55,7 +55,7 @@ 'getitem_kernel_language': 'Assembly', 'iter_matches_keys': True, 'keys_is_list': True, - 'len': 337, + 'len': 338, }) # --- # name: test_solution_construction @@ -295,6 +295,7 @@ 'StreamKAtomic', 'StreamKFixupTreeReduction', 'StreamKForceDPOnly', + 'StreamKMulticast', 'StreamKXCCMapping', 'SubGroup0', 'SubGroup1', @@ -400,7 +401,7 @@ 'tailLoopOptB', ]), 'name': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_SN_LDSB0_AA0_AFC1_AF1_AG0_AGGSUA0_AGNTAB0_AAIGTEn1_AAILTEn1_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DPKLF0_DSK0_DU32_DTL0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA1_GRVWB1_GSU1_GSUAMB_GSUC0_GSUWGMRR0_GLS0_HPLR0_ISA942_ICIW0_IU1_IA0_KLA_LDSSI0_LDSTI0_LBSPPA128_LBSPPB256_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA4_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW4_LRVWA4_LRVWB4_LWPMn1_MIAV0_MIWT1_1_MXLIBL_MXSFNS_MDA2_MI16_16_16_1_MLDS65536_MO40_MPM0_MGRIPM1_NR0_NTn1_NTA0_NTB0_NTC0_NTD4_NTE0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS2_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR1_PKA1_SFCWGM1_1_1_1_SGROB0_SGR1_SIA3_SLW1_SS1_SU32_SUM1_SUS256_SPO1_SRVW0_SSO0_SVW1_SK0_SKA0_SKFTR0_SKFDPO0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TT1_1_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UCMLS0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WSK0_WS64_WG32_8_1_WGM8_WGMXCC1_WGMXCCGn1_WGR0', - 'num_keys': 337, + 'num_keys': 338, 'problem_type': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs', 'stable': dict({ 'DepthU': 32, diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr index 99ff757917be..0c2e30b5d753 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr @@ -1432,6 +1432,7 @@ 'MaxOccupancy', 'MbskPrefetchMethod', 'MinGRIncPerMfma', + 'Multicast', 'NoReject', 'NonTemporal', 'NonTemporalA', @@ -2174,6 +2175,15 @@ 'len': 10, 'type': 'list', }), + 'Multicast': dict({ + 'head': list([ + -1, + 0, + 1, + ]), + 'len': 3, + 'type': 'list', + }), 'NoReject': dict({ 'head': list([ False, diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_gfx1250_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_gfx1250_char.ambr new file mode 100644 index 000000000000..5a24c0a83984 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_gfx1250_char.ambr @@ -0,0 +1,21 @@ +# serializer version: 1 +# name: test_streamk_cluster_multicast_gfx1250_golden + list([ + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg5n1_UAVQwNvJKjxzq3oOtWaugDzT0nK_8g8EArf7NFg=', + 'err': 0, + }), + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgAuSSqpdzu7alrTkBrgFPyFkqyqpG9E_Ty8RQ1xyXrY8=', + 'err': 0, + }), + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgchM4NognQc6wGYmCDZqnoSwpxn315nvCI1kv46p3XiY=', + 'err': 0, + }), + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgiDkOmo2IiGAh_e-YoomKdpTWhfgU7cQn5f7fD7vSzWM=', + 'err': 0, + }), + ]) +# --- diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_pgr2_gfx1250_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_pgr2_gfx1250_char.ambr new file mode 100644 index 000000000000..b63a612e4882 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_pgr2_gfx1250_char.ambr @@ -0,0 +1,21 @@ +# serializer version: 1 +# name: test_streamk_cluster_multicast_pgr2_gfx1250_golden + list([ + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg5tRpHXBjULkMIXSNVBHJHFmD1t3k0UwLm2yJZP4oexA=', + 'err': 0, + }), + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgGNWGjVlzVS5wHYrKb7__g3mxK0KG3m4i23NRg-javuY=', + 'err': 0, + }), + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgRlD1yh-2ZySyGsxGWXEw0e-4y6cBU5fJKfN-NANMhEA=', + 'err': 0, + }), + dict({ + 'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgSiq6ezRM6zYVfmOWJNJGQ9FXlCFA4RsbFcRyfLeirjs=', + 'err': 0, + }), + ]) +# --- diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_coop_load.yaml similarity index 85% rename from projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster.yaml rename to projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_coop_load.yaml index 50a6abddc9fa..0b757d2d3f1b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_coop_load.yaml @@ -2,10 +2,6 @@ # Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. # SPDX-License-Identifier: MIT ################################################################################ -# gfx1250 StreamK + ClusterDim=[2,1] codegen characterization config. -# Normal StreamK=3 (full tree-reduction). Exercises the enableCluster guard in -# StreamKTwoTileDPFirst.preLoop. Mirrors Tests/common/streamk/gfx1250/core/ -# sk_mxf4gemm_tdm.yaml with ClusterDim added. GlobalParameters: SyncsPerBenchmark: 0 MinimumRequiredVersion: 5.0.0 @@ -16,7 +12,7 @@ GlobalParameters: BenchmarkProblems: - - - # ProblemType, TN, MX-FP4 + - OperationType: GEMM DataType: F4 DestDataType: s @@ -31,7 +27,7 @@ BenchmarkProblems: UseBias: 0 MXBlockA: 32 MXBlockB: 32 - - # BenchmarkProblemSizeGroup: StreamK=3, ClusterDim=[2,1] + - InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] @@ -55,7 +51,7 @@ BenchmarkProblems: - LDSTrInst: [False] - PrefetchGlobalRead: [1] - PrefetchLocalRead: [1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [4] - SourceSwap: [False] - StoreRemapVectorWidth: [0] - GlobalSplitU: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast.yaml new file mode 100644 index 000000000000..4d2560a57cbb --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast.yaml @@ -0,0 +1,79 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +GlobalParameters: + SyncsPerBenchmark: 0 + MinimumRequiredVersion: 5.0.0 + NumElementsToValidate: 128 + DataInitTypeBeta: 0 + DataInitTypeAlpha: 1 + Device: 0 + +BenchmarkProblems: + - + - + OperationType: GEMM + DataType: F8 + DestDataType: s + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + MXBlockA: 32 + MXBlockB: 32 + - + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - MatrixInstruction: + - [16, 16, 128, 1, 1, 1, 1, 2, 2] + - [16, 16, 128, 1, 1, 2, 2, 2, 2] + - [16, 16, 128, 1, 1, 4, 4, 2, 2] + - [16, 16, 128, 1, 1, 1, 1, 2, 1] + - DepthU: [256] + - ClusterLocalRead: [0] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMXSA: [-1] + - LdsPadMXSB: [-1] + - LdsBlockSizePerPadMXSA: [-1] + - LdsBlockSizePerPadMXSB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - PrefetchGlobalRead: [1] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [4] + - SourceSwap: [False] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [0] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - ExpandPointerSwap: [False] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - StoreVectorWidth: [-1] + - StreamK: [3] + - StreamKAtomic: [0] + - StreamKXCCMapping: [0] + - ClusterDim: [[4, 1]] + - StaggerU: [0] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - WorkGroupMapping: [1] + - TDMInst: [3] + - MXLoadInst: ["TDM"] + - MXScaleFormat: ["InMemorySwizzle"] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 256] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast_pgr2.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast_pgr2.yaml new file mode 100644 index 000000000000..f51da5ac1c19 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast_pgr2.yaml @@ -0,0 +1,79 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +GlobalParameters: + SyncsPerBenchmark: 0 + MinimumRequiredVersion: 5.0.0 + NumElementsToValidate: 128 + DataInitTypeBeta: 0 + DataInitTypeAlpha: 1 + Device: 0 + +BenchmarkProblems: + - + - + OperationType: GEMM + DataType: F8 + DestDataType: s + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + MXBlockA: 32 + MXBlockB: 32 + - + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - MatrixInstruction: + - [16, 16, 128, 1, 1, 1, 1, 2, 2] + - [16, 16, 128, 1, 1, 2, 2, 2, 2] + - [16, 16, 128, 1, 1, 4, 4, 2, 2] + - [16, 16, 128, 1, 1, 1, 1, 2, 1] + - DepthU: [256] + - ClusterLocalRead: [0] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMXSA: [-1] + - LdsPadMXSB: [-1] + - LdsBlockSizePerPadMXSA: [-1] + - LdsBlockSizePerPadMXSB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - PrefetchGlobalRead: [2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [4] + - SourceSwap: [False] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [0] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - ExpandPointerSwap: [False] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - StoreVectorWidth: [-1] + - StreamK: [3] + - StreamKAtomic: [0] + - StreamKXCCMapping: [0] + - ClusterDim: [[4, 1]] + - StaggerU: [0] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - WorkGroupMapping: [1] + - TDMInst: [3] + - MXLoadInst: ["TDM"] + - MXScaleFormat: ["InMemorySwizzle"] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 512] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_coop_load_gfx1250_char.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_coop_load_gfx1250_char.py new file mode 100644 index 000000000000..cbfc0f6167f9 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_coop_load_gfx1250_char.py @@ -0,0 +1,110 @@ +################################################################################ +# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +################################################################################ +"""gfx1250 StreamK=3 + ClusterDim cooperative-load codegen characterization. + +A StreamK=3 kernel with ClusterDim != [1, 1] now +AUTO-ENABLES the DP cooperative B-multicast fast path (the "bare StreamK cluster" +state was collapsed; StreamKMulticast is a derived-only internal state with no +YAML opt-in). This test therefore verifies BOTH halves of that path from a config +that only sets ClusterDim: + + * the cluster WG-id decode (RemapWorkGroupDone) and the skipped ttmp9/ttmp7 + preLoop reread ("workaround" absent) under clustering -- defineAndResources + leaves the cluster-decoded rank in WorkGroup0/1/2, so the reread guard skips + it; and + * the auto-derived multicast masks: the split B-broadcast / A-self masks, the + runtime clusterMulticastValid predicate, and the DP->SK boundary clear. + +Uses MX-FP4 with ClusterDim=[2, 1] (C=2), complementing the MX-FP8 C=4 coverage +in the sibling test_streamk_cluster_multicast_gfx1250_char.py. +""" + +import os + +import pytest + +from config_harness import emit_kernels_from_config + +pytestmark = pytest.mark.unit + +_ARCH = "gfx1250" + +_CONFIG = os.path.join( + os.path.dirname(__file__), + "data", + "test_data", + "_designed", + "gfx1250", + "streamk_cluster_coop_load.yaml", +) + + +def test_streamk_cluster_coop_load_gfx1250_emits_assembly(): + """StreamK=3 + ClusterDim=[2,1]: cluster decode present, preLoop reread + skipped, and the auto-derived DP cooperative B-multicast path emitted.""" + results = emit_kernels_from_config(_CONFIG, limit=8, arch=_ARCH) + assert len(results) >= 1, f"Expected >=1 kernel, got {len(results)}" + assert all(err == 0 for (_b, _s, err) in results), ( + "All kernels must emit with err==0; " + + str([(b, e) for (b, _s, e) in results if e != 0]) + ) + for base, src, _err in results: + assert src and len(src.splitlines()) > 50, ( + f"Kernel {base!r} emitted suspiciously short source" + ) + assert ".amdgcn_target" in src, f"Kernel {base!r} missing .amdgcn_target" + assert "gfx1250" in src, f"Kernel {base!r} missing gfx1250 arch marker" + # Cluster WG-id decode arm. + assert "RemapWorkGroupDone" in src, ( + f"Kernel {base!r}: missing cluster WG-id decode ('RemapWorkGroupDone')" + ) + # preLoop ttmp reread must be skipped under clustering ("workaround" is + # unique to that reread block). + assert "workaround" not in src, ( + f"Kernel {base!r}: ttmp reread emitted under ClusterDim != [1, 1]" + ) + # Auto-derived DP cooperative B-multicast: B descriptor carries the + # broadcast mask, A the self-only mask (the split topology). + assert "s[sgprtdmBGroup1], s[sgprtdmBGroup1], s[sgprMulticastMaskB]" in src, ( + f"Kernel {base!r}: missing auto-derived B-broadcast mask (StreamKMulticast)" + ) + assert "s[sgprtdmAGroup1], s[sgprtdmAGroup1], s[sgprMulticastMaskA]" in src, ( + f"Kernel {base!r}: missing self-only mask on the A descriptor" + ) + # Runtime clusterMulticastValid predicate gates the broadcast. + assert "nWG0 aligned to C?" in src, ( + f"Kernel {base!r}: missing multicast M-alignment predicate" + ) + assert "cluster fully populated?" in src, ( + f"Kernel {base!r}: missing multicast cluster-population predicate" + ) + # DP->SK boundary clear drops the B broadcast for the SK region. + assert "DP->SK: drop B broadcast -> self-only" in src, ( + f"Kernel {base!r}: missing DP->SK boundary mask clear" + ) + # The multicast loads carry the cluster-scope barrier handshake + # (s_barrier_signal/wait -3) that keeps the C cluster peers in lockstep. + assert "s_barrier_signal -3" in src, ( + f"Kernel {base!r}: missing cluster-scope barrier signal (-3)" + ) + assert "s_barrier_wait -3" in src, ( + f"Kernel {base!r}: missing cluster-scope barrier wait (-3)" + ) + # Cluster-scope split-barrier balance. Every arrive + # (s_barrier_signal -3) must be consumed by a completion + # (s_barrier_wait -3) on every control-flow path. The prologue wave-0 + # arrive is consumed by exactly one of two mutually exclusive cluster + # waits: the last-iteration guard's zero-iteration skip-edge wait, or + # the first-load wait on the >=1-iteration fall-through. Both waits are + # emitted statically but only one executes on any given path, so the + # static wait count is exactly one greater than the signal count; every + # other arrive is a self-contained arrive/wait pair. Any other imbalance + # would leave a cluster wait unpaired and stall the cluster waves. + n_signal = src.count("s_barrier_signal -3") + n_wait = src.count("s_barrier_wait -3") + assert n_wait == n_signal + 1, ( + f"Kernel {base!r}: unexpected cluster barrier balance: " + f"{n_signal} signal(-3) vs {n_wait} wait(-3) (expected wait == signal + 1)" + ) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_gfx1250_char.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_gfx1250_char.py deleted file mode 100644 index 36ebb6bcf094..000000000000 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_gfx1250_char.py +++ /dev/null @@ -1,59 +0,0 @@ -################################################################################ -# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. -# SPDX-License-Identifier: MIT -################################################################################ -"""gfx1250 StreamK + ClusterDim codegen characterization. - -Checks that a StreamK kernel with ClusterDim != [1, 1] emits the cluster -WG-id decode (RemapWorkGroupDone) and skips the raw ttmp9/ttmp7 reread in the -StreamK preLoop. Under clustering, defineAndResources leaves the cluster-decoded -rank in WorkGroup0/1/2; the reread would overwrite it, so the enableCluster -guard skips it. - -Uses normal StreamK=3 with ClusterDim=[2, 1]; the guard applies equally to -StreamK modes 3, 4, and 5. -""" - -import os - -import pytest - -from config_harness import emit_kernels_from_config - -pytestmark = pytest.mark.unit - -_ARCH = "gfx1250" - -_CONFIG = os.path.join( - os.path.dirname(__file__), - "data", - "test_data", - "_designed", - "gfx1250", - "streamk_cluster.yaml", -) - - -def test_streamk_cluster_gfx1250_emits_assembly(): - """StreamK=3 + ClusterDim=[2,1]: cluster decode present, preLoop reread skipped.""" - results = emit_kernels_from_config(_CONFIG, limit=8, arch=_ARCH) - assert len(results) >= 1, f"Expected >=1 kernel, got {len(results)}" - assert all(err == 0 for (_b, _s, err) in results), ( - "All kernels must emit with err==0; " - + str([(b, e) for (b, _s, e) in results if e != 0]) - ) - for base, src, _err in results: - assert src and len(src.splitlines()) > 50, ( - f"Kernel {base!r} emitted suspiciously short source" - ) - assert ".amdgcn_target" in src, f"Kernel {base!r} missing .amdgcn_target" - assert "gfx1250" in src, f"Kernel {base!r} missing gfx1250 arch marker" - # Cluster WG-id decode arm. - assert "RemapWorkGroupDone" in src, ( - f"Kernel {base!r}: missing cluster WG-id decode ('RemapWorkGroupDone')" - ) - # preLoop ttmp reread must be skipped under clustering ("workaround" is - # unique to that reread block). - assert "workaround" not in src, ( - f"Kernel {base!r}: ttmp reread emitted under ClusterDim != [1, 1]" - ) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_gfx1250_char.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_gfx1250_char.py new file mode 100644 index 000000000000..de80f48908df --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_gfx1250_char.py @@ -0,0 +1,117 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +"""StreamK DP cooperative B-multicast -- gfx1250 characterization (CPU-only). + +Exercises the StreamK DP cooperative-load fast path added to +``Tensile/Components/StreamK.py`` + ``Tensile/Components/ClusterLoad.py`` +(the derived StreamKMulticast on-state: ClusterDim=[C,1] +on SK3, which auto-derives StreamKMulticast=1). It drives the same config -> +Solutions -> emit path as the sibling +``test_streamk_cluster_coop_load_gfx1250_char.py``, but through the (explicitly +named) multicast designed config ``_designed/gfx1250/streamk_cluster_multicast.yaml``. + +Asserts: + * every kernel emits real gfx1250 assembly with ``err == 0``; + * the B TDM descriptor carries the split B-broadcast mask + (``MulticastMaskB`` OR'd into ``tdmBGroup1``) while the A descriptor carries + the self-only ``MulticastMaskA`` -- i.e. B is multicast, A is per-WG; + * the runtime ``clusterMulticastValid`` predicate gates the broadcast (M + alignment + fully-populated cluster, else self-only B); and + * the DP->SK boundary clear drops the B broadcast for the SK region; and + * an order-invariant ``{basename, err}`` syrupy snapshot is pinned. + +CPU-only: no GPU required. The emit harness instantiates rocisa and runs +Python+rocisa codegen without compiling or launching any GPU kernels. +""" + +import os + +import pytest + +from config_harness import emit_kernels_from_config + +pytestmark = pytest.mark.unit + +_ARCH = "gfx1250" + +_CONFIG = os.path.join( + os.path.dirname(__file__), + "data", + "test_data", + "_designed", + "gfx1250", + "streamk_cluster_multicast.yaml", +) + + +def test_streamk_cluster_multicast_gfx1250_emits_assembly(): + """gfx1250 SK3 cluster config (StreamKMulticast auto-derived on) emits real + assembly, err==0, with the split B-broadcast mask, runtime predicate, and + DP->SK boundary clear.""" + results = emit_kernels_from_config(_CONFIG, limit=8, arch=_ARCH) + assert len(results) >= 1, "Expected >=1 kernel, got 0" + assert all(err == 0 for (_b, _s, err) in results), ( + f"Expected all err==0, got: {[(b, e) for b, _s, e in results if e != 0]}" + ) + for base, src, _err in results: + assert src and len(src.splitlines()) > 50, ( + f"Kernel {base!r} emitted suspiciously short source" + ) + assert ".amdgcn_target" in src, f"Kernel {base!r} missing .amdgcn_target" + assert "gfx1250" in src, f"Kernel {base!r} missing gfx1250 target" + assert base.startswith("Cijk_"), f"Kernel {base!r} has unexpected prefix" + # Split A/B multicast: B broadcast, A self-only. + assert "s[sgprtdmBGroup1], s[sgprtdmBGroup1], s[sgprMulticastMaskB]" in src, ( + f"Kernel {base!r} missing B-broadcast mask on the B descriptor" + ) + assert "s[sgprtdmAGroup1], s[sgprtdmAGroup1], s[sgprMulticastMaskA]" in src, ( + f"Kernel {base!r} missing self-only mask on the A descriptor" + ) + # Runtime clusterMulticastValid predicate. + assert "nWG0 aligned to C?" in src, ( + f"Kernel {base!r} missing multicast M-alignment predicate" + ) + assert "cluster fully populated?" in src, ( + f"Kernel {base!r} missing multicast cluster-population predicate" + ) + # DP->SK boundary clear (inline comment survives canonicalization; the + # addComment0 banner does not). + assert "DP->SK: drop B broadcast -> self-only" in src, ( + f"Kernel {base!r} missing DP->SK boundary mask clear" + ) + # The multicast tensor_load_to_lds is wrapped by the cluster-scope barrier + # handshake (s_barrier_signal/wait -3) that keeps the C cluster peers in + # lockstep on the multicast loads. + assert "s_barrier_signal -3" in src, ( + f"Kernel {base!r} missing cluster-scope barrier signal (-3)" + ) + assert "s_barrier_wait -3" in src, ( + f"Kernel {base!r} missing cluster-scope barrier wait (-3)" + ) + # Cluster-scope split-barrier balance. Every arrive + # (s_barrier_signal -3) must be consumed by a completion + # (s_barrier_wait -3) on every control-flow path. The prologue wave-0 + # arrive is consumed by exactly one of two mutually exclusive cluster + # waits: the last-iteration guard's zero-iteration skip-edge wait, or + # the first-load wait on the >=1-iteration fall-through. Both waits are + # emitted statically but only one executes on any given path, so the + # static wait count is exactly one greater than the signal count; every + # other arrive is a self-contained arrive/wait pair. Any other imbalance + # would leave a cluster wait unpaired and stall the cluster waves. + n_signal = src.count("s_barrier_signal -3") + n_wait = src.count("s_barrier_wait -3") + assert n_wait == n_signal + 1, ( + f"Kernel {base!r} has unexpected cluster barrier balance: " + f"{n_signal} signal(-3) vs {n_wait} wait(-3) (expected wait == signal + 1)" + ) + + +def test_streamk_cluster_multicast_gfx1250_golden(snapshot): + """Golden: order-invariant {basename, err} digest of the multicast emit.""" + results = emit_kernels_from_config(_CONFIG, limit=8, arch=_ARCH) + digest = sorted( + ({"basename": b, "err": e} for (b, _s, e) in results), + key=lambda d: d["basename"], + ) + assert digest == snapshot diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_pgr2_gfx1250_char.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_pgr2_gfx1250_char.py new file mode 100644 index 000000000000..19bdd856b789 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_pgr2_gfx1250_char.py @@ -0,0 +1,120 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +"""StreamK DP cooperative B-multicast, double-buffered prologue prefetch +(PrefetchGlobalRead=2) -- gfx1250 characterization (CPU-only). + +Companion to ``test_streamk_cluster_multicast_gfx1250_char.py`` (which pins +PrefetchGlobalRead=1). This exercises the ``PrefetchGlobalRead=2`` path with +K > DepthU, so the prologue emits the second, double-buffered ("LDS1") +cooperative multicast prefetch load. That prefetch load is emitted inside the +single-iteration guard branch, past the generic per-load cluster-barrier +bracketing boundary, so ``StreamK.streamKMulticastProloguePrefetchHandshake`` +must bracket it with a dedicated cluster-scope split-barrier handshake. + +Asserts, in addition to the sibling config's split-mask / predicate checks: + * the double-buffered prologue prefetch load is bracketed by a wave-0-elected + cluster-scope ``s_barrier_signal -3`` + ``s_barrier_wait -3`` handshake + (label ``SKMC_SkipPrefetchSignal``); and + * the cluster-scope split barrier stays balanced: every ``s_barrier_signal -3`` + is matched by an ``s_barrier_wait -3``. + +CPU-only: no GPU required. +""" + +import os + +import pytest + +from config_harness import emit_kernels_from_config + +pytestmark = pytest.mark.unit + +_ARCH = "gfx1250" + +_CONFIG = os.path.join( + os.path.dirname(__file__), + "data", + "test_data", + "_designed", + "gfx1250", + "streamk_cluster_multicast_pgr2.yaml", +) + + +def _skip_prefetch_handshake_brackets_load(src): + """Return True iff the prologue prefetch (LDS1) multicast load is bracketed. + + The double-buffered prologue prefetch load sits in the ``skipPGR2`` guard + segment. The dedicated handshake elects wave 0 (branch to + ``SKMC_SkipPrefetchSignal``), signals ``-3``, then all waves wait ``-3`` + immediately before the LDS1 ``tensor_load_to_lds`` group. + """ + lines = src.splitlines() + for i, ln in enumerate(lines): + if "label_SKMC_SkipPrefetchSignal:" not in ln: + continue + # A cluster-scope wait must follow the skip label, before the LDS1 load. + window = lines[i : i + 6] + has_wait = any("s_barrier_wait -3" in w for w in window) + has_load = any("tensor_load_to_lds" in w for w in window) + # A wave-0 signal must precede the skip label. + pre = lines[max(0, i - 4) : i] + has_signal = any("s_barrier_signal -3" in p for p in pre) + if has_wait and has_load and has_signal: + return True + return False + + +def test_streamk_cluster_multicast_pgr2_gfx1250_emits_assembly(): + """gfx1250 SK3 cluster multicast config at PrefetchGlobalRead=2 (K>DepthU) + emits real assembly with err==0, and the double-buffered prologue prefetch + multicast load is bracketed by a balanced cluster-scope -3 handshake.""" + results = emit_kernels_from_config(_CONFIG, limit=8, arch=_ARCH) + assert len(results) >= 1, "Expected >=1 kernel, got 0" + assert all(err == 0 for (_b, _s, err) in results), ( + f"Expected all err==0, got: {[(b, e) for b, _s, e in results if e != 0]}" + ) + for base, src, _err in results: + assert ".amdgcn_target" in src, f"Kernel {base!r} missing .amdgcn_target" + assert "gfx1250" in src, f"Kernel {base!r} missing gfx1250 target" + # Split A/B multicast: B broadcast, A self-only. + assert "s[sgprtdmBGroup1], s[sgprtdmBGroup1], s[sgprMulticastMaskB]" in src, ( + f"Kernel {base!r} missing B-broadcast mask on the B descriptor" + ) + # The PGR>=2 prologue double-buffer prefetch region must exist for K>DepthU. + assert "skipPGR2" in src, ( + f"Kernel {base!r} missing the PGR2 prologue double-buffer region" + ) + # The double-buffered prologue prefetch multicast load must be bracketed + # by a dedicated cluster-scope handshake (the fix under test). + assert _skip_prefetch_handshake_brackets_load(src), ( + f"Kernel {base!r} PGR2 prologue prefetch load is NOT bracketed by a " + f"cluster-scope -3 handshake" + ) + # Cluster-scope split-barrier balance. The prologue wave-0 arrive + # (s_barrier_signal -3) is consumed by exactly one of two mutually + # exclusive cluster waits: the last-iteration guard's zero-iteration + # skip-edge wait, or the first-load wait on the >=1-iteration + # fall-through. Both are emitted statically but only one executes on any + # control-flow path, so the static wait count is exactly one greater + # than the signal count; every other arrive (including this config's + # dedicated prologue-prefetch handshake) is a self-contained arrive/wait + # pair. Any other imbalance would leave a cluster wait unpaired and + # stall the cluster waves. + n_signal = src.count("s_barrier_signal -3") + n_wait = src.count("s_barrier_wait -3") + assert n_wait == n_signal + 1, ( + f"Kernel {base!r} has unexpected cluster barrier balance: " + f"{n_signal} signal(-3) vs {n_wait} wait(-3) (expected wait == signal + 1)" + ) + + +def test_streamk_cluster_multicast_pgr2_gfx1250_golden(snapshot): + """Golden: order-invariant {basename, err} digest of the PGR2 multicast emit.""" + results = emit_kernels_from_config(_CONFIG, limit=8, arch=_ARCH) + digest = sorted( + ({"basename": b, "err": e} for (b, _s, e) in results), + key=lambda d: d["basename"], + ) + assert digest == snapshot diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PrefetchAcrossPersistent.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PrefetchAcrossPersistent.py index a337cb36f015..e1c6f6987db3 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PrefetchAcrossPersistent.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PrefetchAcrossPersistent.py @@ -164,6 +164,11 @@ def __init__(self): memTokenLdsBuffer1=1, staggerUCode=False, unrollIdx=0, + # Capability/kernel state consumed by ClusterLoadTDM.find()'s + # PartialMatch (asmCaps HasTDM + kernel TDMInst==3), mirroring the + # real writer.states on a gfx1250 TDM path so the component matches. + asmCaps={"HasTDM": True}, + kernel={"TDMInst": 3}, ) self.do = {"executeToInitEnd": False} self.dontAppendCode = False diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_cluster_load_component.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_cluster_load_component.py new file mode 100644 index 000000000000..b42fc9b2cefd --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_cluster_load_component.py @@ -0,0 +1,306 @@ +#!/usr/bin/env python3 +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +# Unit tests for the ClusterLoad (TDM multicast) component. +# +# Covers Tensile/Components/ClusterLoad.py: capability-based selection, the +# topology decision (usesCombinedMask / maskSgprName), the cooperative-thread +# partition, the SGPR declare/undeclare, and the emitted assembly of +# computeMasks / applyToDescriptor. These emit no GPU work themselves, so the +# asm string is the contract -- easy to break silently. +# +# Usage: +# pytest test_cluster_load_component.py -v +################################################################################ + +import os +import shutil +import sys + +import pytest +from types import SimpleNamespace + +pytestmark = pytest.mark.unit + +SCRIPT_DIR = os.path.dirname(os.path.abspath(__file__)) +TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", "..")) +sys.path.insert(0, TENSILE_ROOT) + +WAVESIZE_32 = 32 + + +def _init_rocisa_gfx1250(): + from rocisa import rocIsa + from Tensile.Common.Architectures import gfxToIsa + ri = rocIsa.getInstance() + isa = gfxToIsa("gfx1250") + asmpath = shutil.which('amdclang++') or '/usr/bin/amdclang++' + ri.init(isa, asmpath) + ri.setKernel(isa, WAVESIZE_32) + + +class _StubWriter: + """Minimal writer: capability map for find() + defineSgpr/undefineSgpr sinks.""" + + def __init__(self, has_tdm=True, tdm_inst=3): + self.states = SimpleNamespace( + asmCaps={"HasTDM": has_tdm}, + archCaps={}, + kernel={"TDMInst": tdm_inst}, + ) + self.defined = [] + self.undefined = [] + + def defineSgpr(self, name, numSgprs, align=1): + self.defined.append((name, numSgprs)) + + def undefineSgpr(self, name): + from rocisa.code import ValueSet + self.undefined.append(name) + return ValueSet(name="sgpr" + name, value="UNDEF", format=-1) + + +def _kernel(*, multicast=True, clusterDim=(2, 2), tdmA=True, tdmB=True, + numWaves=4, useSubtile=False, sparse=0, tdmMeta=False, tdmInst=3): + return { + "Multicast": multicast, + "ClusterDim": list(clusterDim), + "enableTDMA": tdmA, + "enableTDMB": tdmB, + "enableTDMMetadata": tdmMeta, + "NumWaves": numWaves, + "UseSubtileImpl": useSubtile, + "TDMInst": tdmInst, + "ProblemType": {"Sparse": sparse}, + } + + +# --- selection ------------------------------------------------------------- + +class TestFind: + # Mirrors the production call sites, which resolve the component via the + # concrete ClusterLoadTDM.find(writer) (capability-gated selection). + def test_find_returns_tdm_impl_on_gfx1250(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + comp = ClusterLoadTDM.find(_StubWriter(has_tdm=True, tdm_inst=3)) + assert isinstance(comp, ClusterLoadTDM) + + def test_find_returns_none_without_tdm(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + assert ClusterLoadTDM.find(_StubWriter(has_tdm=False, tdm_inst=0)) is None + + def test_find_returns_none_when_tdm_inst_not_3(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + assert ClusterLoadTDM.find(_StubWriter(has_tdm=True, tdm_inst=0)) is None + + +# --- topology decision ----------------------------------------------------- + +class TestUsesCombinedMask: + def _c(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + return ClusterLoadTDM() + + def test_combined_when_both_tdm_multiwave_non_subtile(self): + assert self._c().usesCombinedMask(_kernel(tdmA=True, tdmB=True, numWaves=4, useSubtile=False)) + + def test_split_when_subtile(self): + assert not self._c().usesCombinedMask(_kernel(useSubtile=True)) + + def test_split_when_single_wave(self): + assert not self._c().usesCombinedMask(_kernel(numWaves=1)) + + def test_split_when_single_tensor(self): + assert not self._c().usesCombinedMask(_kernel(tdmA=True, tdmB=False)) + + +class TestMaskSgprName: + def _c(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + return ClusterLoadTDM() + + def test_wave_separated_combined_name(self): + k = _kernel() + assert self._c().maskSgprName(k, "A", waveSeparated=True) == "MulticastMask" + assert self._c().maskSgprName(k, "B", waveSeparated=True) == "MulticastMask" + + def test_dense_split_names(self): + k = _kernel() + assert self._c().maskSgprName(k, "A") == "MulticastMaskA" + assert self._c().maskSgprName(k, "B") == "MulticastMaskB" + + def test_dense_strips_mxs_prefix(self): + k = _kernel() + assert self._c().maskSgprName(k, "MXSA") == "MulticastMaskA" + assert self._c().maskSgprName(k, "MXSB") == "MulticastMaskB" + + def test_metadata_name(self): + k = _kernel() + assert self._c().maskSgprName(k, "Metadata") == "MulticastMaskMetadata" + + def test_subtile_split_name(self): + k = _kernel(useSubtile=True) + assert self._c().maskSgprName(k, "A", subtile=True) == "MulticastMaskA" + assert self._c().maskSgprName(k, "B", subtile=True) == "MulticastMaskB" + + +# --- SGPR declare / undeclare ---------------------------------------------- + +class TestDeclareUndeclare: + def _c(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + return ClusterLoadTDM() + + def test_declare_combined(self): + w = _StubWriter() + self._c().declareSgprs(w, _kernel()) # combined + assert [n for n, _ in w.defined] == ["MulticastMask"] + + def test_declare_split(self): + w = _StubWriter() + self._c().declareSgprs(w, _kernel(useSubtile=True)) # split + assert [n for n, _ in w.defined] == ["MulticastMaskA", "MulticastMaskB"] + + def test_declare_metadata(self): + w = _StubWriter() + self._c().declareSgprs(w, _kernel(useSubtile=True, sparse=1, tdmMeta=True)) + assert w.defined[-1] == ("MulticastMaskMetadata", 1) + + def test_declare_noop_when_multicast_off(self): + w = _StubWriter() + self._c().declareSgprs(w, _kernel(multicast=False)) + assert w.defined == [] + + def test_undeclare_combined(self): + _init_rocisa_gfx1250() + w = _StubWriter() + self._c().undeclareSgprs(w, _kernel()) + assert w.undefined == ["MulticastMask"] + + def test_undeclare_split(self): + _init_rocisa_gfx1250() + w = _StubWriter() + self._c().undeclareSgprs(w, _kernel(useSubtile=True)) + assert w.undefined == ["MulticastMaskA", "MulticastMaskB"] + + def test_undeclare_metadata(self): + # The metadata SGPR is freed alongside the split A/B masks on the sparse + # TDM path (enableTDMMetadata). + _init_rocisa_gfx1250() + w = _StubWriter() + self._c().undeclareSgprs(w, _kernel(useSubtile=True, sparse=1, tdmMeta=True)) + assert w.undefined == ["MulticastMaskA", "MulticastMaskB", "MulticastMaskMetadata"] + + def test_undeclare_noop_when_multicast_off(self): + _init_rocisa_gfx1250() + w = _StubWriter() + self._c().undeclareSgprs(w, _kernel(multicast=False)) + assert w.undefined == [] + + +# --- computeMasks emitted asm ---------------------------------------------- + +class TestComputeMasks: + def _c(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + return ClusterLoadTDM() + + def test_combined_parity_branch(self): + _init_rocisa_gfx1250() + # ClusterDim=[2,2] -> maskA = 1 | (1<<2) = 5 (0x5); maskB = (1<<2)-1 = 3 (0x3). + mod = self._c().computeMasks(_StubWriter(), _kernel(clusterDim=(2, 2)), + sgprWgX=61, sgprWgY=62, sgprNWgX=63, sTmp=60) + src = str(mod) + assert "Calculate multicast mask" in src + # Parity election on WaveIdx + even/odd label blocks. + assert "s_bitcmp1_b32 s[sgprWaveIdx], 0" in src + assert "setMulticastMask_OddWave" in src + assert "setMulticastMask_EvenWave" in src + # Combined mask target, both maskA (even) and maskB (odd) into MulticastMask. + assert "s_lshl_b32 s[sgprMulticastMask], 0x5, s61" in src + assert "s_lshl_b32 s[sgprMulticastMask], 0x3, s62" in src + # No split names on the combined path. + assert "MulticastMaskA" not in src + assert "MulticastMaskB" not in src + + def test_split_ab_branch(self): + _init_rocisa_gfx1250() + mod = self._c().computeMasks(_StubWriter(), _kernel(clusterDim=(2, 2), useSubtile=True), + sgprWgX=61, sgprWgY=62, sgprNWgX=63, sTmp=60) + src = str(mod) + assert "s_lshl_b32 s[sgprMulticastMaskA], 0x5, s61" in src + assert "s_lshl_b32 s[sgprMulticastMaskB], 0x3, s62" in src + # Split path has no wave-parity election. + assert "setMulticastMask_OddWave" not in src + + def test_noop_when_multicast_off(self): + _init_rocisa_gfx1250() + mod = self._c().computeMasks(_StubWriter(), _kernel(multicast=False), + sgprWgX=61, sgprWgY=62, sgprNWgX=63, sTmp=60) + assert str(mod).strip() == "" + + def test_metadata_mask_sparse_a(self): + # Sparse==1: the metadata mask follows sparse A -- shift maskA (0x5 for + # ClusterDim=[2,2]) by wg_x into MulticastMaskMetadata. + _init_rocisa_gfx1250() + mod = self._c().computeMasks( + _StubWriter(), _kernel(clusterDim=(2, 2), sparse=1, tdmMeta=True), + sgprWgX=61, sgprWgY=62, sgprNWgX=63, sTmp=60) + src = str(mod) + assert "Setting metadata mask (follows sparse A)" in src + assert "s_lshl_b32 s[sgprMulticastMaskMetadata], 0x5, s61" in src + + def test_metadata_mask_sparse_b(self): + # Sparse==2: the metadata mask follows sparse B -- shift maskB (0x3) by + # (wg_y * nwg_x) computed into the sTmp+4 scratch slot. + _init_rocisa_gfx1250() + mod = self._c().computeMasks( + _StubWriter(), _kernel(clusterDim=(2, 2), sparse=2, tdmMeta=True), + sgprWgX=61, sgprWgY=62, sgprNWgX=63, sTmp=60) + src = str(mod) + assert "Shift factor: wg_y * nwg_x (metadata)" in src + assert "Setting metadata mask (follows sparse B)" in src + assert "s_lshl_b32 s[sgprMulticastMaskMetadata], 0x3, s64" in src + + +# --- applyToDescriptor emitted asm ----------------------------------------- + +class TestApplyToDescriptor: + def _c(self): + from Tensile.Components.ClusterLoad import ClusterLoadTDM + return ClusterLoadTDM() + + def test_dense_split_or(self): + _init_rocisa_gfx1250() + w = _StubWriter() + mod = self._c().applyToDescriptor(w, _kernel(), "tdmAGroup1", "A") + assert "s_or_b32 s[sgprtdmAGroup1], s[sgprtdmAGroup1], s[sgprMulticastMaskA]" in str(mod) + + def test_wave_separated_combined_or(self): + _init_rocisa_gfx1250() + w = _StubWriter() + mod = self._c().applyToDescriptor(w, _kernel(), "tdmAGroup1", "A", waveSeparated=True) + assert "s_or_b32 s[sgprtdmAGroup1], s[sgprtdmAGroup1], s[sgprMulticastMask]" in str(mod) + + def test_subtile_split_or(self): + _init_rocisa_gfx1250() + w = _StubWriter() + mod = self._c().applyToDescriptor(w, _kernel(useSubtile=True), "tdmBGroup1", "B", subtile=True) + assert "s_or_b32 s[sgprtdmBGroup1], s[sgprtdmBGroup1], s[sgprMulticastMaskB]" in str(mod) + + def test_empty_when_multicast_off(self): + _init_rocisa_gfx1250() + w = _StubWriter() + mod = self._c().applyToDescriptor(w, _kernel(multicast=False), "tdmAGroup1", "A") + assert str(mod).strip() == "" + + def test_empty_when_cluster_disabled(self): + _init_rocisa_gfx1250() + w = _StubWriter() + mod = self._c().applyToDescriptor(w, _kernel(clusterDim=(1, 1)), "tdmAGroup1", "A") + assert str(mod).strip() == "" + + +if __name__ == "__main__": + sys.exit(pytest.main([__file__, "-v"])) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_multicast_legacy_coercion.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_multicast_legacy_coercion.py new file mode 100644 index 000000000000..431b166163c0 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_multicast_legacy_coercion.py @@ -0,0 +1,229 @@ +#!/usr/bin/env python3 +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +# Regression tests for the legacy bool -> int Multicast coercion. +# +# Multicast became an int solution parameter (valid values [-1, 0, 1]: -1 auto, +# 0 off, 1 on). But 600+ shipped library-logic artifacts still serialize the +# pre-existing derived value as `Multicast: false` (BOOL). Because msgpack +# serializes bool and int as different wire types, a bool value trips the strict +# create-library type gate (raiseIfTypeMismatches) -> ConfigTypeError, and would +# cause std::bad_cast at runtime. +# +# Fix: on the serialized-state loading path (raiseProblemTypeOnTypeMismatch= +# False), coerce a bool Multicast to the tri-state int (false -> 0 "off", +# true -> 1 "on"; never -1 "auto"), so both the type gate and the emitted +# msgpack see an int. The config/derivation path now emits an int Multicast on +# all branches directly (no post-derivation coercion needed). +# +# Usage: +# pytest test_multicast_legacy_coercion.py -v +################################################################################ + +import copy +import importlib + +import pytest + +pytestmark = pytest.mark.unit + +from Tensile.SolutionStructs.Solution import ( + coerceLegacyMulticastType, + validateParameterTypes, +) + +S = importlib.import_module("Tensile.SolutionStructs.Solution") +Solution = S.Solution + + +# --- pure coercion helper -------------------------------------------------- + +class TestCoerceHelper: + def test_false_maps_to_int_zero_off(self): + state = {"Multicast": False} + coerceLegacyMulticastType(state) + assert state["Multicast"] == 0 + assert type(state["Multicast"]) is int + # Must be "off" (0), never "auto" (-1). + assert state["Multicast"] != -1 + + def test_true_maps_to_int_one_on(self): + state = {"Multicast": True} + coerceLegacyMulticastType(state) + assert state["Multicast"] == 1 + assert type(state["Multicast"]) is int + + def test_int_values_untouched(self): + for v in (-1, 0, 1): + state = {"Multicast": v} + coerceLegacyMulticastType(state) + assert state["Multicast"] == v + assert type(state["Multicast"]) is int + + def test_absent_key_is_noop(self): + state = {"SomethingElse": 5} + coerceLegacyMulticastType(state) + assert "Multicast" not in state + + +# --- coercion clears the strict type gate ---------------------------------- + +class TestGateIntegration: + def test_bool_multicast_flagged_before_coercion(self): + """A bool Multicast is what the create-library gate flags as + 'found bool ... expected int'.""" + records = validateParameterTypes({"Multicast": False}) + assert records, "expected a type-mismatch record for bool Multicast" + (name, actual, expected), _value, _src = records[0] + assert name == "Multicast" + assert actual == "bool" + assert "int" in expected + + def test_no_mismatch_after_coercion(self): + """After coercion the same state passes validateParameterTypes clean, + i.e. raiseIfTypeMismatches() would not fire.""" + for legacy in (False, True): + state = {"Multicast": legacy} + coerceLegacyMulticastType(state) + assert validateParameterTypes(state) == [] + + +# --- end-to-end through the real Solution constructor ---------------------- +# +# Mirrors the create-library / library-logic loading path +# (raiseProblemTypeOnTypeMismatch=False), which is where the CI build broke. +# Requires amdclang++ for the ISA-info map (same as the other Solution +# characterization tests); skipped gracefully if the toolchain is absent. + +@pytest.fixture(scope="module") +def _toolchain(): + from Tensile.Common.Architectures import gfxToIsa + from Tensile.Common.Capabilities import makeIsaInfoMap + from Tensile.Toolchain.Assembly import makeAssemblyToolchain + from Tensile.Toolchain.Validators import validateToolchain, ToolchainDefaults + + try: + cxx = validateToolchain("amdclang++") + bundler = validateToolchain(ToolchainDefaults.OFFLOAD_BUNDLER) + except Exception as e: # pragma: no cover - environment dependent + pytest.skip(f"amdclang++ toolchain unavailable: {e}") + isa = gfxToIsa("gfx950") + iim = makeIsaInfoMap([isa], cxx) + assembler = makeAssemblyToolchain(cxx, bundler, "default").assembler + return iim, assembler + + +@pytest.fixture(scope="module") +def _gp_assigned(_toolchain): + from Tensile.Common.GlobalParameters import globalParameters, assignGlobalParameters + from Tensile.Common.ValidParameters import validParameters + + iim, _ = _toolchain + saved_gp = copy.deepcopy(dict(globalParameters)) + saved_vp = copy.deepcopy(dict(validParameters)) + assignGlobalParameters({}, iim) + yield + globalParameters.clear() + globalParameters.update(saved_gp) + validParameters.clear() + validParameters.update(saved_vp) + + +def _make_solution(iim, assembler, multicast, *, loadingSerializedState): + """Build & derive a minimal gfx950 MX FP8 solution carrying a bool Multicast. + + ``loadingSerializedState=True`` drives the library-logic/artifact path + (raiseProblemTypeOnTypeMismatch=False); False drives the config path. + """ + from Tensile.BenchmarkProblems import matrixInstructionToMIParameters + + isa = list(iim.keys())[0] + params = { + "ProblemType": { + "OperationType": "GEMM", + "DataType": "F8", + "DestDataType": "s", + "ComputeDataType": "s", + "HighPrecisionAccumulate": True, + "MXBlockA": 32, + "MXBlockB": 32, + "TransposeA": True, + "TransposeB": False, + "UseBeta": True, + "Batched": True, + }, + "ISA": isa, + "MatrixInstruction": [16, 16, 128, 1, 1, 2, 2, 2, 2], + "WorkGroup": [16, 16, 1], + "WavefrontSize": 64, + "DepthU": 256, + "KernelLanguage": "Assembly", + "PrefetchGlobalRead": 1, + "PrefetchLocalRead": 1, + "ScheduleIterAlg": 3, + "DirectToLds": 1, + "StaggerU": 0, + "StreamK": 3, + "UseSubtileImpl": True, + "LocalReadVectorWidth": -1, + "GlobalReadVectorWidthA": 16, + "GlobalReadVectorWidthB": 16, + "TransposeLDS": 1, + "LdsPadA": -1, + "LdsPadB": -1, + "LdsBlockSizePerPadA": -1, + "LdsBlockSizePerPadB": -1, + "1LDSBuffer": -1, + "VectorWidthA": -1, + "VectorWidthB": -1, + "StoreVectorWidth": -1, + "SourceSwap": True, + "ExpandPointerSwap": True, + "GlobalSplitU": 1, + "InnerUnroll": 1, + "DebugStreamK": 0, + # The legacy artifact wire type under test: a serialized bool. + "Multicast": multicast, + } + mi = params["MatrixInstruction"] + params.update(matrixInstructionToMIParameters( + mi, isa, params["WavefrontSize"], params["ProblemType"], params["WorkGroup"], iim)) + + return Solution( + params, False, False, False, assembler, iim, + raiseProblemTypeOnTypeMismatch=not loadingSerializedState, + ) + + +class TestSolutionConstructorPath: + @pytest.mark.parametrize("legacy,expected_int", [(False, 0), (True, 1)]) + def test_library_logic_path_coerces_to_int( + self, _toolchain, _gp_assigned, legacy, expected_int + ): + """On the serialized-state path a bool Multicast is coerced to the + tri-state int, and the state passes the strict type gate clean.""" + iim, assembler = _toolchain + sol = _make_solution(iim, assembler, legacy, loadingSerializedState=True) + mc = sol._state["Multicast"] + assert type(mc) is int, f"Multicast must be int on the load path, got {type(mc)}" + assert mc == expected_int + assert mc != -1 # coerced off/on, never auto + # The gate the create-library path runs must see no Multicast mismatch. + records = validateParameterTypes(sol._state) + assert all(rec[0][0] != "Multicast" for rec in records), records + + def test_config_path_derives_int(self, _toolchain, _gp_assigned): + """The config/derivation path now emits an int Multicast on all branches + (the post-derivation coercion was removed as redundant): a Multicast=False + input derives to int 0 (the mc==0 force-off branch), not a bool.""" + iim, assembler = _toolchain + sol = _make_solution(iim, assembler, False, loadingSerializedState=False) + mc = sol._state["Multicast"] + assert type(mc) is int, f"config path must derive int, got {type(mc)}" + assert mc == 0, mc + + +if __name__ == "__main__": + import sys + sys.exit(pytest.main([__file__, "-v"])) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_multicast_tristate.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_multicast_tristate.py new file mode 100644 index 000000000000..c28c516f48d1 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_multicast_tristate.py @@ -0,0 +1,126 @@ +#!/usr/bin/env python3 +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +# Unit tests for the decoupled tri-state Multicast solution parameter. +# +# Multicast used to be a derived-only state var, unconditionally forced on for +# ClusterDim != [1,1] (except Stream-K). It is now an explicit +# tri-state control: +# -1 = auto (legacy coupling), 0 = force off, 1 = force on. +# Default -1 reproduces the historic derivation exactly. These tests pin: +# * registration (valid values + default), and +# * the derivation semantics (-1 legacy == old behavior; 0 off; 1 on), +# driven through the real config -> Solution derivation path. +# +# Usage: +# pytest test_multicast_tristate.py -v +################################################################################ + +import copy +import os +import sys + +import pytest + +pytestmark = pytest.mark.unit + +SCRIPT_DIR = os.path.dirname(os.path.abspath(__file__)) +TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", "..")) +sys.path.insert(0, TENSILE_ROOT) +sys.path.insert(0, os.path.join( + TENSILE_ROOT, "Tensile", "Tests", "unit", "characterization", "_codegen")) + +_DESIGNED = os.path.join( + TENSILE_ROOT, "Tensile", "Tests", "unit", "characterization", + "_codegen", "data", "test_data", "_designed", "gfx1250") +_XCCREMAP = os.path.join(_DESIGNED, "xccremap.yaml") +_STREAMK_CLUSTER = os.path.join(_DESIGNED, "streamk_cluster_coop_load.yaml") + + +# --- registration ---------------------------------------------------------- + +class TestRegistration: + def test_valid_values(self): + from Tensile.Common.ValidParameters import validParameters + assert validParameters["Multicast"] == [-1, 0, 1] + + def test_default_is_legacy_auto(self): + from Tensile.Common.GlobalParameters import defaultSolution + assert defaultSolution["Multicast"] == -1 + + +# --- derivation (through real config -> Solution) -------------------------- + +def _write_variant(tmp_path, base_yaml, name, *, multicast=None): + """Copy a designed config, optionally injecting a Multicast fork value.""" + from Tensile import LibraryIO + import yaml + + cfg = copy.deepcopy(LibraryIO.read(base_yaml)) + if multicast is not None: + fork = cfg["BenchmarkProblems"][0][1]["ForkParameters"] + fork.append({"Multicast": [multicast]}) + out = tmp_path / name + with open(out, "w") as f: + yaml.safe_dump(cfg, f, default_flow_style=None) + return str(out) + + +def _derive_states(cfg_path): + from config_harness import solutions_from_config + sols = solutions_from_config(cfg_path, arch="gfx1250", limit_solutions=8) + states = [] + for s in sols: + st = s._state if hasattr(s, "_state") else s + states.append(st) + return states + + +class TestDerivation: + def test_legacy_auto_cluster_on(self, tmp_path): + # -1 (omitted) + ClusterDim=[2,2], no StreamK reduction -> Multicast on. + cfg = _write_variant(tmp_path, _XCCREMAP, "legacy.yaml") + states = _derive_states(cfg) + assert states, "expected >=1 derived solution" + assert all(st["Multicast"] == 1 for st in states), ( + [st["Multicast"] for st in states]) + + def test_explicit_off(self, tmp_path): + # Multicast=0 forces off even with ClusterDim=[2,2]. + cfg = _write_variant(tmp_path, _XCCREMAP, "off.yaml", multicast=0) + states = _derive_states(cfg) + assert states, "expected >=1 derived solution" + assert all(st["Multicast"] == 0 for st in states), ( + [st["Multicast"] for st in states]) + # ClusterBarrier is gated on Multicast, so it must also be off. + assert all(st["ClusterBarrier"] is False for st in states) + + def test_explicit_on(self, tmp_path): + # Multicast=1 forces on. + cfg = _write_variant(tmp_path, _XCCREMAP, "on.yaml", multicast=1) + states = _derive_states(cfg) + assert states, "expected >=1 derived solution" + assert all(st["Multicast"] == 1 for st in states), ( + [st["Multicast"] for st in states]) + + def test_streamk_cluster_auto_multicast(self, tmp_path): + # Collapse: -1 (omitted Multicast) + StreamK=3 + ClusterDim != [1,1] now + # AUTO-ENABLES the cooperative-load path. StreamKMulticast is derived to + # 1 and Multicast to True -- the bare index-only StreamK cluster state no + # longer exists. + cfg = _write_variant(tmp_path, _STREAMK_CLUSTER, "sk_auto_mc.yaml") + states = _derive_states(cfg) + assert states, "expected >=1 derived solution" + assert all(st["StreamKMulticast"] == 1 for st in states), ( + [st.get("StreamKMulticast") for st in states]) + assert all(st["Multicast"] == 1 for st in states), ( + [st["Multicast"] for st in states]) + # The cooperative multicast now pairs the masks with the cluster-scope + # barrier handshake, so ClusterBarrier is derived on. + assert all(st["ClusterBarrier"] is True for st in states), ( + [st["ClusterBarrier"] for st in states]) + + +if __name__ == "__main__": + sys.exit(pytest.main([__file__, "-v"])) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_cluster_sk45_reject.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_cluster_sk45_reject.py new file mode 100644 index 000000000000..72e98d5040d7 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_cluster_sk45_reject.py @@ -0,0 +1,96 @@ +#!/usr/bin/env python3 +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +# Unit test: StreamK dynamic (SK4) / hybrid (SK5) modes reject ClusterDim. +# +# WG-cluster support is StreamK==3-only (the [C,1] barrier reduction and the DP +# cooperative B-multicast are both SK3 features). There is no cluster-load / +# reduction implementation for the dynamic (SK4) or hybrid (SK5) work-queue +# modes, so Solution.assignDerivedParameters rejects StreamK in {4,5} with +# ClusterDim != [1,1] outright (rather than emitting an unusable cluster kernel +# whose decoded cluster WG-id no feature consumes). +# +# These tests drive the real config -> Solution derivation path and assert: +# * SK4/SK5 + ClusterDim != [1,1] -> 0 derived solutions (the reject fires); and +# * the SAME config with ClusterDim = [1,1] still derives solutions, proving +# the differentiator is the cluster (not some unrelated SK4/SK5 reject). +# +# Usage: +# pytest test_streamk_cluster_sk45_reject.py -v +################################################################################ + +import copy +import os +import sys + +import pytest + +pytestmark = pytest.mark.unit + +SCRIPT_DIR = os.path.dirname(os.path.abspath(__file__)) +TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", "..")) +sys.path.insert(0, TENSILE_ROOT) +sys.path.insert(0, os.path.join( + TENSILE_ROOT, "Tensile", "Tests", "unit", "characterization", "_codegen")) + +_DESIGNED = os.path.join( + TENSILE_ROOT, "Tensile", "Tests", "unit", "characterization", + "_codegen", "data", "test_data", "_designed", "gfx1250") +# A known-good gfx1250 StreamK=3 + ClusterDim config; we only re-fork StreamK +# (and, for the control, ClusterDim) on top of it. +_BASE = os.path.join(_DESIGNED, "streamk_cluster_coop_load.yaml") + +_ARCH = "gfx1250" + + +def _write_variant(tmp_path, name, overrides): + """Copy _BASE, replacing/appending the given fork parameter values.""" + from Tensile import LibraryIO + import yaml + + cfg = copy.deepcopy(LibraryIO.read(_BASE)) + fork = cfg["BenchmarkProblems"][0][1]["ForkParameters"] + for key, val in overrides.items(): + replaced = False + for entry in fork: + if key in entry: + entry[key] = val + replaced = True + break + if not replaced: + fork.append({key: val}) + out = tmp_path / name + with open(out, "w") as f: + yaml.safe_dump(cfg, f, default_flow_style=None) + return str(out) + + +def _derive_states(cfg_path): + from config_harness import solutions_from_config + sols = solutions_from_config(cfg_path, arch=_ARCH, limit_solutions=8) + return [s._state if hasattr(s, "_state") else s for s in sols] + + +@pytest.mark.parametrize("sk", [4, 5]) +def test_sk45_cluster_rejected(tmp_path, sk): + """StreamK dynamic/hybrid + ClusterDim != [1,1] derives no solutions.""" + cfg = _write_variant(tmp_path, f"sk{sk}_cluster.yaml", {"StreamK": [sk]}) + assert _derive_states(cfg) == [], ( + f"StreamK={sk} with ClusterDim != [1,1] must be rejected " + "(cluster support is SK3-only)") + + +@pytest.mark.parametrize("sk", [4, 5]) +def test_sk45_without_cluster_still_valid(tmp_path, sk): + """Control: the same config with ClusterDim=[1,1] still derives solutions, + so the reject above is caused by the cluster, not an unrelated SK4/SK5 + constraint.""" + cfg = _write_variant(tmp_path, f"sk{sk}_nocluster.yaml", + {"StreamK": [sk], "ClusterDim": [[1, 1]]}) + assert _derive_states(cfg), ( + f"StreamK={sk} without ClusterDim should still derive solutions") + + +if __name__ == "__main__": + sys.exit(pytest.main([__file__, "-v"])) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_multicast.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_multicast.py new file mode 100644 index 000000000000..5a1ef212f02f --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_multicast.py @@ -0,0 +1,338 @@ +#!/usr/bin/env python3 +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT +################################################################################ +# Unit tests for the gfx1250 StreamK DP cooperative B-multicast fast path +# (StreamKMulticast). +# +# StreamKMulticast co-locates C consecutive StreamK DP workgroups in a 1-D +# workgroup cluster (ClusterDim = [C, 1]); in the DP region those C WGs process +# M-adjacent tiles that share the same B (N-block) over full K, so B is loaded +# once and TDM-multicast to the whole cluster while A stays per-WG. +# +# These tests pin (CPU-only, no GPU): +# * internalization (StreamKMulticast is derived-only: absent from +# validParameters/defaultSolution, never user-settable); +# * the validation matrix (accepted only for SK3 + ClusterDim=[C,1] pow2 2..16 +# + gfx1250 HasTDM/TDMInst + XCC=0 + not atomic; rejected else); +# * the emitted asm: DP loads carry the split B-broadcast mask +# (MulticastMaskB OR'd into the B descriptor Group1), A carries the self-only +# mask, the runtime clusterMulticastValid predicate is present, and the +# DP->SK boundary clear drops the B broadcast for the SK region. +# +# Usage: +# pytest test_streamk_multicast.py -v +################################################################################ + +import copy +import os +import sys + +import pytest + +pytestmark = pytest.mark.unit + +SCRIPT_DIR = os.path.dirname(os.path.abspath(__file__)) +TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", "..")) +sys.path.insert(0, TENSILE_ROOT) +sys.path.insert(0, os.path.join( + TENSILE_ROOT, "Tensile", "Tests", "unit", "characterization", "_codegen")) + +_DESIGNED = os.path.join( + TENSILE_ROOT, "Tensile", "Tests", "unit", "characterization", + "_codegen", "data", "test_data", "_designed", "gfx1250") +_STREAMK_MULTICAST = os.path.join(_DESIGNED, "streamk_cluster_multicast.yaml") +_STREAMK_CLUSTER_BARE = os.path.join(_DESIGNED, "streamk_cluster_coop_load.yaml") + +_ARCH = "gfx1250" + + +# --- registration ---------------------------------------------------------- + +class TestRegistration: + def test_not_a_valid_parameter(self): + """StreamKMulticast is derived-only (ClusterBarrier precedent): it must + NOT be a user/benchmark-settable validParameter.""" + from Tensile.Common.ValidParameters import validParameters + assert "StreamKMulticast" not in validParameters + + def test_not_a_default_benchmark_parameter(self): + """Not in defaultSolution either -- it is seeded/derived on state only by + Solution.assignProblemIndependentDerivedParameters.""" + from Tensile.Common.GlobalParameters import defaultSolution + assert "StreamKMulticast" not in defaultSolution + + +# --- config -> Solution derivation helpers --------------------------------- + +def _write_variant(tmp_path, name, *, fork_overrides=None): + """Copy the designed multicast config, overriding fork param values. + + ``fork_overrides`` maps a fork parameter name to its single-element value + list; an existing fork entry is replaced, otherwise appended. + """ + from Tensile import LibraryIO + import yaml + + cfg = copy.deepcopy(LibraryIO.read(_STREAMK_MULTICAST)) + if fork_overrides: + fork = cfg["BenchmarkProblems"][0][1]["ForkParameters"] + for key, val in fork_overrides.items(): + replaced = False + for entry in fork: + if key in entry: + entry[key] = val + replaced = True + break + if not replaced: + fork.append({key: val}) + out = tmp_path / name + with open(out, "w") as f: + yaml.safe_dump(cfg, f, default_flow_style=None) + return str(out) + + +def _derive_states(cfg_path): + from config_harness import solutions_from_config + sols = solutions_from_config(cfg_path, arch=_ARCH, limit_solutions=8) + states = [] + for s in sols: + st = s._state if hasattr(s, "_state") else s + states.append(st) + return states + + +# --- validation matrix ----------------------------------------------------- + +class TestValidation: + def test_accepted_baseline(self, tmp_path): + """The designed SK3 cluster config (ClusterDim=[4,1]) derives valid + solutions with the internal StreamKMulticast auto-derived to 1 and + Multicast on.""" + cfg = _write_variant(tmp_path, "ok.yaml") + states = _derive_states(cfg) + assert states, "expected >=1 derived solution for the valid config" + for st in states: + assert st["StreamKMulticast"] == 1 + assert st["Multicast"] == 1, st["Multicast"] + assert st["ClusterDim"] == [4, 1] + # The cooperative multicast pairs the B-broadcast masks with the + # cluster-scope barrier handshake, so ClusterBarrier is derived on. + assert st["ClusterBarrier"] is True, st.get("ClusterBarrier") + + # NB: test_auto_enable_from_bare_cluster was removed as a strict subset of + # test_multicast_tristate.py::TestDerivation::test_streamk_cluster_auto_multicast, + # which derives the same bare SK3 + ClusterDim config and asserts the same + # StreamKMulticast==1 / Multicast==1 (plus ClusterBarrier is True). + + def test_reject_multicast_force_off(self, tmp_path): + """StreamKMulticast auto-enabled by ClusterDim on SK3 is incompatible with + an explicit Multicast=0 (force off): the mask SGPRs are gated on Multicast + while the predicate/boundary-clear emitters are gated on StreamKMulticast, + so Multicast=0 would reference undeclared MulticastMaskA/B. Reject.""" + cfg = _write_variant(tmp_path, "mc_off.yaml", + fork_overrides={"Multicast": [0]}) + assert _derive_states(cfg) == [] + + # NB: test_control_multicast_auto_enabled was removed as a subset of + # test_accepted_baseline above, which derives the same SK3 [4,1] cluster + # config with the default Multicast=-1 (auto) and already asserts + # StreamKMulticast==1 / Multicast==1 (plus ClusterDim/ClusterBarrier). The + # negative-path partner test_reject_multicast_force_off is retained. + + def test_reject_atomic(self, tmp_path): + cfg = _write_variant(tmp_path, "atomic.yaml", + fork_overrides={"StreamKAtomic": [1]}) + assert _derive_states(cfg) == [] + + def test_accept_pgr2(self, tmp_path): + """The DP cooperative multicast path supports double-buffered global + prefetch (PrefetchGlobalRead > 1): the prologue double-buffer prefetch + multicast load is bracketed by a cluster-scope handshake in codegen, so + both PrefetchGlobalRead 1 and 2 are accepted.""" + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + + class _Info: + asmCaps = {"HasTDM": True, "HasClusterBarrier": True} + isaInfoMap = {(12, 5, 0): _Info()} + + def _state(pgr): + return { + "StreamKMulticast": 1, + "Multicast": 1, + "StreamK": 3, + "StreamKAtomic": 0, + "StreamKXCCMapping": 0, + "ClusterDim": [4, 1], + "ISA": [12, 5, 0], + "TDMInst": 3, + "PrefetchGlobalRead": pgr, + } + + assert _validateStreamKMulticast(_state(2), False, isaInfoMap) is True + assert _validateStreamKMulticast(_state(1), False, isaInfoMap) is True + + cfg = _write_variant(tmp_path, "pgr2.yaml", + fork_overrides={"PrefetchGlobalRead": [2]}) + states = _derive_states(cfg) + assert states, "expected the PrefetchGlobalRead=2 multicast config to be accepted" + for st in states: + assert st["StreamKMulticast"] == 1 + + def test_xcc_mapping_forced_to_zero(self, tmp_path): + """StreamKXCCMapping is coerced to 0 (not rejected) under StreamK+ClusterDim. + + The general Stream-K + ClusterDim reconciliation force-sets + StreamKXCCMapping = 0 (the WGM/XCC WorkGroup0 remap has no cluster + awareness) *before* _validateStreamKMulticast runs. That coerced value is + exactly what StreamKMulticast requires (XCC == 0), so the solution is + accepted with the remap disabled rather than rejected. Our + _validateStreamKMulticast XCC check remains as redundant safety.""" + cfg = _write_variant(tmp_path, "xcc.yaml", + fork_overrides={"StreamKXCCMapping": [3]}) + states = _derive_states(cfg) + assert states, "expected the XCC=3 config to be accepted with XCC coerced to 0" + for st in states: + assert st["StreamKMulticast"] == 1 + assert st["StreamKXCCMapping"] == 0, st["StreamKXCCMapping"] + + def test_reject_non_1d_cluster(self, tmp_path): + # ClusterDim = [2, 2] is not the [C, 1] spatial DP cluster. + cfg = _write_variant(tmp_path, "cd22.yaml", + fork_overrides={"ClusterDim": [[2, 2]]}) + assert _derive_states(cfg) == [] + + def test_reject_non_pow2_cluster(self, tmp_path): + cfg = _write_variant(tmp_path, "cd3.yaml", + fork_overrides={"ClusterDim": [[3, 1]]}) + assert _derive_states(cfg) == [] + + # NB: C > 16 is not an expressible ClusterDim (validParameters caps + # ClusterDim x at 16), so the "> 16" branch of the validator is defensive + # and unreachable through valid params -- no test drives it here. + + # --- direct _validateStreamKMulticast reject branches ------------------ + # Several reject branches are unreachable through the config-derivation path + # (the collapse only auto-derives StreamKMulticast for SK3 and force-coerces + # StreamKXCCMapping=0, and the designed configs are always gfx1250 with full + # caps), so drive them directly with a hand-built state -- the same pattern + # test_accept_pgr2 uses. + @staticmethod + def _direct_state(**overrides): + st = { + "StreamKMulticast": 1, "Multicast": 1, "StreamK": 3, + "StreamKAtomic": 0, "StreamKXCCMapping": 0, "ClusterDim": [4, 1], + "ISA": [12, 5, 0], "TDMInst": 3, "PrefetchGlobalRead": 1, + } + st.update(overrides) + return st + + @staticmethod + def _isa_map(has_tdm=True, has_cluster_barrier=True): + class _Info: + asmCaps = {"HasTDM": has_tdm, "HasClusterBarrier": has_cluster_barrier} + return {(12, 5, 0): _Info()} + + def test_reject_streamk_not_3_direct(self): + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + assert _validateStreamKMulticast( + self._direct_state(StreamK=4), False, self._isa_map()) is False + + def test_reject_xcc_mapping_direct(self): + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + assert _validateStreamKMulticast( + self._direct_state(StreamKXCCMapping=3), False, self._isa_map()) is False + + def test_reject_non_gfx1250_isa(self): + # The ISA gate rejects before indexing isaInfoMap, so a foreign ISA need + # not be present in the map. + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + assert _validateStreamKMulticast( + self._direct_state(ISA=[9, 4, 2]), False, self._isa_map()) is False + + def test_reject_missing_hastdm(self): + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + assert _validateStreamKMulticast( + self._direct_state(), False, self._isa_map(has_tdm=False)) is False + + def test_reject_missing_hasclusterbarrier(self): + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + assert _validateStreamKMulticast( + self._direct_state(), False, self._isa_map(has_cluster_barrier=False)) is False + + +class TestTDMInstValidation: + """The tightened TDMInst check: StreamKMulticast requires TDMInst == 3 (the + only TDMInst a ClusterLoadTDM component matches), so TDMInst in {1,2} is + rejected even on gfx1250 HasTDM -- otherwise the masks would silently drop.""" + + @staticmethod + def _state(tdminst, pgr=1): + return { + "StreamKMulticast": 1, + # StreamKMulticast on always co-derives Multicast on (real invariant). + "Multicast": 1, + "StreamK": 3, + "StreamKAtomic": 0, + "StreamKXCCMapping": 0, + "ClusterDim": [4, 1], + "ISA": [12, 5, 0], + "TDMInst": tdminst, + "PrefetchGlobalRead": pgr, + } + + @staticmethod + def _isa_map(has_tdm=True): + class _Info: + asmCaps = {"HasTDM": has_tdm, "HasClusterBarrier": True} + return {(12, 5, 0): _Info()} + + @pytest.mark.parametrize("tdminst", [1, 2]) + def test_reject_non_tdm3(self, tdminst): + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + st = self._state(tdminst) + assert _validateStreamKMulticast(st, False, self._isa_map()) is False + assert st.get("Valid") is False + + def test_accept_tdm3(self): + from Tensile.SolutionStructs.Solution import _validateStreamKMulticast + st = self._state(3) + assert _validateStreamKMulticast(st, False, self._isa_map()) is True + + +# --- emitted assembly ------------------------------------------------------ + +class TestEmit: + # NB: the byte-exact emitted assembly is pinned by the characterization + # golden test_streamk_cluster_multicast_gfx1250_char.py (+ its .ambr / + # cmpasm snapshot). The former positive string assertions here (emits + # assembly, split-mask bindings, clusterMulticastValid predicate, cluster + # barrier handshake, DP->SK boundary clear) duplicated that golden and were + # removed. Only the two checks the golden does NOT express as a single + # readable assertion are kept: the C=4 mask arithmetic and the negative + # combined-mask-leak scan. + def _emit(self, cfg=_STREAMK_MULTICAST): + from config_harness import emit_kernels_from_config + return emit_kernels_from_config(cfg, limit=8, arch=_ARCH) + + def test_broadcast_mask_value(self): + """maskB = (1< 1) + { + size_t c = sizeMapping.clusterDim.x; + sk.grid = ((sk.grid + c - 1) / c) * c; + } } GSUSettings gsuSettings; @@ -4064,6 +4078,21 @@ namespace TensileLite } } + // StreamKMulticast (gfx1250): DP cooperative B-multicast. The multicast + // cluster is SPATIAL: the C peers of a cluster process C DISTINCT, + // M-adjacent tiles and share the B (N-block) tile. v1 is single-round, so + // the launched grid is `tiles` rounded UP to a multiple of C = clusterDim.x. + // Rounding up keeps every launched HW cluster full while giving each WG one + // tile in the single DP round; a trailing partial cluster (tiles % C != 0) + // has idle-but-present tail WGs whose cluster is disabled by the kernel's + // clusterMulticastValid runtime predicate (so a masked target is never left + // without a matching load). See docs/design/cluster-load-component-and-streamk-multicast.md. + if(self.sizeMapping.streamKMulticast && self.sizeMapping.clusterDim.x > 1) + { + size_t c = self.sizeMapping.clusterDim.x; + skGrid = ((tiles + c - 1) / c) * c; + } + return skGrid; } } // namespace diff --git a/shared/stinkytofu/include/stinkytofu/bindings/python/Module.hpp b/shared/stinkytofu/include/stinkytofu/bindings/python/Module.hpp index 74fe7d2d6f08..b453b734130e 100644 --- a/shared/stinkytofu/include/stinkytofu/bindings/python/Module.hpp +++ b/shared/stinkytofu/include/stinkytofu/bindings/python/Module.hpp @@ -75,6 +75,7 @@ X(EnableSwPrefetchInsertion, bool) \ X(SwPrefetchScratchSgpr, int) \ X(ClusterBarrier, bool) \ + X(StreamKMulticast, bool) \ X(PrefetchGlobalRead, int) \ X(PrefetchLocalRead, int) \ X(RemoveInstructions, std::string) \ diff --git a/shared/stinkytofu/include/stinkytofu/transforms/asm/InsertClusterBarrierPass.hpp b/shared/stinkytofu/include/stinkytofu/transforms/asm/InsertClusterBarrierPass.hpp index 4d1ad2e1f1db..3c66ee18a7b3 100644 --- a/shared/stinkytofu/include/stinkytofu/transforms/asm/InsertClusterBarrierPass.hpp +++ b/shared/stinkytofu/include/stinkytofu/transforms/asm/InsertClusterBarrierPass.hpp @@ -29,8 +29,7 @@ namespace stinkytofu { class Pass; -STINKYTOFU_EXPORT std::unique_ptr createInsertClusterBarrierPass(bool isKernelScope = true, - int pgrValue = 1, - int plrValue = 1); +STINKYTOFU_EXPORT std::unique_ptr createInsertClusterBarrierPass( + bool isKernelScope = true, int pgrValue = 1, int plrValue = 1, bool streamKMulticast = false); } // namespace stinkytofu diff --git a/shared/stinkytofu/src/pipeline/backend/Gfx1250Backend.cpp b/shared/stinkytofu/src/pipeline/backend/Gfx1250Backend.cpp index 79dc373392c8..bed154bc13b2 100644 --- a/shared/stinkytofu/src/pipeline/backend/Gfx1250Backend.cpp +++ b/shared/stinkytofu/src/pipeline/backend/Gfx1250Backend.cpp @@ -156,9 +156,11 @@ bool buildGfx1250Pipeline(PassManager& pm, StinkyAsmModule& module, const PassBu // the module opts in. Must precede InsertVgprMsbPass so the new // branches/labels are present when MSB configuration is materialized. if (moduleOptions.ClusterBarrier) { - pm.addPass(createInsertClusterBarrierPass(/*isKernelScope=*/true, - /*pgrValue=*/moduleOptions.PrefetchGlobalRead, - /*plrValue=*/moduleOptions.PrefetchLocalRead)); + pm.addPass(createInsertClusterBarrierPass( + /*isKernelScope=*/true, + /*pgrValue=*/moduleOptions.PrefetchGlobalRead, + /*plrValue=*/moduleOptions.PrefetchLocalRead, + /*streamKMulticast=*/moduleOptions.StreamKMulticast)); } // Build the CFG after the flat region splice-backs so RegionClonePass can match its diff --git a/shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp b/shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp index f0a0076cd440..23e6e6e2d6bd 100644 --- a/shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp +++ b/shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp @@ -610,6 +610,30 @@ void insertLoopCounterLGatedClusterBarrierWaitBefore(IRBase* anchor, AsmIRBuilde void insertClusterBarrierWaitBefore(IRBase* anchor, const char* comment, AsmIRBuilder& irBuilder, GfxArchID archId); +/// Emit `s_wait_tensorcnt 0` immediately before \p anchor, where \p anchor is +/// the instruction right after a cooperative `tensor_load_to_lds` group. Under +/// PGR>=2 the cluster/broadcast round runs at the TOP of the mainloop, so a +/// pre-round broadcast drain (before the wave-gated cluster-scope arrive) would +/// only retire the PREVIOUS iteration's load -- one iteration too late for the consumer +/// that reads the freshly cooperatively-loaded LDS half at the next loop head +/// (the producing wave is a peer wave, so the consumer's own tensor counter is +/// already zero and cannot order the peer's async write). Draining right after +/// the load issues makes the cooperative broadcast coherent before the back +/// 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) { + const HwInstDesc* waitDesc = getMCIDByUOp(GFX::s_wait_tensorcnt, archId); + assert(waitDesc && "s_wait_tensorcnt opcode is not supported on this architecture"); + StinkyInstruction* w = irBuilder.create(waitDesc, anchor); + w->addSrcReg(StinkyRegister(0)); + SWaitTensorCntData d; + d.tlcnt = 0; + w->addModifier(d); + w->addModifier( + CommentData{"retire cooperative tensor_load_to_lds before back-edge (PGR>=2 coherence)"}); +} + /// Emit Rule 4's cluster-barrier handshake before `anchor` (the iterator /// position right after the load's anchoring `s_barrier_wait -1`). /// @@ -806,8 +830,12 @@ class InsertClusterBarrierPassImpl : public Pass { public: static char ID; - InsertClusterBarrierPassImpl(bool isKernelScope, int pgrValue, int plrValue) - : isKernelScope_(isKernelScope), pgrValue_(pgrValue), plrValue_(plrValue) {} + InsertClusterBarrierPassImpl(bool isKernelScope, int pgrValue, int plrValue, + bool streamKMulticast) + : isKernelScope_(isKernelScope), + pgrValue_(pgrValue), + plrValue_(plrValue), + streamKMulticast_(streamKMulticast) {} const char* getName() const override { return "Insert Cluster Barrier"; @@ -851,6 +879,10 @@ class InsertClusterBarrierPassImpl : public Pass { std::tuple> pending; std::unordered_set seenTriggers; + // Anchors (instruction right after each cooperative tensor_load + // group) for the mode-(c) producer-side drain; see + // insertProducerTensorDrainBefore. + std::vector producerDrainAnchors; auto segBegin = bb.begin(); for (auto it = bb.begin(); it != bb.end(); ++it) { @@ -886,6 +918,27 @@ class InsertClusterBarrierPassImpl : public Pass { sumLoopCounterLDecrementsBeforeInSegment(segBegin, trigger); pending.emplace_back(trigger, std::next(BasicBlock::iterator(trigger)), liveLclCmp, lclPreDecrement); + + // Mode (c) only: record the instruction right after this + // cooperative tensor_load group so a producer-side tensor + // drain can be planted there (retire the load before the back + // edge / next publishing barrier). Advance past any + // immediately-following tensor_load(s) so the drain covers the + // whole group (e.g. the A/B operand load plus its MX-scale + // load) rather than landing between them. + if (kRule4ForceUngatedSignalMode && streamKMulticast_ && pgrValue_ >= 2) { + auto postIt = std::next(it); + while (postIt != bb.end()) { + auto* pinst = dyn_cast(postIt.getNodePtr()); + if (pinst != nullptr && isTensorLoad(*pinst)) { + ++postIt; + continue; + } + break; + } + producerDrainAnchors.push_back((postIt != bb.end()) ? postIt.getNodePtr() + : nullptr); + } } // Rule 1: signal-only handshake immediately AFTER each @@ -1101,6 +1154,12 @@ class InsertClusterBarrierPassImpl : public Pass { liveLclCmp, lclPreDecrement); (void)trigger; // queued for ordering only; insertion uses `anchor` } + // Mode (c): producer-side drain right after each cooperative + // tensor_load group (retire the async cooperative load before the + // back-edge so the next loop-head barrier publishes it coherently). + for (IRBase* postAnchor : producerDrainAnchors) { + insertProducerTensorDrainBefore(postAnchor, irBuilder, archId); + } for (IRBase* anchor : gsu1Anchors) { insertLoopCounterLGatedClusterBarrierSignalBefore( anchor, irBuilder, archId, @@ -1157,15 +1216,17 @@ class InsertClusterBarrierPassImpl : public Pass { const bool isKernelScope_; const int pgrValue_; const int plrValue_; + const bool streamKMulticast_; }; char InsertClusterBarrierPassImpl::ID = 0; } // namespace -std::unique_ptr createInsertClusterBarrierPass(bool isKernelScope, int pgrValue, - int plrValue) { - return std::make_unique(isKernelScope, pgrValue, plrValue); +std::unique_ptr createInsertClusterBarrierPass(bool isKernelScope, int pgrValue, int plrValue, + bool streamKMulticast) { + return std::make_unique(isKernelScope, pgrValue, plrValue, + streamKMulticast); } } // namespace stinkytofu