feat(ck-tile): enable gfx1250 Stream-K on the comp_tdm and comp_tdm_v2 pipelines by ozturkosu · Pull Request #15083 · ROCm/rocm-libraries

ozturkosu · 2026-11-04T18:18:45Z

JIRA ID: AICK-2558

Builds on #12340 (native gfx1250 Stream-K) and #13975 (comp_tdm_pc bridge), both in develop. Independent of #14962 (Stream-K for comp_tdm_pc): either PR can merge first, see "Relation to other PRs". The remaining pipelines are in the follow-up TODO list at the end.

Motivation

On gfx1250, Stream-K only runs the mem/compv3/compv4 pipelines. The TDM pipelines are the fastest universal GEMM pipelines there, but StreamKKernel could not host them:

comp_tdm and comp_tdm_v2 have no call form for one A window, one B window and a runtime hot-loop/tail, which is how StreamKKernel calls a pipeline.
comp_tdm does not define UsePersistentKernel, and StreamKKernel requires it to match the tile partitioner.
The dispatcher rejects every gfx1250-only pipeline for stream_k and only builds the cshuffle epilogue.

The shapes Stream-K is meant for (small M x N, very large K) gain the most. On H (crr 192x448x393216) and M (crr 118x102x368640) the best existing Stream-K kernels reach 228.9 and 36.8 TFLOPS, and the best universal kernels 30.1 and 3.27 TFLOPS. Stream-K on comp_tdm_v2 reaches 297.2 and 64.3 TFLOPS.

Changes

Pipelines (comp_tdm, comp_tdm_v2)

New operator()(a_window, b_window, num_loop, has_hot_loop, tail_number, p_smem), enabled only for non-tuple windows. It wraps the windows in tuples and forwards to the existing tuple form, so Stream-K runs the same pipeline entry as universal GEMM (hot loop and tail are derived from num_loop the same way). comp_tdm_v2 restates it because its operator() overloads hide comp_tdm's.
comp_tdm: UsePersistentKernel = Traits::UsePersistentKernel && !MX, as in comp_async. comp_tdm_v2 inherits it.

Kernel (`StreamKKernel`)

After each Stream-K loop iteration, comp_tdm and comp_tdm_v2 wait for the pipeline's last TDM prefetch (s_wait_tensorcnt<0>()) before the next tile refills the same LDS. A workgroup that only stores a partial tile runs no epilogue, so nothing else waits for it there. comp_tdm_pc (wavelet) is not affected.
Compile-time checks: no TdmEpilogue with atomic reduction (it overwrites C); no plain-GEMM-only pipelines (comp_tdm_v3); no TDM cluster launch (every workgroup has its own K range); comp_tdm and comp_tdm_v2 need CShuffleEpilogue or TdmEpilogue (the two that wait for the last prefetch) and row-major C.
IsSupportedArgument, for TDM pipelines: check for a gfx125 device (the TDM instructions are no-ops elsewhere). With a non-TDM epilogue, also reject a C whose contiguous dimension is not a multiple of the block tile. UniversalGemmKernel skips its tile checks for TDM pipelines because the TDM loads clip, but the CShuffle store does not. Without this check my prototype gave wrong results on N=448 and crashed on N=102.
grouped_gemm_kernel.hpp: removed a comment that said the TDM pipelines do not define UsePersistentKernel.

Dispatcher

codegen_common.streamk_trait_reject_reason(pipeline, epilogue, reduction, pads): the one Stream-K trait rule set. Atomic needs cshuffle; linear and tree take cshuffle or tdm; TDM pipelines keep the no-pad rule; gfx1250-only pipelines outside STREAMK_GFX1250_PIPELINES = ("comp_tdm", "comp_tdm_v2") are rejected.
gfx1250_pipeline_reject_reason(..., streamk=True): accepts only STREAMK_GFX1250_PIPELINES and skips the universal "requires epilogue=tdm" and "no persistent kernel" rules. The intrawave, exact-gfx1250, pad and comp_tdm_v2 4-wave rules stay.
unified_gemm_codegen.py (STREAM_K branch), gemm_utils.expand_sweep and arch_filter (GEMM_STREAMK) all use these two functions, so codegen, sweep expansion and arch filtering agree. The Stream-K codegen now emits the TDM epilogue for linear and tree.
tile_engine/ops/gemm/gemm_vector_fallback._tile_fits: a TDM pipeline with a non-TDM epilogue needs the contiguous C extent to be a tile multiple, the same as the kernel check above.
comp_tdm_pc and comp_async stay rejected for stream_k, with the same reason text as before. Other variants and the compv3/compv4/mem Stream-K kernel set are unchanged.

I copied the gate design from #14962's dispatcher change, so the two PRs differ only in the pipeline tuple.

Tests

test/ck_tile/gemm_streamk/test_gemm_streamk_tdm_gfx1250_wmma.cpp (TestGfx1250StreamKTdm), 36 configs. The core is bf16 row/col/row with a 128x128x64 tile: comp_tdm and comp_tdm_v2 x {atomic + CShuffle, linear + CShuffle/TDM, tree + CShuffle/TDM} x persistent and non-persistent. On top: fp16, a 64x64x64 tile, col/row, row/row and col/col A/B, and comp_tdm on eight waves (4x2x1, 2x4x1).
Each config runs 8 tests: data-parallel only, split-K segments of 1-5 iterations (both tails, with and without the hot loop), segments that cross tiles, mixed data-parallel and Stream-K, ragged K, ragged M/N, a full-device grid, and a misaligned-pointer rejection.
CMake: the file goes into the existing test_ck_tile_streamk_gfx1250_wmma (REGRESSION_TESTS, 288 new tests) and _smoke targets. The smoke build keeps 3 atomic-only configs (21 tests), because the emulator runs one workgroup at a time and a linear or tree tile owner would wait forever there (feat(ck-tile): enable gfx1250 GEMM Stream-K reductions #12340's split). The _legacy target does not build it: legacy buffer addressing has no LDS transpose load for the TDM pipelines, which is already the case on develop.
dispatcher/tests/test_streamk_tdm_gfx1250.py (11 tests, 951 subtests), registered in dispatcher/tests/CMakeLists.txt. It checks the full pipeline x epilogue x reduction x pads x persistent x waves matrix against an independent oracle and the reject reasons. It runs the real codegen and checks it against expand_sweep and arch_filter, then checks the generated headers.
test_gfx1250_tdm_pipelines.py: the Stream-K golden test now expects comp_tdm and comp_tdm_v2. It also passes when STREAMK_GFX1250_PIPELINES also contains comp_tdm_pc (after feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962).
CHANGELOG entry.

Validation

The main runs were on a gfx1250 A0 board (asicRevision 0, ROCm 10.2, clang bdc35841c97) and the gtests were repeated on a B0 board (asicRevision 1, 96-CU compute partition, ROCm 10.2, clang c7ceccf9a57), building with the Jenkins flags: BUILD_DEV=On, -Werror, -O3, GPU_TARGETS=gfx1250. All builds finished with 0 warnings.

gtests (A0 board):

test_ck_tile_streamk_gfx1250_wmma: 379/379 (288 new + 91 existing).
_smoke: 34/34 (21 new).
_legacy: 91/91.
The new suite with --gtest_repeat=5: 5 x 288 passed, 0 failed.
Other targets that include the edited headers all pass: tdm_subtile 24, tdm_atomic 4, tdm_lds_padding 54, grouped_gemm_tdm 70, multi_d_tdm 360, batched_contraction_tdm 80, mx_tdm_gates 19, mx_splitk_support 5, quant_pipelines_gfx1250 30, grouped_gemm_mx_tdm 8 and grouped_gemm_mx_flatmm_tdm 24. tdm_wmma (144) and mx_tdm_smoke (111) skip on asicRevision 0, the same as on develop.

gtests (B0 board, asicRevision 1, 96-CU compute partition, ROCm 10.2 clang c7ceccf9a57):

test_ck_tile_streamk_gfx1250_wmma: 379/379; --gtest_repeat=3: 1137/1137.
_smoke: 34/34; repeat 3: 102/102. _legacy: 91/91; repeat 3: 273/273.
tdm_wmma (universal comp_tdm V1+V2): 122 passed, 22 skipped (pk_fp4 configs, same behaviour as develop baseline). tdm_subtile 24/24, tdm_atomic 4/4, tdm_lds_padding 54/54, grouped_gemm_tdm 70/70, multi_d_tdm 360/360. Build: rc 0, 0 warnings. Total failures: 0.
Persistent side effect of the comp_tdm flag: I ran the 03_gemm TDM example with -persistent=0/1 for comp_tdm and comp_tdm_v2, fp16 and bf16, plus 1000^3. All runs report "GPU verification result is: correct".

Emulator (Jenkins gfx1250 smoke path):

In the FFM image I ran HSA_MODEL_LIB=/libhsakmtmodel.so, HSA_MODEL_TOPOLOGY=/topology/mi450, ctest -L SMOKE_TEST -R streamk: 1/1 passed in 66 s, with no deadlock. The 21 new atomic tests take about 28 s.
SMOKE_TEST holds only the _smoke target. The linear/tree suites are REGRESSION_TEST.

No existing kernel changed. I compared device-only gfx1250 assembly between develop and this branch:

10 test TUs (199 kernels) are identical: Stream-K full, smoke and legacy; universal comp_tdm V1+V2; TDM subtile, atomic and lds_padding; comp_tdm_pc; grouped TDM; multi-D TDM. After the last kernel change I re-checked the Stream-K, universal comp_tdm, TDM atomic and grouped TDM TUs, and they are still identical.
16 dispatcher-built kernels are identical: compv3/compv4 Stream-K (atomic/linear/tree, persistent), and universal compv4, comp_tdm and comp_tdm_v2.

Dispatcher on the GPU, with #10919 applied (23bf4ae):

120 new Stream-K kernels built: comp_tdm/comp_tdm_v2 x 128/64 tiles x rcr/crr/rrr x (CShuffle atomic/linear/tree + TDM linear/tree) x persistent on/off.
240 bf16 verify runs, tolerance 0.02, on the shapes in the table below:
- 203 verified.
- 36 rejected by the C-tail check, as intended: CShuffle on H with the 128 tile, and CShuffle on M.
- 1 bf16 atomic precision miss on H (max_rel 0.0212), the known atomic class at K=393216.
- 0 wrong results, crashes or hangs.
After the last kernel change I rebuilt a subset (comp_tdm 8-wave, persistent tree, comp_tdm_v2 atomic). It verified 8/8; CShuffle kernels were rejected on ragged N, as intended.
streamk_gemm_full_benchmark.py end to end on the control shape: 20/20 verified.

Performance: bf16, 5 interleaved A/B rounds, warmup 50, repeat 100. The table compares the best new kernel with the best existing Stream-K kernel and the best universal kernel. "Rounds" counts the rounds where the new kernel was faster. The timing comes from the bridge build before the last review fixes; those fixes added only compile-time asserts, host-side checks and comments.

Shape	Best new kernel	TFLOPS	Best existing Stream-K	vs existing (rounds)	Best universal	vs universal (rounds)
H crr 192x448x393216	comp_tdm_v2, TDM epilogue, tree	297.2	compv3 linear, 228.9	1.30x (5/5)	compv4, 30.1	9.87x (5/5)
M crr 118x102x368640	comp_tdm_v2, TDM epilogue, tree	64.3	compv3 tree, 36.8	1.75x (5/5)	comp_tdm_v2, 3.27	19.7x (5/5)
J rcr 2048x32768x9984	comp_tdm_v2, TDM epilogue, tree	1056.6	compv3 linear, 935.3	1.13x (5/5)	comp_tdm_v2, 1102.6	0.958x (0/5)
G rrr 1024x1024x68032	comp_tdm_v2, TDM epilogue, linear	706.1	compv4 linear, 724.8	0.974x (0/5)	compv4, 348.1	2.03x (5/5)
A rcr 3072x512x2048	comp_tdm_v2, CShuffle, atomic	128.1	compv4 atomic, 126.7	1.01x (4/5)	compv4, 316.5	0.405x (0/5)
control rcr 1280x2048x1024	comp_tdm_v2, CShuffle, atomic	104.4	compv4 linear, 98.6	1.06x (4/5)	compv4, 430.5	0.242x (0/5)

comp_tdm (V1) matches V2 on H, A and control, but reaches only 0.74x of V2 on G and 0.83x on J.

CPU Python: I ran every test_*.py under dispatcher/tests and tile_engine. On develop: 2831 passed, 0 failed, 139 skipped. On this branch: 2842 passed, 0 failed, 139 skipped. The 11 extra are the new file. No shared file changed its result.

All C++ changes pass clang-format 18.1.4, and all added text is ASCII.

Limitations

Stream-K with comp_tdm or comp_tdm_v2 supports row-major C only, enforced by a static_assert. CShuffleEpilogue supports only row-major C. TdmEpilogue with column-major C compiles, but returned wrong values (97-98% of them) in my tests, including on data-parallel-only tiles, which run the universal GEMM path.
With the CShuffle epilogue, N (the contiguous C dimension) must be a multiple of the block tile. Padding is not accepted, because TDM pads stay rejected, as for universal TDM and in feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962. Ragged N runs on the TDM epilogue with linear or tree reduction. Atomic reduction on ragged N has no TDM kernel: neither H with the 128 tile nor any of M. H with CShuffle tops out at 238 TFLOPS; with the TDM epilogue it reaches 297.
bf16 atomic reduction at very large K can miss a 0.02 tolerance, as with the existing Stream-K pipelines.
The drain is defensive: with it compiled out, 20 x 288 runs still passed on my board. I keep it because nothing else orders that last prefetch against the next tile's LDS writes.
Side effect of the comp_tdm UsePersistentKernel flag: a universal or grouped kernel built with persistent traits now takes the persistent entry with these pipelines. All current universal, grouped, MX, multi-D and quant TDM users are non-persistent. Their device code is unchanged and their tests pass. Grouped GEMM still rejects a persistent TDM pipeline with its existing static_assert.
On my board the universal comp_tdm/comp_tdm_v2 and MX TDM gtests skip (asicRevision 0), so identical device code is the evidence for them. The universal comp_tdm_pc gtest hangs on its first test on develop as well, and its device code is identical, so it is unrelated to this PR.

Relation to other PRs

feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962 (draft, Stream-K for comp_tdm_pc): independent. I left alone the call sites it rewrites (RunGemm, the linear/tree call, the wavelet RunGemmPipeline) and the Stream-K test util. A trial merge against its head (ab8ab43) gives 4 small textual conflicts:
- streamk_gemm_kernel.hpp: both PRs add a block at the top of IsSupportedArgument. Keep both (ours first) and drop our duplicate device check.
- codegen_common.py: make STREAMK_GFX1250_PIPELINES the union ("comp_tdm", "comp_tdm_v2", "comp_tdm_pc") and merge the comment.
- test_gfx1250_tdm_pipelines.py, Stream-K golden test: combine both bodies, keeping feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962's comp_tdm_pc persistent check.
- gemm_vector_fallback.py: one docstring line; either side works.
With the union resolved our tests pass. Two feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962 tests (test_only_comp_tdm_pc_among_gfx1250_pipelines, test_gfx1250_gate_for_stream_k) assert that comp_tdm and comp_tdm_v2 are rejected and need updating.
feat(ck-tile): enable gfx1250 Stream-K Dispatcher-Bridge #10919 (gfx1250 Stream-K dispatcher bridge): no conflicts. I ran the bridge validation above with it applied.
fix(ck-tile): run misaligned Stream-K problems through the vector-width fallback (gfx1250) #14163 (misaligned Stream-K through the vector-width fallback): one conflict in the unified_gemm_codegen.py STREAM_K loop. Keep our per-strategy loop and add fix(ck-tile): run misaligned Stream-K problems through the vector-width fallback (gfx1250) #14163's streamk_atomic_c_reject_reason check after streamk_trait_reject_reason. Its other conflicts already exist against develop.

Follow-up TODO

Remaining pipelines:

comp_tdm_pc Stream-K: in progress in feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962. Whichever PR lands second makes STREAMK_GFX1250_PIPELINES the union and updates the tests that assert the other pipelines are rejected.
comp_tdm_v3 (fp8/bf8): only in the Tile Engine today, and plain-GEMM-only (kPlainGemmOnly), which StreamKKernel now rejects at compile time. Needs the same Stream-K call form, a Stream-K exception to the plain-GEMM-only rule, tests, and the dispatcher bridge.
comp_async: needs a runtime-tail call adapter (it has no non-tuple runtime-tail overload) and Stream-K pad rules (non-MX comp_async on gfx1250 must be fully padded).

Other follow-ups from this work:

Stream-K grid size / partition choice per shape: a smaller grid gave +21-61% on existing Stream-K kernels.
nopad Stream-K configs for the gfx1250 Tile Engine.
Packed bf16 atomic add (in feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962).
Tile Engine gfx1250 Stream-K TDM configs once feat(ck-tile): enable gfx1250 Stream-K Dispatcher-Bridge #10919 lands. The old Tile Engine Stream-K instance builder offers only mem/compv3/compv4.
After feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962 lands:
- remove the comp_tdm/comp_tdm_v2 6-argument call form if Stream-K calls the tuple form;
- drop our C-tail check in favor of its UniversalGemmKernel check;
- merge our TdmEpilogue/atomic static_assert with its epilogue_can_accumulate_v assert;
- keep the drain or move it into its RunTilePipeline;
- move the new gtest onto its extended test_gemm_streamk_util.hpp fixture.
Atomic reduction on a ragged contiguous C dimension with a TDM pipeline: a padding-aware C rule agreed with feat(ck_tile): Enable Stream-K for Producer Consumer pipeline #14962's check, or a CShuffle C vector-width fallback.
Column-major C with TdmEpilogue gives wrong results (universal TDM GEMM probably too). Fix or reject it, and make the dispatcher reject column-major C for Stream-K TDM; today such a kernel stops at the static_assert.
Performance:
- J reaches 0.958x of universal comp_tdm_v2; candidates are the 128x128 main loop and the per-segment drain.
- G reaches 0.974x of compv4 Stream-K.
- comp_tdm (V1) Stream-K reaches 0.74-0.83x of V2 on G and J.
- A persistent data-parallel section on TDM measured 729-737 TFLOPS against about 1000 non-persistent on J (prototype).

Checklist

Added GPU regression tests and CPU dispatcher tests.
Kept linear/tree tests out of the emulator smoke stage.
Added a release note.
clang-format 18 and ASCII checks on all changed files.

Add the call form StreamKKernel uses (one A and one B window, hot-loop flag, tail number) to comp_tdm and comp_tdm_v2. It forwards to the tuple form, so Stream-K runs the same pipeline entry as universal GEMM. V2 restates it because its operator() overloads hide V1's. comp_tdm now exposes UsePersistentKernel (non-MX problems only, as in comp_async), which StreamKKernel requires to match its partitioner. comp_tdm_v2 inherits it. StreamKKernel: - wait for the last TDM prefetch after a partial tile, which runs no epilogue, before the next tile refills the same LDS; - reject TdmEpilogue with atomic reduction (it overwrites C), the plain-GEMM-only pipelines and TDM cluster launch at compile time; - require CShuffleEpilogue or TdmEpilogue and row-major C for comp_tdm and comp_tdm_v2; - check for a gfx125 device, and with a non-TDM epilogue reject a C whose contiguous dimension is not a multiple of the block tile (the CShuffle store does not clip). Existing Stream-K, universal, grouped and multi-D kernels compile to the same device code.

TestGfx1250StreamKTdm covers 36 configs: both pipelines x atomic, linear and tree reduction x CShuffle and TDM epilogue (TDM with linear/tree only) x persistent and non-persistent, plus fp16, a 64x64 tile, the other A/B layouts and comp_tdm on eight waves. Each config runs data-parallel-only, split-K segments of 1-5 iterations, segments that cross tiles, mixed data-parallel and Stream-K, ragged K, ragged M/N, a full-device grid and a misaligned-pointer rejection. The suite is built into the gfx1250 Stream-K full and smoke targets. The smoke build keeps three atomic configs, because the emulator runs one workgroup at a time and a linear or tree tile owner would wait forever there. The legacy buffer-addressing target does not build it: that mode has no LDS transpose load for the TDM pipelines.

codegen_common.streamk_trait_reject_reason holds the Stream-K trait rules: atomic reduction needs the cshuffle epilogue, linear and tree take cshuffle or tdm, and TDM pipelines keep the no-pad rule. The codegen, expand_sweep and arch_filter all use it together with gfx1250_pipeline_reject_reason(streamk=True), which accepts only STREAMK_GFX1250_PIPELINES (comp_tdm, comp_tdm_v2) and skips the universal "epilogue=tdm" and "no persistent kernel" rules. comp_tdm_pc and comp_async stay rejected with the same reasons as before. The Stream-K codegen now emits the TDM epilogue for linear and tree reduction. Vector-width pairing (_tile_fits) follows the kernel's new C-tail check for a TDM pipeline with a CShuffle epilogue. Add test_streamk_tdm_gfx1250.py: the trait matrix against an independent oracle, the reject reasons, and agreement between codegen, expand_sweep and arch_filter, plus checks on the generated headers.

…gelog

… pipelines Since #15005, UniversalGemmKernel::IsSupportedArgument runs IsCTensorSupported for a TDM pipeline with a non-TDM epilogue: a ragged contiguous C dimension needs kPadN (kPadM), and that dimension must be a multiple of the C vector size. StreamKKernel::IsSupportedArgument calls it after its TDM block, so the Stream-K copy of the C-tail check is redundant for the unpadded TDM configs (the gtest and the dispatcher never pad a TDM pipeline), and its comment, which says the universal kernel skips that check, no longer holds. Drop the copy and keep the gfx125 device check. A padded C goes through the same padded C window as the other Stream-K pipelines.

therock-pr-bot · 2026-11-04T18:25:24Z

✅ 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-11-04T18:34:24Z

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

Copilot

🟡 Changes recommended

The TDM Stream-K gate still admits unsupported column-major layouts and CShuffle wave tuples that can cause compilation failures or output corruption.

4 open issues

Filter column-major C for TDM Stream-K reductions · New
Reject multi-warp-K TDM Stream-K with CShuffle · New
Allow ragged dimensions for TDM Stream-K kernels · New
Use pipeline-aware A/B widths in vector fallback · New

What changed in this PR

Enables gfx1250 CK Tile Stream-K GEMM on comp_tdm and comp_tdm_v2, extending existing native Stream-K support to TDM pipelines for workloads with few output tiles and large K dimensions.

Changes:

Adds pipeline call adapters, persistent-kernel support, synchronization, and kernel compatibility checks.
Shares dispatcher acceptance rules and enables TDM epilogues for linear/tree reduction.
Adds GPU regression and emulator-smoke coverage, CPU dispatcher tests, and a release note.

File	Description
`projects/composablekernel/tile_engine/ops/gemm/gemm_vector_fallback.py`	Adds the contiguous-C tile check for TDM with non-TDM epilogues.
`projects/composablekernel/test/ck_tile/gemm_streamk/test_gemm_streamk_tdm_gfx1250_wmma.cpp`	Adds TDM Stream-K GPU coverage across pipelines, reductions, persistence, and tails.
`projects/composablekernel/test/ck_tile/gemm_streamk/CMakeLists.txt`	Includes the new suite in builtin and smoke targets, excluding legacy addressing.
`projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_tdm_v2.hpp`	Adds V2’s runtime Stream-K call adapter.
`projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_tdm_v1.hpp`	Adds the persistent trait and runtime Stream-K call adapter.
`projects/composablekernel/include/ck_tile/ops/gemm/kernel/streamk_gemm/streamk_gemm_kernel.hpp`	Adds TDM compatibility checks, device validation, and prefetch draining.
`projects/composablekernel/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp`	Removes an outdated comment about TDM persistent traits; no functional change.
`projects/composablekernel/dispatcher/tests/test_streamk_tdm_gfx1250.py`	Adds CPU rule-matrix, codegen-equivalence, pairing, and generated-header tests.
`projects/composablekernel/dispatcher/tests/test_gfx1250_tdm_pipelines.py`	Updates expected Stream-K pipeline coverage.
`projects/composablekernel/dispatcher/tests/CMakeLists.txt`	Registers the new CPU regression suite with CTest.
`projects/composablekernel/dispatcher/python/gemm_utils.py`	Applies shared Stream-K rules during sweep expansion.
`projects/composablekernel/dispatcher/codegen/unified_gemm_codegen.py`	Expands supported reduction/epilogue combinations and emits TDM epilogues.
`projects/composablekernel/dispatcher/codegen/codegen_common.py`	Centralizes Stream-K trait rules and the gfx1250 pipeline allowlist.
`projects/composablekernel/dispatcher/codegen/arch_filter.py`	Applies the shared gfx1250 gate to Stream-K configurations.
`projects/composablekernel/CHANGELOG.md`	Documents the new TDM Stream-K support and restrictions.

_{Review effort: Max}

+    if streamk and pipeline not in STREAMK_GFX1250_PIPELINES:
+        return f"pipeline={pipeline} is not supported for stream_k"


-        if epilogue != "tdm":
+        # StreamKKernel takes a TDM pipeline with either epilogue and persistent or not;
+        # streamk_trait_reject_reason decides the epilogue per reduction strategy.
+        if not streamk and epilogue != "tdm":


+# comp_tdm and comp_tdm_v2 are the gfx1250-only pipelines StreamKKernel has been
+# validated with.
+STREAMK_REDUCTION_STRATEGIES = ("atomic", "linear", "tree")
+STREAMK_GFX1250_PIPELINES = ("comp_tdm", "comp_tdm_v2")


+        if cfg.epilogue == "tdm":
+            return True
+        if cfg.layout[2] == "r":
+            return dims["n"] % cfg.tile_n == 0


ozturkosu added 4 commits November 4, 2026 07:05

docs(ck-tile): note gfx1250 Stream-K on the TDM pipelines in the chan…

7938fa7

…gelog

ozturkosu self-assigned this Nov 4, 2026

github-actions Bot added documentation project: composablekernel labels Nov 4, 2026

ozturkosu added 2 commits November 4, 2026 13:21

Merge branch 'develop' into users/muozturk/ck/streamk-tdm-pipelines

27127fb

assistant-librarian Bot added the organization: ROCm label Nov 4, 2026

ozturkosu requested a max review from Copilot November 4, 2026 19:41

Copilot started reviewing on behalf of ozturkosu November 4, 2026 19:42

View session

Copilot AI reviewed Nov 4, 2026

View reviewed changes

Merge branch 'develop' into users/muozturk/ck/streamk-tdm-pipelines

e295da3

ozturkosu mentioned this pull request Nov 5, 2026

feat(ck-tile): enable gfx1250 Stream-K on the comp_tdm_v3 pipeline #15118

Draft

ozturkosu added this pull request to stack #15120 November 5, 2026 05:36

Navigation Menu

Repository navigation

feat(ck-tile): enable gfx1250 Stream-K on the comp_tdm and comp_tdm_v2 pipelines - #15083

feat(ck-tile): enable gfx1250 Stream-K on the comp_tdm and comp_tdm_v2 pipelines#15083
ozturkosu wants to merge 7 commits into
developfrom
users/muozturk/ck/streamk-tdm-pipelines

ozturkosu commented Nov 4, 2026

Uh oh!

therock-pr-bot Bot commented Nov 4, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Nov 4, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

		if streamk and pipeline not in STREAMK_GFX1250_PIPELINES:
		return f"pipeline={pipeline} is not supported for stream_k"

Conversation

ozturkosu commented Nov 4, 2026

Motivation

Changes

Pipelines (comp_tdm, comp_tdm_v2)

Kernel (StreamKKernel)

Dispatcher

Tests

Validation

Limitations

Relation to other PRs

Follow-up TODO

Checklist

Uh oh!

therock-pr-bot Bot commented Nov 4, 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 Nov 4, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Kernel (`StreamKKernel`)

therock-pr-bot Bot commented Nov 4, 2026 •

edited

Loading