Repository navigation
feat(ck-tile): comp_tdm_pc producer/consumer GEMM TE to dispatcher bridge (gfx1250) - #13975
Merged
illsilin merged 41 commits intoNov 2, 2026
Merged
Conversation
…ings No consumers yet and never run on hardware; the gfx1250 barrier ABI it assumes is from a compile probe, not documentation. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Never executed: no gfx1250 hardware was available, so it is compile- and ISA-verified only. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
The store reads LDS asynchronously, so a caller that reuses the buffer, such as a persistent kernel's next tile, could corrupt C. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
They leave grouped_convolution_utils.hpp so GEMM kernels can use them too; the grouped-conv names stay, so no behaviour changes. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
run() owns the priming, the ragged tail and the final waits; init() now compiles away off gfx1250 so kernels can arm unconditionally. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Records the handshakes run() issues and checks every wave's log against a model of the barrier semantics under many interleavings. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Adds a two-producer, two-consumer fan-out test and pins down the signal-without-join semantics every handshake rests on. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
…mmKernel Launches LaunchBlockSize threads, arms the pipeline's barriers and runs the epilogue on the math waves only. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Loader waves stream A and B into a ring of LDS slots that the compute waves drain, synchronised over gfx1250 hardware named barriers. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Replaces the stringified TEST_SUITE_NAME comparisons in the shared GEMM pipeline cases with a TestFixture::IsTdm trait. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Runs every layout, a persistent case that reuses the barriers across tiles, and an fp8 case whose K block spans several WMMA K steps. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
…r ring PublishLag matches the pipeline policy; arm() uses static_ford, for_each_step a while loop; the group-count assert says why it stays. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Assert per element, use get_warp_id()/get_lane_id(), pin the hand-unrolled loops to two slots, and follow the PublishLag rename. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
The ring must end each tile at rest for persistent tiles to reuse it without re-arming; also follow the PublishLag rename. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
…er tile InitBarriers() arms them before the persistent tile loop; RunGemm takes its token, which only named-barrier pipelines must pass. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Reject padding, which would carry TDM transfers past the tensor; match the B loader's wave explicitly; rename k_is_dim0 to k_is_outer. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Append the label after the relabel loop so the test keeps SMOKE_TEST, the only label gfx1250 CI runs. Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
Without them a benchmark pass hung 70 shapes. Also caps barrier groups at 4, the real per-work-group limit. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…t-out Also repairs the GetOutputOffset detector, which was silently always false. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…e limits Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The relabel loop replaces LABELS rather than appending, so the test was built but never run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
A positional split reported pipeline "comp" for every multi-token name and shifted later fields. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Budgets LDS by slot count, bounds tiles by accumulator VGPRs, and corrects a stale comp_async count pin. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
…idge (gfx1250) Route the gfx1250 producer/consumer TDM pipeline through the dispatcher: Pipeline::CompTDMPC enum, codegen type maps, and the old-TE acceptance rules in the shared gfx1250 gate (tdm epilogue, intrawave, no pads, no persistent, accumulator cap of 256 VGPRs per lane, quarter-LDS budget for the four slots). Batched comp_tdm_pc is rejected because BatchedGemmKernel does not arm the producer/consumer barriers.
Resolve conflicts with the landed producer/consumer pipeline while preserving the bridge's split-K rejection, K-wave validation, and transpose-aware LDS accounting. Keep a single producer/consumer gtest registration with its gfx1250 filter, smoke label, and timeout. Validation: 129 tests passed, 7 skipped, and 433 subtests passed across the focused dispatcher/tile-engine suites; full CK clang-format-18 (18.1.3), ASCII, and CRLF checks passed; configure-only CMake registration checks passed for mixed, skipped, and non-gfx1250 configurations. No local HIP/device build or execution.
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Unresolved correctness issues and shared hardware-synchronization changes require fixes and human review alongside the pending gfx1250 bridge device-parity results.
6 open issues
Review effort: Balanced
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Shared kernel validation and named-barrier pipeline integration need final human review with the pending gfx1250 on-device bridge parity results.
6 open issues
Review effort: Balanced
The standalone SplitKRejectedForFloatOutput test does not run the typed fixture's SetUp(), so on a non-gfx1250 device IsSupportedArgument rejects the positive control at its named-barrier check and the test fails while the typed cases skip. Move the fixture's gate into CK_TILE_SKIP_UNLESS_GFX125() and use it in both places, keeping the CK_TILE_REQUIRE_GFX125=1 failure behavior.
ozturkosu
requested review from
andriy-ca and
illsilin
and
a balanced review from Copilot
October 22, 2026 18:06
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
End-to-end gfx1250 bridge correctness and parity with the existing Tile Engine remain pending and require device validation and human review.
No open issues 🎉
6 resolved since last review
Review effort: Balanced
1 task done
…comp-tdm-pc-dispatcher-bridge Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 53fc044c-7bc8-41f2-b5ae-b03bfde3b135
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 53fc044c-7bc8-41f2-b5ae-b03bfde3b135
illsilin
added a commit
that referenced
this pull request
Oct 30, 2026
…a split barrier (#14466) JIRA ID : AICK-2347 ## Summary A fast fp8 GEMM for gfx1250 in CK Tile. This PR adds `comp_tdm_v3`, an opt-in GEMM pipeline: a 16-wave tile fed by the Tensor Data Mover (TDM), with the per-tile workgroup barrier split into arrive and wait so the WMMAs overlap the wait. At the lead shape, 4096 x 4096 x 65536 (RCR, fp8 e4m3 in, fp32 accumulate, bf16 out), it matches the fastest known gfx1250 fp8 GEMM. Everything is opt-in: existing kernels compile byte-for-byte the same. ## What's new - **`comp_tdm_v3` pipeline** (gemm): 16 waves with a waves-aware register budget; one split barrier per K tile, with B held in registers and A streamed per row slice; the A row windows built outside the K loop; TENSORcnt drained on exit. The joint barrier stays available (`kUseSplitBarrier = false`), and an early-refill policy (`RefillPoint::AfterArrive`) is available opt-in; it brings no gain at the lead shape, so it is off by default. - **Split workgroup barrier** (core): fenced, pinned `workgroup_barrier_arrive` / `_wait`. - **tile_engine**: `comp_tdm_v3` registered for `gemm_universal` (`COMPUTE_TDM_V3`, `[4,4,1]` and CShuffle for this pipeline), the fp8/bf8 32x32x128 warp tile on gfx1250, opt-in bf16 C for fp8/bf8, `-init 3` (normal) and `-init 4` (integers), and `-verify_exact`. Its targets build with the post-RA scheduler. - **Example** `53_gemm_comp_tdm_v3_fp8`: a timed benchmark and verifier for the lead shape. ## Performance Ratio to the fastest known gfx1250 fp8 GEMM at the lead shape (the rocKE kernel of rocKE PR #14707, with the split barrier and alternating TDM issuers), random data, builds interleaved sample by sample in one session on one gfx1250 host: | | Peak | Sustained | |---|---|---| | `comp_tdm_v3`, default policy | 1.008x | 1.001x | Differences of about 1% or less are within session noise, so this is parity. Every timed build was bit-exact at 2048³ and passed integer and random checks at the lead shape first. ## Scope of `comp_tdm_v3` Supported and tested: - fp8 or bf8 A and B, including mixed fp8 x bf8; bf16 C; - A row-major, B column-major, C row-major (RCR); - 16 waves (256x256x256) and 4 waves (128x128x256), warp tile 32x32x128; - the plain GEMM kernel with the CShuffle epilogue; - `k_batch = 1`. Rejected: - **Split-K.** `IsSupportedArgument` rejects `k_batch != 1`. - **Non-RCR layouts.** The pipeline has a `static_assert`, and tile_engine rejects them. The other A/B layouts spill VGPRs in the K loop at the 16-wave tile. - **Other kernels.** The multi-D, grouped and batched-contraction kernels have a `static_assert` against the pipeline's new `kPlainGemmOnly` marker. - **tile_engine.** It offers the pipeline for fp8/bf8 and `gemm_universal` only. Not in scope: bf16/fp16 A/B (they compile but are not tested), LDS padding, other C types, MX GEMM and group quant. `comp_tdm_v2` remains the TDM path for MX GEMM and group quant. ## Default behaviour changes - The new pipeline and every tile_engine rule are opt-in. No default or CI config lists `comp_tdm_v3`. - Existing kernels are byte-identical before and after this PR (gfx1250, plus gfx942 and gfx950 cases), on TheRock 10.1.0 and 10.2.0. - **One host-side change.** `UniversalGemmKernel::IsSupportedArgument` used to return `true` for every TDM pipeline right after the general k_batch checks. It now returns `IsTdmArgumentSupported(...)`, the check the multi-D, grouped, batched-contraction and MX kernels already apply: `k_batch == 1`, `K > 0`, element-aligned A/B/E. - For comp_tdm / comp_tdm_v2 / comp_tdm_pc in the plain GEMM kernel, split-K was rejected only for fp16/bf16 C. With fp32 C, k_batch > 1 passed the generic split-K check, although the TDM epilogue's tensor store overwrites the output tile; it is now rejected. A/B/E pointers must also be element-aligned, which TDM addressing requires. No test used k_batch != 1 or misaligned pointers with a TDM pipeline. - Device code is unchanged: on ROCm 10.1.0, the device ISA is identical for 235 kernels (the TDM tests, compv3 and example 53). ## Testing On TheRock 10.1.0 (CK CI compiler) and 10.2.0, gfx1250 under the rocJITsu emulator; no hardware CI exists for gfx1250 yet: - gtests for the split barrier and `comp_tdm_v3` (both barrier forms, 4x4 and 2x2 wave maps, alternating and fixed issuers, long K), bit-exact. - The example at 256 x 256 x {1024, 65536} and odd K-tile counts: bit-exact on integers, within tolerance on random data. - ISA checks: 0 in-loop scratch, VGPRs <= 256, the expected barrier and wait placement; 13 of 13 planted sync bugs caught. - Other GPU targets still parse the headers (gfx942, gfx950, gfx1201). On gfx1250 hardware, at the head of this PR: - `comp_tdm_v3` gtests: 55 of 55 short and 16 of 16 long-K passed, including bf8, mixed fp8/bf8 and the split-K rejection check. - The scratch-use check passed: 9 kernels, no scratch, no VGPR spills. Compile-only checks (ROCm 10.1.0, gfx1250): - `comp_tdm_v3` with RRR, CRR, CCR and RCC each fails on the layout `static_assert`; RCR compiles. - `comp_tdm_v3` in the multi-D, grouped and batched-contraction kernels each fails on the `kPlainGemmOnly` `static_assert`; the plain GEMM kernel compiles. - The existing multi-D, grouped and batched-contraction TDM gtests still compile, and the `comp_tdm_v3` gtests build with `-Werror -Weverything`. ## Notes - **Merge.** Squash and merge, per rocm-libraries convention. The commits form a reviewable stack; each message has a `Gate:` paragraph with its evidence. - **Not in this PR.** Dispatcher registration of `comp_tdm_v3` (and bf16 C for fp8 there) conflicts with open #13975; it follows after #13975 lands. - **Credit.** The split-barrier K-loop, the B-resident / A-streamed order and the post-RA scheduler finding are Jeffrey Yang's, from the rocKE gfx1250 fp8 GEMM (rocKE PR #14707). He co-authored several commits and authored the scope gates. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Co-authored-by: Jeffrey Yang <jeffreyj.yang@amd.com> Co-authored-by: Illia Silin <Illia.Silin@amd.com> Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
…comp-tdm-pc-dispatcher-bridge
illsilin
approved these changes
Nov 2, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Nov 4, 2026
feat(ck-tile): comp_tdm_pc producer/consumer GEMM TE to dispatcher bridge (gfx1250) (#13975) > Stacked on #12524 (producer/consumer TDM pipeline). Please review/merge that one first; this PR only adds the dispatcher/bridge side. ## Summary This routes the gfx1250 producer/consumer TDM GEMM pipeline (`comp_tdm_pc`, added to old-TE in #12524) through the TE to dispatcher bridge. The **Dispatcher is the single source of truth for codegen -> build -> runtime**, while the **Tile Engine keeps only the config search space and the benchmark loop**. After this, `comp_tdm_pc` kernels can be generated, filtered and registered by the dispatcher with the same acceptance rules old-TE applies. ## Why comp_tdm_pc needs special handling - **Four LDS slots.** The pipeline rotates `NumSlots = 4` LDS buffers, so the A+B staging budget is a quarter of the hardware LDS (gfx1250: 80 KB), not the half that comp_async / comp_tdm / comp_tdm_v2 share. It gets its own `pipeline_lds_budget` entry instead of the comp_async alias. - **Accumulator cap.** Producer waves eat into the register file, so old-TE rejects tiles with `ceil(tile_m * tile_n / (waves * 32)) > 256` accumulator VGPRs per lane. I mirror that rule in the shared gfx1250 gate, so codegen, the arch filter and `gemm_utils.expand_sweep` all agree. - **No batched variant.** `BatchedGemmKernel` does not arm the producer/consumer barriers (old-TE batched CMake also refuses it), so batched `comp_tdm_pc` is rejected with an explicit reason instead of generating a kernel that hangs. - Everything else matches the other gfx1250 TDM pipelines: gfx1250 only, tdm epilogue, intrawave only, no padding, no persistent, fp32 rcr only, `TransposeC=false`, `DoubleSmemBuffer=true`. ## Design note No launch change is needed. The kernel derives its own `kBlockSize` (producer + consumer waves) and arms its barriers inside `UniversalGemmKernel::operator()`, and the dispatcher already launches with `GemmKernel::BlockSize()`. So the dispatcher-side work is the pipeline enum, the type maps, and the acceptance rules. All of them live in the existing shared gfx1250 gate (`gfx1250_pipeline_reject_reason`) and `arch_specs.json`, not in per-call-site copies. ## Changes **Modified** - `codegen/arch_specs.json`: `comp_tdm_pc` LDS budget (fraction 0.25) and interwave trait rejects. - `codegen/generate_arch_specs.py`: `CompTDMPC` enum mapping; included in the C++ trait-unsupported check. - `codegen/arch_specs_generated.py`, `include/ck_tile/dispatcher/arch_specs_generated.hpp`: regenerated. This also collapses a duplicate gfx1250 `fp32_fp32_fp32` warp-tile key in the JSON; the value is identical. - `include/ck_tile/dispatcher/kernel_key.hpp`: `Pipeline::CompTDMPC` with `to_string` / `string_to_pipeline`. - `codegen/codegen_common.py`: type maps (`GemmPipelineAgBgCrCompTDMProducerConsumer`, `BaseGemmPipelineAgBgCrCompTDM`), `TDM_PIPELINES`, accumulator-cap and batched rejects in the shared gfx1250 gate. - `codegen/arch_filter.py`: trait rejects, fallback LDS budget, passes the tile M/N to the shared gate. - `codegen/unified_gemm_codegen.py`: `DoubleSmemBuffer=true` for comp_tdm_pc; passes batched / tile M/N to the gate. - `python/gemm_utils.py`: same gate inputs for `expand_sweep`. **Tests** - `tests/test_gfx1250_tdm_pipelines.py`: accumulator cap, four-slot LDS budget per arch, type maps, generated header contents, batched exclusion, layout coverage, `gemm_utils` rules, and a C++ probe of the enum / LDS / trait helpers. - `tests/test_gfx1250_gate_equivalence.py`: comp_tdm_pc added to the codegen / arch filter / gemm_utils equivalence sweep (0 mismatches against old-TE's rule). - `tests/test_lds_capacity_arch_aware.py`: deliberate quarter-capacity entry for comp_tdm_pc. ## Testing - Dispatcher Python tests (`dispatcher/tests`): all green except the GPU-only quant correctness tests, which fail identically on the #12524 base commit on a node without HIP headers (environment, not this change). - Old-TE gfx1250 tests (`test_gfx1250_pipelines.py`, `gemm_universal/test_gemm_universal_gfx1250.py`): 56 passed. - gfx1250 on-device bridge vs old-TE parity: TODO. Related bridge PRs: #8997 (GEMM), #9000 (grouped), #9028 (stream-K), #8887 (fp8/bf8/int8).
4 of 15 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
This routes the gfx1250 producer/consumer TDM GEMM pipeline (
comp_tdm_pc, added to old-TE in #12524) through the TE to dispatcher bridge. The Dispatcher is the single source of truth for codegen -> build -> runtime, while the Tile Engine keeps only the config search space and the benchmark loop. After this,comp_tdm_pckernels can be generated, filtered and registered by the dispatcher with the same acceptance rules old-TE applies.Why comp_tdm_pc needs special handling
NumSlots = 4LDS buffers, so the A+B staging budget is a quarter of the hardware LDS (gfx1250: 80 KB), not the half that comp_async / comp_tdm / comp_tdm_v2 share. It gets its ownpipeline_lds_budgetentry instead of the comp_async alias.ceil(tile_m * tile_n / (waves * 32)) > 256accumulator VGPRs per lane. I mirror that rule in the shared gfx1250 gate, so codegen, the arch filter andgemm_utils.expand_sweepall agree.BatchedGemmKerneldoes not arm the producer/consumer barriers (old-TE batched CMake also refuses it), so batchedcomp_tdm_pcis rejected with an explicit reason instead of generating a kernel that hangs.TransposeC=false,DoubleSmemBuffer=true.Design note
No launch change is needed. The kernel derives its own
kBlockSize(producer + consumer waves) and arms its barriers insideUniversalGemmKernel::operator(), and the dispatcher already launches withGemmKernel::BlockSize(). So the dispatcher-side work is the pipeline enum, the type maps, and the acceptance rules. All of them live in the existing shared gfx1250 gate (gfx1250_pipeline_reject_reason) andarch_specs.json, not in per-call-site copies.Changes
Modified
codegen/arch_specs.json:comp_tdm_pcLDS budget (fraction 0.25) and interwave trait rejects.codegen/generate_arch_specs.py:CompTDMPCenum mapping; included in the C++ trait-unsupported check.codegen/arch_specs_generated.py,include/ck_tile/dispatcher/arch_specs_generated.hpp: regenerated. This also collapses a duplicate gfx1250fp32_fp32_fp32warp-tile key in the JSON; the value is identical.include/ck_tile/dispatcher/kernel_key.hpp:Pipeline::CompTDMPCwithto_string/string_to_pipeline.codegen/codegen_common.py: type maps (GemmPipelineAgBgCrCompTDMProducerConsumer,BaseGemmPipelineAgBgCrCompTDM),TDM_PIPELINES, accumulator-cap and batched rejects in the shared gfx1250 gate.codegen/arch_filter.py: trait rejects, fallback LDS budget, passes the tile M/N to the shared gate.codegen/unified_gemm_codegen.py:DoubleSmemBuffer=truefor comp_tdm_pc; passes batched / tile M/N to the gate.python/gemm_utils.py: same gate inputs forexpand_sweep.Tests
tests/test_gfx1250_tdm_pipelines.py: accumulator cap, four-slot LDS budget per arch, type maps, generated header contents, batched exclusion, layout coverage,gemm_utilsrules, and a C++ probe of the enum / LDS / trait helpers.tests/test_gfx1250_gate_equivalence.py: comp_tdm_pc added to the codegen / arch filter / gemm_utils equivalence sweep (0 mismatches against old-TE's rule).tests/test_lds_capacity_arch_aware.py: deliberate quarter-capacity entry for comp_tdm_pc.Testing
dispatcher/tests): all green except the GPU-only quant correctness tests, which fail identically on the feat(ck-tile): producer consumer pipeline #12524 base commit on a node without HIP headers (environment, not this change).test_gfx1250_pipelines.py,gemm_universal/test_gemm_universal_gfx1250.py): 56 passed.Related bridge PRs: #8997 (GEMM), #9000 (grouped), #9028 (stream-K), #8887 (fp8/bf8/int8).