feat(tensilelite): StreamKForceDPOnly cluster multicast for gfx1250 by jaopaulolc · Pull Request #10189 · ROCm/rocm-libraries

jaopaulolc · 2026-07-30T16:46:24Z

JIRA ID : AIHPBLAS-3929

Motivation

gfx1250 StreamK GEMMs can co-locate consecutive data-parallel workgroups into a hardware workgroup cluster and share operands across the cluster instead of re-loading them per workgroup. This PR enables that for StreamKForceDPOnly=1 (FDPO=1) StreamK, driven purely by ClusterDim = [Cs, Ck]:

Cs = ClusterDim[0] peers sit on M-adjacent output tiles and therefore read the same B over full K, so B is loaded once and TDM-multicast to them;
Ck = ClusterDim[1] peers sit on N-adjacent output tiles and read the same A, so A is multicast along that axis.

Ck == 1 is not a separate feature: it is the degenerate case where A has no peers and its mask collapses to the self bit, while B multicasts exactly as it does for Ck > 1. Both shapes go through one code path.

StreamK=3 with StreamKForceDPOnly=0 is out of scope: this PR leaves that path exactly as develop has it, verified below.

Technical Details

One path for every cluster shape

Everything below fires for any FDPO=1 cluster, [Cs, 1] and [Cs, Ck] alike:

Launch (ContractionSolution::generateSingleCall) spans the real tile space: gridX = nWG0 (M-tiles), gridY = nWG1 * gsu (N-tiles), gridZ = batch, then rounds each extent up to its ClusterDim factor because the HIP driver rejects a cluster launch that is not divisible. getSKGridImpl returns skGrid = tiles (one work-group per output tile, no K-split). The round-up is gated to streamK == 0 || StreamKForceDPOnly != 0, i.e. it does not reach an FDPO=0 Stream-K cluster: that path's 1-D sk.grid is not a tile space and has no pad-exit, so padding it would only add work-groups that run the entire Stream-K prologue before falling out on an empty iteration range. Develop excluded it deliberately and keeps its launch unchanged.
Tile-index fold (StreamK.preLoop) turns the hardware coords back into the linear index the DP decode expects: StreamKIdx = WorkGroup2*(nWG0*nWG1) + WorkGroup1*nWG0 + WorkGroup0, M-fastest so X-peers land M-adjacent.
Padded-peer exit (streamKClusterPadEarlyExit) runs before the fold and before the first cluster barrier: a work-group whose raw M- or N-tile coord is past the real tile count s_endpgms, and its WAVEDONE frees the -3 barrier slot for its surviving peers.
Mask reduction (computeMulticastMaskReduction) trims each broadcast mask to the peers actually present in a boundary cluster, so ld_bcst never waits out the multicast timeout on a lane that exited.

Because absent peers are handled structurally, no runtime clusterMulticastValid selection guard is emitted at all: the padded-peer exit plus reduced masks subsume it on the FDPO=1 path, and the FDPO=0 path derives Multicast=False and so has no masks to guard. For the same reason the prologue cluster arrive is emitted pre-loop: the no-work peers are already gone by the time it executes.

Validation applies to solutions the derivation has already identified as cluster multicast (StreamK=3, Cs > 1, FDPO=1): Cs and Ck each a power of two with 2 <= Cs*Ck <= 16, TDMInst=3 on gfx1250 with HasTDM/HasClusterBarrier, no atomic Stream-K, no XCC remap. A cluster with no B-sharing peers ([1, C]) is rejected. SK4 and SK5 clusters are rejected outright.

FDPO=1 is part of the detector, not a per-call-site check

streamKMulticast() now derives the path from StreamK=3 and Cs > 1 and StreamKForceDPOnly=1, instead of returning true for any SK3 cluster and relying on each call site to add the FDPO conjunction. Two emitters were missing that conjunction and were reaching the FDPO=0 cluster: the zero-iteration cluster-barrier wait and the prologue prefetch handshake. Folding the condition into the detector closes both and removes the redundant check from the five sites that had it.

Both XCC remaps are neutralized under a cluster

A cluster's peers are only adjacent in the unremapped work-group id space, so any remap of WorkGroup0 breaks the tile adjacency that the cooperative load depends on. Develop already clears StreamKXCCMapping for a clustered Stream-K; this PR additionally forces WorkGroupMappingXCC — the second WorkGroup0 remap, the wgmXCC CU-count one — to identity in the same block, so "any XCC remapping disabled with ClusterDim != [1,1] + Stream-K" holds for both. It is set before the auto-WGMXCC check so a clustered kernel never reaches that check still holding -1.

This is behaviour-neutral today and was measured rather than assumed: across all ten clustered Stream-K configs (develop's and this PR's, 617 solutions), every solution already derived the identity value 1, no config sets the parameter under a clustered Stream-K, and the full derived state and solution counts are identical with and without the change. It matters for the multi-XCD architectures CI exercises.

Workspace SRD is computed in 64 bits

computeWorkspaceSrd advanced the partials SRD with a 32-bit s_mul_i32 and then folded only the carry: s_addc_u32 SrdWS+1, SrdWS+1, 0. The product is MacroTile0 * MacroTile1 * bpeCinternal * partialIdx, which exceeds 2^32 for large Stream-K grids — at MT256x256 with fp32 partials the per-tile stride is 256 KiB, so the offset wraps past a partial index of 16384. On wrap the SRD silently aliases the wrong workspace slot: a peer writes over another work-group's partials and the owner reads a slot that was never written, i.e. a wrong answer with no fault. The fix computes the high word with s_mul_hi_u32 and folds it, plus the low-add carry, into SrdWS+1.

This is not cluster-specific: computeWorkspaceSrd is emitted by every Stream-K kernel that writes partials — SK3/SK4/SK5, cluster and non-cluster, on every architecture. It is not emitted at all under FDPO=1, which is DP-only and never writes partials (confirmed: 0 of the 38 FDPO=1 characterization kernels contain it, against 96 of 96 for FDPO=0). It rides in this PR by maintainer decision; coverage on the other architectures comes from CI.

Supporting pieces

ClusterLoad TDM component (Tensile/Components/ClusterLoad.py) owns the multicast-mask SGPR lifecycle (declareSgprs / undeclareSgprs / computeMasks / applyToDescriptor), used by SubtileGREmit, KernelWriter and KernelWriterAssembly instead of hardcoded mask SGPRs. Under PrefetchAcrossPersistent the masks must survive the persistent refresh; the one exception is a Ck == 1 A mask, which is self-only, so its SGPR is freed to stay inside the 106-SGPR budget.
Derived Multicast / ClusterBarrier. Both stay derived-only bools, as on develop, and are not user-settable. The single change is that the derivation now also turns them on for the FDPO=1 SK3 cluster this PR introduces. A differential enumeration over all 3200 reachable (ClusterDim, StreamK, StreamKForceDPOnly, TDMInst, HasClusterBarrier) combinations confirms the derived pair is identical to develop everywhere except exactly the 272 combinations with StreamK=3 and ClusterDim[0]>1 and StreamKForceDPOnly=1. As on develop, ClusterBarrier is only ever enabled within the multicast case.
Detectors streamKMulticast (SK3, Cs > 1, FDPO=1) and streamK2DMulticast (Ck > 1, i.e. A has peers) are derived from existing state; there is no separate knob.
PrefetchGlobalRead=2 emits a second, double-buffered prologue prefetch inside the single-iteration guard, past the generic bracketing boundary, so streamKMulticastProloguePrefetchHandshake gives it its own cluster-scope handshake.
stinkytofu (additive): StreamKMulticast module option and an InsertClusterBarrierPass producer-drain for PGR>=2, plumbed through Gfx1250Backend.

Configs removed

sk_mxf4gemm_pap_cluster.yaml and sk_mxf8gemm_pap_cluster.yaml are removed. Both forked StreamKForceDPOnly over [0, 1] beneath a ClusterDim. Cluster multicast is derived only for StreamKForceDPOnly=1, so the FDPO=0 half of each config cannot exercise the path this PR adds: deriving the merge-base configs against this tree, all 24 FDPO=0 permutations of the mxf4 config and all 18 of the mxf8 config come out Multicast=False, ClusterBarrier=False. Their FDPO=1 half is covered, over a wider ClusterDim set, by sk_mxf4_force_dp_only_pap_cluster_multicast.yaml and sk_mxf8_force_dp_only_pap_cluster_multicast.yaml. Develop's own FDPO=0 cluster coverage is untouched: sk_mxf4gemm_tdm_cluster.yaml and streamk_cluster.yaml are carried unchanged and verified below.

An earlier commit message on this branch (d7e03de7) described these two configs as hanging. That description is withdrawn: it was investigated and did not reproduce. These configs are removed for the scope reason above, not because of a hardware failure.

Test Plan

CPU unit: ClusterLoad component, streamKMulticast detection and the validation matrix, SK4/SK5-cluster reject.
gfx1250 codegen characterization: test_streamk_cluster_multicast_gfx1250_char.py over PrefetchGlobalRead in {1, 2} x ClusterDim in {[4,1], [2,2]}, test_streamk_pap_cluster_multicast_gfx1250_char.py for the PAP mask-liveness contrast, and test_streamk_cluster_coop_load_gfx1250_char.py. Develop's FDPO=0 test_streamk_cluster_gfx1250_char.py is kept unchanged as the FDPO=0 guard.
gfx1250 FFM functional simulation of the eight FDPO=1 core configs — four cluster multicast and four non-cluster baselines, each cluster config sweeping ClusterDim [2,1] [4,1] [8,1] [2,2] [2,4] [4,2] — plus develop's FDPO=0 cluster config, run sequentially with zero skipped configs. The sizes deliberately include tile counts that are not cluster multiples, so the padded-peer exit and reduced masks are exercised rather than assumed.

MT256x256 is intentionally excluded for F8: that kernel needs 1066 VGPRs, above the 1024 hardware limit, and fails assembly generation. That is a datatype register-budget limit, not a cluster predicate.

Test Result

CPU unit suite: 6165 passed, 29 skipped, 0 failed.
gfx1250 characterization: all cluster drivers pass, 760 snapshots passed. No develop-owned snapshot moves.
gfx1250 FFM: 7584 validations over the eight FDPO=1 configs, all 7584 passed — 0 failed, 0 did-not-satisfy, 0 invalid, no hangs or aborts, and no kernel over its register budget. Develop's FDPO=0 cluster config: 192 passed, 0 failed.

Develop's FDPO=0 cluster configs emit assembly unchanged from the merge-base apart from the workspace SRD fix, verified by emitting both trees.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Risk level

Low for the cluster work: every path is gated on gfx1250, StreamK=3, StreamKForceDPOnly=1 and a ClusterDim other than [1, 1]; non-cluster and FDPO=0 configs emit unchanged assembly apart from the workspace SRD fix, verified by the derived-state enumeration and by an assembly and FFM comparison against the merge-base. The change is largely additive (a new ClusterLoad component, new unit and characterization tests, new configs), and the host-side launch change is confined to the same gate.

The workspace SRD fix is the one change with reach beyond that gate: it touches every Stream-K kernel that writes partials, on every architecture. It is a strict correctness fix for a silent-aliasing wrong answer, it adds one s_mul_hi_u32 per SRD setup, and gfx1250 Stream-K is verified here; the other architectures are covered by CI.

therock-pr-bot · 2026-07-30T16:51:06Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

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

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

🙋 Wish to Override Policy?

therock-pr-bot · 2026-07-30T16:51:09Z

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

codecov · 2026-07-31T04:16:10Z

Codecov Report

❌ Patch coverage is 84.05316% with 48 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
.../hipblaslt/tensilelite/src/ContractionSolution.cpp	14.29%	17 Missing and 1 partial ⚠️
...lt/tensilelite/Tensile/SolutionStructs/Solution.py	70.27%	8 Missing and 3 partials ⚠️
...ipblaslt/tensilelite/Tensile/Components/StreamK.py	87.80%	6 Missing and 4 partials ⚠️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	84.38%	0 Missing and 5 partials ⚠️
...lelite/Tensile/Components/Subtile/SubtileGREmit.py	0.00%	3 Missing ⚠️
...aslt/tensilelite/Tensile/Components/ClusterLoad.py	99.05%	0 Missing and 1 partial ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #10189      +/-   ##
===========================================
+ Coverage    69.35%   69.36%   +0.01%     
===========================================
  Files         2777     2778       +1     
  Lines       457902   458092     +190     
  Branches     67514    67544      +30     
===========================================
+ Hits        317557   317735     +178     
- Misses      117214   117224      +10     
- Partials     23131    23133       +2

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`38.05% <14.29%> (-0.01%)`	⬇️
TensileLite-Unit	`76.57% <89.29%> (+0.07%)`	⬆️
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 45d1ba8
hipBLASLt	`34.89% <ø> (ø)`
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 45d1ba8
hipDNN	`86.80% <ø> (ø)`		Carriedforward from 45d1ba8
hipFFT	`43.60% <ø> (ø)`		Carriedforward from 45d1ba8
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 45d1ba8
hipSOLVER	`69.03% <ø> (ø)`		Carriedforward from 45d1ba8
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 45d1ba8
rocBLAS	`48.26% <ø> (ø)`		Carriedforward from 45d1ba8
rocFFT	`47.64% <ø> (ø)`		Carriedforward from 45d1ba8
rocRAND	`57.05% <ø> (ø)`		Carriedforward from 45d1ba8
rocSOLVER	`77.12% <ø> (ø)`		Carriedforward from 45d1ba8
rocSPARSE	`74.61% <ø> (ø)`		Carriedforward from 45d1ba8
rocThrust	`91.60% <ø> (ø)`		Carriedforward from 45d1ba8

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

Files with missing lines	Coverage Δ
.../hipblaslt/tensilelite/Tensile/Common/Utilities.py	`100.00% <100.00%> (ø)`
...rojects/hipblaslt/tensilelite/Tensile/Component.py	`95.74% <100.00%> (+0.03%)`	⬆️
...pblaslt/tensilelite/Tensile/Components/__init__.py	`100.00% <ø> (ø)`
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py	`72.93% <100.00%> (-0.01%)`	⬇️
...aslt/tensilelite/Tensile/Components/ClusterLoad.py	`99.05% <99.05%> (ø)`
...lelite/Tensile/Components/Subtile/SubtileGREmit.py	`70.65% <0.00%> (-0.11%)`	⬇️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	`70.75% <84.38%> (+0.07%)`	⬆️
...ipblaslt/tensilelite/Tensile/Components/StreamK.py	`83.58% <87.80%> (+0.12%)`	⬆️
...lt/tensilelite/Tensile/SolutionStructs/Solution.py	`58.97% <70.27%> (+0.06%)`	⬆️
.../hipblaslt/tensilelite/src/ContractionSolution.cpp	`7.81% <14.29%> (+0.05%)`	⬆️

... and 1 file with indirect coverage changes

🚀 New features to boost your workflow:

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

Enable hardware workgroup-cluster operand sharing for StreamK with StreamKForceDPOnly=1 on gfx1250, driven entirely by ClusterDim = [Cs, Ck]. Cs peers sit on M-adjacent output tiles and so read the same B over full K, letting B be loaded once and TDM-multicast to them; Ck peers sit on N-adjacent tiles and read the same A, multicast along that axis. Ck == 1 is the degenerate case where A has no peers and its mask collapses to the self bit, so both cluster shapes go through a single code path. - Launch spans the real tile space (gridX = nWG0, gridY = nWG1 * gsu, gridZ = batch), rounding each extent up to its ClusterDim factor because the HIP driver rejects a cluster launch that is not divisible. The round-up is gated so it does not reach an FDPO=0 StreamK cluster, whose 1-D grid is not a tile space and has no pad exit. - StreamK.preLoop folds the hardware coords back into the linear index the data-parallel decode expects, M-fastest so X-peers land M-adjacent. - A padded-peer exit runs before the fold and before the first cluster barrier, and mask reduction trims each broadcast mask to the peers actually present, so a boundary cluster neither waits out the multicast timeout on an exited lane nor holds a barrier slot. - streamKMulticast() derives the path from StreamK=3, Cs > 1 and StreamKForceDPOnly=1 rather than leaving the FDPO conjunction to each call site, which closes two emitters that previously reached an FDPO=0 cluster: the zero-iteration barrier wait and the prologue prefetch handshake. - Add a ClusterLoad TDM component owning the multicast-mask SGPR lifecycle, used by SubtileGREmit, KernelWriter and KernelWriterAssembly in place of hardcoded mask SGPRs. - Force WorkGroupMappingXCC to identity under a clustered StreamK, next to develop's existing StreamKXCCMapping clear, since a cluster's peers are only adjacent in the unremapped workgroup id space. - Compute the workspace SRD partials offset in full 64 bits; the previous 32-bit product could wrap for large StreamK grids and silently alias another workgroup's slot, giving a wrong answer with no fault. - stinkytofu gains a StreamKMulticast module option and an InsertClusterBarrierPass producer-drain for PrefetchGlobalRead >= 2. - Add and update gfx1250 configs and characterization drivers covering the FDPO=1 cluster multicast shapes and the non-cluster FDPO=1 baselines. StreamK with StreamKForceDPOnly=0 is out of scope and is left as develop has it; its emitted assembly is unchanged apart from the workspace SRD fix.

…icast

Copilot

Pull request overview

Enables gfx1250 StreamK “ForceDPOnly” (FDPO=1) cluster multicast for StreamK=3 by deriving the multicast path from ClusterDim=[Cs,Ck] (B multicast along Cs, A multicast along Ck), including launch/grid padding, padded-peer early-exit, and boundary-mask reduction so cooperative ld_bcst doesn’t wait on exited peers.

Changes:

Add StreamK FDPO cluster-multicast detection/validation and codegen plumbing (derived Multicast/ClusterBarrier, cluster tile-index fold, padded-peer early-exit, 64-bit workspace SRD fix).
Introduce a ClusterLoad TDM component to own multicast mask SGPR lifecycle, mask computation, and descriptor attachment (including PAP mask-liveness handling).
Add/extend gfx1250 CPU-only unit + characterization coverage and update StreamK gfx1250 test configs (including PAP and 2-D cluster shapes).

Reviewed changes

Copilot reviewed 37 out of 37 changed files in this pull request and generated no comments.

Show a summary per file

File	Description
shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp	Add optional StreamKMulticast-gated producer-side `s_wait_tensorcnt 0` drain placement for PGR>=2 cooperative loads.
shared/stinkytofu/src/pipeline/backend/Gfx1250Backend.cpp	Pass `StreamKMulticast` module option into the cluster-barrier insertion pass.
shared/stinkytofu/include/stinkytofu/transforms/asm/InsertClusterBarrierPass.hpp	Extend pass factory signature to accept `streamKMulticast` gating flag (defaulted).
shared/stinkytofu/include/stinkytofu/bindings/python/Module.hpp	Expose `StreamKMulticast` module option in Python bindings.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_multicast.py	New unit tests for StreamK multicast derivation, validation matrix, and mask arithmetic.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_cluster_sk45_reject.py	New unit tests asserting SK4/SK5 reject `ClusterDim != [1,1]` (control vs. reject).
projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PrefetchAcrossPersistent.py	Update unit harness state so `ClusterLoadTDM.find()` can match in PAP-related tests.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_cluster_load_component.py	New unit tests for `ClusterLoad` component selection, SGPR lifecycle, and emitted-asm snippets.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/conftest.py	Add characterization `_codegen` directory to `sys.path` for shared harness imports.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_pap_cluster_multicast_gfx1250_char.py	New CPU-only characterization pinning PAP mask liveness + SGPR budget behavior for 1-D vs 2-D clusters.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_multicast_gfx1250_char.py	New CPU-only characterization pinning StreamK cluster multicast asm expectations + goldens.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_streamk_cluster_coop_load_gfx1250_char.py	New CPU-only characterization for StreamK=3 + ClusterDim cooperative-load path basics.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_pap_cluster_multicast.yaml	New designed config for PAP + multicast characterization.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast.yaml	New designed config for PGR=1 multicast characterization.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_multicast_pgr2.yaml	New designed config for PGR=2 multicast characterization.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_cluster_coop_load.yaml	New designed config for basic StreamK+cluster cooperative-load characterization.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/config_harness.py	Extend harness with `cluster_dim` filtering and shared StreamK-cluster assertions/helpers.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/snapshots/test_streamk_pap_cluster_multicast_gfx1250_char.ambr	New golden digests for PAP multicast characterization.
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/snapshots/test_streamk_cluster_multicast_gfx1250_char.ambr	New golden digests for multicast characterization arms.
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only.yaml	Add/refresh gfx1250 MXFP8 FDPO StreamK test coverage (including PGR sweep).
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only_pap.yaml	Narrow to FDPO=1 and align parameters/sizes with non-PAP counterpart for diffability.
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only_pap_cluster_multicast.yaml	New PAP + cluster-multicast MXFP8 coverage including 2-D cluster shapes.
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf8_force_dp_only_cluster_multicast.yaml	New cluster-multicast MXFP8 coverage including 2-D cluster shapes.
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only.yaml	Expand MXFP4 FDPO coverage (MI/PGR sweep and broader size set).
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_pap.yaml	Align PAP MXFP4 FDPO config with non-PAP baseline; force FDPO=1.
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_pap_cluster_multicast.yaml	New PAP + cluster-multicast MXFP4 coverage including 2-D cluster shapes.
projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx1250/core/sk_mxf4_force_dp_only_cluster_multicast.yaml	Extend MXFP4 cluster-multicast coverage to include 2-D cluster shapes and size set.
projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py	Add StreamK multicast validation + SK4/SK5 cluster rejection; derive `Multicast`/`ClusterBarrier` for FDPO clusters; disable XCC remaps under StreamK clusters.
projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py	Use `ClusterLoadTDM` component to compute/apply masks; enable boundary-mask reduction for StreamK multicast; add zero-iter cluster wait emission.
projects/hipblaslt/tensilelite/Tensile/KernelWriter.py	Centralize mask SGPR lifecycle via `ClusterLoadTDM`; plumb `StreamKMulticast` module option; preserve cluster-scope barriers in barrier pass-2 rebuild.
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGREmit.py	Attach multicast masks via `ClusterLoadTDM` for subtile path.
projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py	Add 64-bit workspace SRD fix; add cluster multicast prologue signal/handshakes, padded-peer early-exit, and tile-index fold for FDPO clusters.
projects/hipblaslt/tensilelite/Tensile/Components/ClusterLoad.py	New component owning multicast mask SGPR naming, compute, attach, and PAP liveness behavior.
projects/hipblaslt/tensilelite/Tensile/Components/init.py	Register new `ClusterLoad` component module.
projects/hipblaslt/tensilelite/Tensile/Component.py	Add `ClusterLoad` component base class.
projects/hipblaslt/tensilelite/Tensile/Common/Utilities.py	Add helpers: `isPow2`, `streamKMulticast`, `streamK2DMulticast`.
projects/hipblaslt/tensilelite/src/ContractionSolution.cpp	Adjust StreamK launch dimensions for FDPO clusters and allow cluster grid padding for FDPO multicast path; set `skGrid=tiles` for that case.

…icast Conflict resolution: adopt develop's InsertClusterBarrierPass rewrite (PR 10210) and re-plant this branch's PGR>=2 producer tensor drain onto its new Rule 3 scan. This merge also repairs a develop regression it would otherwise inherit. Commit 63b77e6 ("Disable staggerU when workgroup cluster is enabled") turns off staggerU codegen whenever ClusterDim != [1,1]. That activates the pre-existing StreamK const-increment optimization, which stops GlobalReadIncs{A,B,MXSA,MXSB} from being declared while two sites still reference them by name, so every StreamK + ClusterDim + TDM/MX kernel fails to assemble with "expected absolute expression". It breaks configs already on develop too, including sk_mxf4gemm_tdm_cluster.yaml (added by PR 9464). The fix keeps real SGPRs for swizzled MX scale increments, whose Size*DepthU*bpe value is only known at runtime and so cannot become a constant, and teaches the wave-separated TDM increment setup to honour useConstSgprGlobalReadIncsForTc. StreamKTwoTileDPFirst's short-lived SKMappingTemp scratch now grows the SGPR pool instead of rejecting the solution when no aligned hole is free. The cluster characterization tests could not catch this class of bug because they only pattern-match generated assembly text. assert_assembles() now feeds each emitted cluster kernel to the assembler.

therock-pr-bot · 2026-08-04T20:25:42Z

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

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

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

…icast

Aligns the ternary continuation in the PGR>=2 producer drain anchor collection, which the pre-commit clang-format hook rejected.

…icast

…rvation Keeping GlobalReadIncsMXSA/MXSB in real SGPRs for swizzled MX scale layouts is only needed where GSU.graIncrements actually emits the SGPR increment. Subtile builds its own descriptors and never reads GlobalReadIncs (the Subtile global read path has no references, and the increment setup is skipped for it), so the reservation there buys nothing and only consumes registers. That is fatal on gfx950, where Solution forces MX kernels onto subtile: seven StreamK + MX kernels with a HostPreSwizzle scale layout ran out of SGPRs and failed codegen with "Generating kernel source resulted in error 2", which aborts TensileCreateLibrary and takes the whole device-library build down. The abort surfaces as a nanobind reference-leak dump, which hides the real cause. Verified by a clean gfx950 build of the device libraries, which now succeeds where it previously failed on those seven kernels, and by re-running the gfx1250 StreamKForceDPOnly FFM matrix: 7584 validations pass with 0 failures across all eight configs, each at its established count, so the cluster multicast kernels still get the SGPR they require.

The cluster configs need hardware that only part of the gfx1250 fleet provides, so they move next to the existing _ext configs and take the same suffix, leaving core to the sweeps that run anywhere. Also drops comments that only recorded how these files relate to each other or which sizes were inherited from where; what remains is the coverage each file is there to provide. No parameter, cluster shape or problem size changes: each config parses identically to its previous revision.

…icast Conflict resolution, SubtileGREmit.py: develop reworked the subtile TDM descriptor for iterate mode and this branch replaced the inline multicast mask write with a ClusterLoad component call. Both are kept. The import line was the only textual conflict: develop's isSubtileIterateMode import is taken, and clusterEnabled is dropped because the block that used it is exactly the block this branch replaced, leaving it with no remaining use in the file.

…icast

…icast Four conflicts, all resolved by keeping both sides: Module.hpp and KernelWriter.py: develop added a TDMLoadWaveSync module option in the same slot this branch uses for StreamKMulticast. Both options are kept. Solution.py: develop began gating Multicast on the HasTDMMulticast arch capability, while this branch had widened multicast to the DP-only Stream-K cluster and tied ClusterBarrier to the multicast decision. The shared condition is now named, so multicast requires both peer-tile sharing and the capability, and ClusterBarrier keeps develop's behaviour of depending only on peer-tile sharing rather than on the capability. Gfx1250Backend.cpp: develop moved cluster-barrier insertion into the main-only pipeline and replaced the kernel-scope block with a per-function CFG rebuild. Develop's structure is taken, with this branch's streamKMulticast and pgrValue arguments re-planted onto the relocated call.

…icast One conflict, test_PrefetchAcrossPersistent.py: develop added a waveIdxReleasedAfterStagger field to the setup-new-tile PAP TDM mock state while this branch added the asmCaps/kernel entries that ClusterLoadTDM.find() matches on. Both are kept.

…icast

…M parameter TDMLoadWaveSync was added as a solution parameter, so it now appears in every kernel name as TDMLWS0 at its default. The names are hashed into the {basename, err} digest these goldens record, so all 36 entries across the two cluster multicast snapshots shifted. Only the basenames moved: err stays 0 for every kernel, the kernel count per arm is unchanged, and the behavioural assertions in the same files -- broadcast masks bound to their descriptors, the tile-index fold, the padded-peer exit, cluster barrier balance, the PGR2 prologue handshake, and that each kernel assembles -- all still pass untouched. Recorded per node rather than by bulk regeneration, and re-run twice afterwards to confirm the digests are stable.

…icast

Conflict resolution: - Tensile/Components/StreamK.py, Tensile/KernelWriter.py, Tensile/KernelWriterAssembly.py and Tensile/Tests/unit/test_streamk_queue_index.py are resolved to develop. The branch carried an earlier version of the count-preserving per-XCD StreamK work-queue counter fix from when this was a stacked PR; an equivalent fix has since landed independently on develop (#10008). That work is out of scope for this PR, so develop's landed version wins and these files are byte-identical to develop. - include/Tensile/ContractionSolution.hpp keeps both sides: develop's CustomKernel, StreamKStaticSplit / streamKStaticSplit() / streamKStaticSplitRowUniform(), UniformSummationOrderError, the four-element WGMParamsCache key and the checkUniformSummationOrder() declaration, plus this branch's StreamKDecisions struct and the computeStreamKDecisions() / printStreamKLaunchSummary() declarations. Post-merge mirror fixes in src/ContractionSolution.cpp, so the diagnostic report still describes the launch it observes (report-only; no launch behaviour changes): - getSKGridImpl() gained a StreamKForceDPOnly cluster-multicast clamp (#10189) that runs after the tree-fixup-bounds fallback and also resets the grid to the tile count. The summary had no way to attribute that clamp, so a launch taking it reported "changedBy = none". Added a clusterDPGridClamped out-param and StreamKDecisions field, reported under "fallbacks" and in the changedBy attribution (ordered latest-clamp-wins). - The static-split mirror now calls the shared streamKStaticSplit() helper develop factored out, instead of repeating its skTiles arithmetic inline. makeArgs() packs skTiles from that same helper, so the report can no longer drift from the kernel arguments.

…stent tiles (#12817) JIRA ID : AIHPBLAS-3929 ## Motivation On gfx1250, TensileLite lets neighboring GPU workgroups share matrix-input loads through multicast. DataParallel previously limited these workgroups to one complete output tile each. This change lets each cluster process successive blocks of tiles while keeping its members synchronized. ## Technical Details - For `DataParallel` with `StaticGrid`, assign whole blocks of tiles matching `ClusterDim` to each cluster. Round the launch grid to complete clusters and reject padded tile counts exceeding the 32-bit scheduling limit. - Keep members beyond the matrix edge involved in shared loads by loading the nearest valid tile and suppressing their output writes. - Synchronize before reusing shared memory for another block. Keep operand and sparse-metadata multicast masks live across tiles, and balance synchronization when `K=0` or `alpha=0` skips the summation. Prefetch between output tiles (`PrefetchAcrossPersistent`) still requires one block per cluster, so clusters do not repeat in that configuration. Operand reuse (`ReuseAcrossPersistent`) and tile remapping (`SpaceFillingAlgo`) are rejected for this path. StreamK clustered reduction retains its existing scheduling. Stack 5/5, based on [PR #12540](#12540). Fixes scheduling and synchronization in the persistent DataParallel multicast support introduced by [PR #10189](#10189). ## Test Plan - Run the full `tox` suite from `tensilelite/` on gfx1250, covering Python/common tests, rocisa tests, and lint. - Run `data_parallel_cluster_multicast_boundaries.yaml` with `TENSILE_PERSISTENT_FIXED_GRID=4` and alpha values 1 and 0 to check repeated tiles, uneven batched boundaries, and skipped summations. ## Test Result Local validation on gfx1250: - Full `tox` functional suite: 10,186 passed, 0 failed, 1,758 skipped; 786 snapshots passed. - rocisa: 126 passed, 4 skipped. - Lint still fails, including 8 new unused-import diagnostics. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium: changes GPU launch geometry, cluster synchronization, and register lifetimes, which can affect correctness and performance.

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast shared: stinkytofu labels Jul 30, 2026

assistant-librarian Bot added the organization: ROCm label Jul 30, 2026

jaopaulolc changed the title ~~feat(tensilelite): StreamK Force-DP-Only cluster multicast (FDPO=1) for gfx1250~~ feat(tensilelite): StreamKForceDPOnly cluster multicast for gfx1250 Jul 30, 2026

jaopaulolc added the rocm:gemm algos Label to identify PRs of the GEMM Algorithms team. label Jul 30, 2026

jaopaulolc force-pushed the users/jolabega/streamk-force-dp-only-multicast branch from 446caf0 to 42152d0 Compare July 31, 2026 20:12

jaopaulolc marked this pull request as ready for review July 31, 2026 20:14

jaopaulolc requested review from a team as code owners July 31, 2026 20:14

jaopaulolc added 2 commits July 31, 2026 20:15

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

88ea31b

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

56a9703

…icast

jichangjichang approved these changes Aug 4, 2026

View reviewed changes

jichangjichang requested a review from Copilot August 4, 2026 02:10

Copilot started reviewing on behalf of jichangjichang August 4, 2026 02:11

View session

Copilot AI reviewed Aug 4, 2026

View reviewed changes

KKyang approved these changes Aug 4, 2026

View reviewed changes

jaopaulolc added 6 commits August 4, 2026 15:43

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

5d75955

…icast

style(stinkytofu): fix clang-format continuation alignment

f38c000

Aligns the ternary continuation in the PGR>=2 producer drain anchor collection, which the pre-commit clang-format hook rejected.

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

67ad914

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

a91cadf

…icast

jaopaulolc added 10 commits August 7, 2026 17:44

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

6b076b1

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

95ad944

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

16b564e

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

9db5785

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

bf6b5dc

…icast

Merge branch 'develop' into users/jolabega/streamk-force-dp-only-mult…

f2579f9

…icast

jaopaulolc enabled auto-merge (squash) August 18, 2026 17:17

jaopaulolc disabled auto-merge August 18, 2026 17:19

jaopaulolc merged commit 0e24fd6 into develop Aug 18, 2026
75 of 80 checks passed

jaopaulolc deleted the users/jolabega/streamk-force-dp-only-multicast branch August 18, 2026 21:07

jaopaulolc mentioned this pull request Aug 18, 2026

feat(tensilelite): StreamK cluster multicast for ForceDPOnly 0 and 1 #10969

Closed

1 task done

ThanHenderson mentioned this pull request Sep 30, 2026

fix(tensilelite): correct DataParallel cluster multicast across persistent tiles #12817

Merged

1 task done

Navigation Menu

Repository navigation

feat(tensilelite): StreamKForceDPOnly cluster multicast for gfx1250 - #10189

feat(tensilelite): StreamKForceDPOnly cluster multicast for gfx1250#10189
jaopaulolc merged 20 commits into
developfrom
users/jolabega/streamk-force-dp-only-multicast

jaopaulolc commented Jul 30, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Jul 30, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Jul 30, 2026

Uh oh!

codecov Bot commented Jul 31, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

therock-pr-bot Bot commented Aug 4, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

jaopaulolc commented Jul 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

One path for every cluster shape

FDPO=1 is part of the detector, not a per-call-site check

Both XCC remaps are neutralized under a cluster

Workspace SRD is computed in 64 bits

Supporting pieces

Configs removed

Test Plan

Test Result

Submission Checklist

Risk level

Uh oh!

therock-pr-bot Bot commented Jul 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Jul 30, 2026

Uh oh!

codecov Bot commented Jul 31, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

therock-pr-bot Bot commented Aug 4, 2026

Pre-commit check failed

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

jaopaulolc commented Jul 30, 2026 •

edited

Loading

therock-pr-bot Bot commented Jul 30, 2026 •

edited

Loading

codecov Bot commented Jul 31, 2026 •

edited

Loading