feat(ck-tile): TE to dispatcher GEMM bridge (fp16/bf16, all layouts) by ozturkosu · Pull Request #8997 · ROCm/rocm-libraries

ozturkosu · 2026-07-01T05:11:08Z

JIRA ID : AICK-2083

Re-opened from #8479 with a compliant branch name (users/muozturk/ck-tile/gemm-bridge-all-layout-bf16-fp16). Supersedes #8479.

Summary

This PR routes the Tile Engine (TE) regular-GEMM sweep through the Dispatcher,
making the Dispatcher the single source of truth for codegen → build → runtime
while the Tile Engine keeps only the config search space and the benchmark
loop. It is the consolidated, single-commit GEMM bridge covering all four
layouts (rcr/rrr/crr/ccr) and both fp16 and bf16.

It is a clean re-roll of the earlier bridge work (previously split across
#8123 + the stacked key/bf16/layouts/parity/example PRs and consolidated in
#8261). Those branches accumulated unrelated cross-project commits through repeated
develop merges; this branch is a single clean commit off the latest develop
containing only the GEMM-bridge files. It supersedes and replaces #8123 / #8261.

Motivation

The Tile Engine historically owned its own codegen/build/runtime for GEMM
(tile_engine/ops/gemm/gemm_universal/). The consolidation goal is for the
Dispatcher to own all of that — exactly as it already does for FMHA and
Grouped Conv — so there is one kernel-generation/build/runtime path and the
TE shrinks to a config+benchmark frontend. This PR brings regular GEMM in line
with that reference binding.

The binding (mirrors the FMHA/Conv reference, six stages)

Config JSON (TE side) — the sweep search space lives in
tile_engine/ops/gemm/configs/ (flat op-root layout, matching the
fmha/ and grouped_conv/ bridges).
Codegen (Dispatcher) — dispatcher/codegen/unified_gemm_codegen.py emits
one fully-typed .hpp per kernel; GemmKernelConfig.name reproduces
KERNEL_NAME byte-for-byte (the thread tying config → kernel → runtime).
Compile to .so — a single static gemm_ctypes_lib.cpp is force-included
(-include <kernel.hpp>); one .so per kernel.
Flat extern "C" ABI — dispatcher_run_gemm(A, B, C, M, N, K, time_ms) +
the kernel-name enumeration entry points. Host-pointer memory model (the C
lib hipMallocs internally) — the FMHA-forward branch of the reference.
Python ctypes wrapper — dispatcher/python/gemm_utils.py
(GemmDispatcherLib + GpuGemmRunner).
TE driver (3 phases) — gemm_full_benchmark.py (parallel codegen+build →
expand_sweep → subprocess-isolated benchmark) + the disposable per-kernel
worker run_one_gemm_kernel.py.

What's included

Bridge core

dispatcher/codegen/unified_gemm_codegen.py — GEMM codegen, byte-exact naming.
dispatcher/bindings/ctypes/gemm_ctypes_lib.cpp — flat C ABI, host-pointer model.
dispatcher/python/gemm_utils.py — GemmKernelConfig, multi-kernel build
(setup_multiple_gemm_dispatchers), expand_sweep, one-.so-per-kernel.
tile_engine/ops/gemm/gemm_full_benchmark.py + run_one_gemm_kernel.py —
3-phase, multi-GPU, subprocess-isolated driver/worker.

Feature surface (the point of this PR)

All four layouts rcr/rrr/crr/ccr (row-major C only — ck_tile rejects
column-major C at build) with layout-aware host transpose.
fp16 + bf16 (bf16 via uint16 byte-encoding; dtype derived from kernel name).
Trait-derived registry KernelKey — replaces the earlier hard-coded
fp16/rcr key so the registry path generalizes across dtype/layout/tile.

Correctness & performance hygiene

--verify opt-in fp32 numpy-reference gate (global max|out-ref|/max|ref|),
verified/max_rel columns in the CSV; a mismatch counts as a failure.
Tile Engine AMDGPU -mllvm codegen-flag parity (without these the kernel
builds with different occupancy and the timing diverges) and
arch-validated tile filtering against the real pipeline/scheduler.
Multi-GPU fan-out across all visible GPUs (--devices, device-pinned
HIP_VISIBLE_DEVICES workers).

Example & tests

dispatcher/examples/gemm/python/12_te_bridge.py — runnable end-to-end example.
dispatcher/tests/test_gemm_parity.py, test_gemm_utils.py, and a parity
regression harness.

Cleanup

Removes the legacy standalone gemm_universal build path
(gemm_universal_instance_builder.py, *_benchmark*.{py,cpp,hpp},
gemm_universal/CMakeLists.txt) and the old test/ck_tile/gemm_tile_engine/
harness; promotes the sweep configs to the flat op-root configs/.

Design decisions (consistent with the reference)

Host-pointer memory ownership (C lib owns device memory) — matches
FMHA-forward; the Python runner passes host numpy arrays straight through.
One .so per kernel — packaging choice; the multi-kernel name ABI is
retained (get_kernel_name_at(0) reports the single kernel), so the Python
enumeration path is unchanged from FMHA/Conv.
Flat configs/ at the op root — matches the fmha//grouped_conv/
convention; the not-yet-bridged variants keep their per-variant configs/
dirs, selected by --variant.

Validation (gfx942 / MI300X)

Bridge build + benchmark + --verify across fp16 and bf16 and all
four layouts, checked against an fp32 numpy reference (A @ B).
Name parity holds end-to-end: each .so's reported runtime name equals
GemmKernelConfig(...).name.
bf16 passes under a widened fp16/bf16 tolerance; fp16 within the standard
max_rel gate.

Test plan

gemm_full_benchmark.py --verify over configs/default_ci_config.json for
fp16 and bf16, each of rcr/rrr/crr/ccr.
unified_gemm_codegen.py emits a header whose stem == GemmKernelConfig.name.
setup_multiple_gemm_dispatchers builds + links each config against
gemm_ctypes_lib.cpp.
pytest dispatcher/tests/test_gemm_parity.py dispatcher/tests/test_gemm_utils.py.
examples/gemm/python/12_te_bridge.py runs end to end.

Notes

Single clean commit off the latest develop; the diff is 35 files, all under
projects/composablekernel/ (dispatcher + tile_engine/ops/gemm + test/ck_tile).
Supersedes [CK_TILE] Add Tile Engine -> Dispatcher bridge for GEMM #8123 and [CK_TILE] TE->Dispatcher GEMM bridge (consolidated: key + bf16 + layouts + parity + example) #8261, which will be closed.
Stream-K ([CK_TILE] Add Tile Engine -> Dispatcher bridge for Stream-K GEMM #8136) and grouped GEMM are separate bridge efforts, not in this PR.

Consolidated, single-commit GEMM bridge routing the Tile Engine regular-GEMM sweep through the Dispatcher (codegen -> build -> runtime), so the Dispatcher is the single source of truth and the Tile Engine owns only the config search space and the benchmark loop. Mirrors the FMHA/Conv reference binding end to end. Scope: - Regular GEMM bridge: unified_gemm_codegen.py, gemm_ctypes_lib.cpp (flat extern "C" ABI, host-pointer model), gemm_utils.py (GemmKernelConfig with byte-exact .name, one-.so-per-kernel build), 3-phase TE driver + subprocess worker (gemm_full_benchmark.py / run_one_gemm_kernel.py). - Trait-derived registry KernelKey (replaces the hard-coded fp16/rcr key). - bf16 support and all four layouts (rcr/rrr/crr/ccr; row-major C only). - Tile Engine AMDGPU -mllvm codegen-flag parity + arch-validated tile filtering. - --verify fp32-reference correctness gate; multi-GPU fan-out. - Runnable example (examples/gemm/python/12_te_bridge.py) and parity/unit tests. - Removes the legacy standalone gemm_universal build path and the old test/ck_tile/gemm_tile_engine harness; promotes sweep configs to the op-root flat configs/ directory (fmha/grouped_conv convention). Validated on gfx942 / MI300X (fp16 + bf16, all four layouts) against an fp32 numpy reference via --verify.

The bridge dispatcher's tile-divisibility gate rejected any problem where M % TileM != 0 for every layout, returning status -2 ("No suitable kernel") at runtime even though the .so built fine. This wrongly excluded bf16 rcr/rrr kernels with a non-power-of-two TileM (e.g. 192) on standard shapes like 1024^3 -- cases Old-TE compiles, runs, and verifies as correct. Root cause: supports() was layout-blind, while the underlying ck_tile::GemmKernel::IsSupportedArgument only constrains a dimension when an operand whose inner axis is that dimension participates without padding: RowMajor A -> K, ColMajor A -> M RowMajor B -> N, ColMajor B -> K RowMajor C -> N, ColMajor C -> M So for rcr (RowMajor A & C) M is never gated, which is why Old-TE runs M=192 tiles on M-indivisible problems. Make supports() compute require_m/n/k from the kernel key's A/B/C layouts so it mirrors IsSupportedArgument exactly (also honoring k_batch in the K grain). Anything it now lets through is still validated by the kernel's own IsSupportedArgument inside launch(), so the bridge stays a strict functional equivalent of Old-TE. Applied to both generated_tile_backend.hpp (the GEMM .so path) and the sibling tile_backend.hpp. Validated on gfx942 (MI300X): 85 previously status-2 rcr/rrr bf16 192-tile .so now run at 1024^3 (Old-TE runs the same, verification correct); the 8 remaining rejects are tile N=192 cases that Old-TE also reports "Arguments not supported" at N=1024 -- parity preserved in both directions.

…oding rcr dispatcher_initialize() in gemm_ctypes_lib.cpp hardcoded the KernelKey layout to rcr (RowMajor/ColMajor/RowMajor) for every kernel. Now that supports() is layout-aware, that wrong key layout makes the dispatcher reject valid problems: a crr kernel does not gate K (neither A=ColMajor nor B=RowMajor has K as its inner axis), but with a hardcoded rcr key supports() applies rcr's K-gate and returns status -2 for TileK=192 problems (e.g. crr 64x64x192 at 1024^3) that Old-TE compiles, runs, and verifies (~87 TFLOPS). Derive signature.layout_a/b/c from the force-included kernel's own ALayout/BLayout/CLayout types via std::is_same_v with tensor_layout::gemm::RowMajor. The key now matches the kernel, so the layout-aware gate is correct for all four layouts. Execution was already layout-correct (the kernel uses its own compile-time layouts); only the host-side selection metadata was wrong. Validated on gfx942 (MI300X): crr 64x64x192 now runs on the bridge (93 TFLOPS), restoring parity with Old-TE.

The >=20% bridge-vs-old-TE perf gaps in the parity sweep are a harness artifact: the sweep timed the bridge in-process but timed old-TE via its separate standalone benchmark binary, which runs the byte-identical kernel at a lower sustained SCLK. Measured through one harness the gap is <1%. ab_same_harness.py removed that artifact but hardcoded the old-TE header dir to fp16/rcr. Derive it per stem as <base>/<dtype>/<layout> so one run covers rcr/rrr/ccr/crr and fp16+bf16, add a --stems-file/--csv resume-aware sweep mode, and use the median (not max) per point.

For a full ~2000-stem sweep on a single GPU: batch all shapes into one worker call per side (5x fewer process startups), cache the compiled old-TE .so, and add a parallel --build-only pre-pass so hipcc compilation uses all CPU cores while GPU measurement stays serial.

… guard) The bridge-vs-old-TE A/B reported phantom regressions from two MEASUREMENT bugs, not real codegen gaps: - ab_same_harness.py built the old-TE side WITHOUT the TE codegen flags the bridge (and real old-TE's own CMake) use, so -enable-post-misched defaulted back on and old-TE ran ~10-40% faster -> the bridge looked regressed when it is at parity. Now both sides build with identical flags. - ab_efficient_sweep.py measured whatever libgemm_<stem>.so existed with no freshness check, so 3-day-old binaries built from an obsolete codegen showed up as -78%/+703% gaps. Added a guard: skip any .so older than its generated header (treated as missing) instead of reporting a phantom gap. With both fixes the 41 former >15% outlier stems measure within +/-10% (median +0.01%); no bridge codegen regression exists. Note: a separate, deliberately UNCOMMITTED perf change in gemm_utils.py (gate -enable-post-misched=0 on persistent) gives non-persistent large tiles ~9-40%; held back pending a broader persistent-kernel no-regression sweep.

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

The bridge compiles each kernel .so with a hand-maintained hipcc flag list (dispatcher/python/gemm_utils.py) that had drifted from Tile Engine's CMake flags, so the bridge .so and the TE benchmark were not compiled apples-to-apples: * MISSING -mllvm -amdgpu-coerce-illegal-types=1 (TE's CMakeLists.txt adds it when the compiler accepts it; the bridge build never did) * EXTRA -mllvm -enable-noalias-to-md-conversion=0 (not a TE GEMM flag; it only appears in standalone CK examples/tests, never the TE gemm path) Align the bridge's backend codegen flags with the exact set the TE gemm_universal benchmark TU is built with. The coerce flag is added through a cached hipcc probe that mirrors TE's check_cxx_compiler_flag, so the bridge stays matched to TE on every toolchain (present where TE has it, skipped where TE's CMake would skip it too). The generated kernel source was already identical between the two engines; this makes their compilation identical as well.

…y_diag/regression Old-TE must remain until the dispatcher bridge implements every datatype Old-TE supports, so revert the Old-TE removal from the bridge commit and re-wire its build: * restore test/ck_tile/gemm_tile_engine/* (10 files) * restore tile_engine/ops/gemm/gemm_universal/* (6 files: benchmark / instance builder / profiler / single-bench / CMakeLists) * re-add `add_subdirectory(gemm_universal EXCLUDE_FROM_ALL)` in tile_engine/ops/gemm/CMakeLists.txt; restore test/ck_tile/CMakeLists.txt to the develop state (gemm_tile_engine entry kept commented, as in develop) Also drop the parity_diag/regression dev scripts that should not ship in the PR: * dispatcher/parity_diag/regression/ab_efficient_sweep.py * dispatcher/parity_diag/regression/ab_same_harness.py

…whitespace - Add AMD copyright/SPDX header to gemm_full_benchmark.py and run_one_gemm_kernel.py (CK requires a header on every source file). - Remove a trailing-whitespace blank line in generated_tile_backend.hpp that would trip the whitespace/clang-format CI gate.

….g. 192) The CShuffle epilogue stores the accumulator back through LDS in power-of-two MRepeat/NRepeat chunks, where MRepeat = tile_m / (wave_m * warp_tile_m) (and likewise N). A tile whose per-wave repeat is not a power of two (or whose tile dim is not divisible by wave*warp_tile) is mis-stored and produces numerically WRONG results at runtime -- yet it still passes the ctypes validator and the epilogue's static_asserts, so it compiles and silently returns garbage. Observed on MI350 for tile_m=192 (MRepeat = 192/(2*32) = 3) and tile_n=192 (e.g. 64x192x64_1x4x1, 192 not divisible by 4*32): both verified incorrect (fp32 reference, max_rel ~1.2-1.4) on the bridge AND Tile Engine, at every shape including shapes divisible by 192. Power-of-two tiles (64/128/256) are unaffected; a control 256-tile verifies cleanly (max_rel ~4e-4). Add a validity gate in both tile-expansion paths: * unified_gemm_codegen.py::_get_tile_configs (codegen CLI path) * gemm_utils.py::expand_sweep (bridge .so build path; this path only ran the ctypes validate_kernel_config, which does not catch this) so invalid tiles are dropped instead of emitted/run. tile_k is unaffected (the K reduction has no CShuffle store constraint).

therock-pr-bot · 2026-07-01T05:15:29Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
🌿 Branch Name	✅ Pass	—
📝 PR Title/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.

therock-pr-bot · 2026-07-01T05:15:30Z

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

ISSUE ID: #8997 ## Motivation The CK Tile dispatcher could already generate and launch regular GEMM through the TileEngine → Dispatcher bridge, but it had no path for the multi-tensor **gemm_multi_abd** op. Multi-ABD is used when a GEMM needs to combine several A and B operands and fuse several D operands in the epilogue (`E = cde_op(a_op(As) @ b_op(Bs), {Ds})`), which is a real Old-TE capability with no dispatcher equivalent. This PR closes that gap so Python callers can drive multi_abd through the dispatcher at parity with the legacy Tile Engine version, without touching C++. It follows the divergent-ABI pattern established by the grouped bridge (#9000) because multi_abd needs **arrays** of A/B/D device pointers, not the single-pointer regular GEMM ABI. The capability set matches the Old-TE `gemm_multi_abd_instance_builder.py` exactly: `fp16`, `rcrr` layout, configurable A/B/D tensor counts, and the element-wise op set `{PassThrough, AddScale, MultiDMultiply, MultiDAdd}`. ## Test Plan - Run the CPU-only unit tests (no GPU required): `python3 -m pytest dispatcher/tests/test_multi_abd_bridge.py -v` - On-GPU numeric verification through the bridge launch path (gfx942 / MI300X), 512x512x512 fp16 rcrr, across the default 2/2/2 all-PassThrough config and non-PassThrough element-wise ops. - Confirm the CI and default config expansions yield the expected kernel counts. ## Test Result - CPU-only unit tests pass (10 passed). - Numeric verification (bridge launch path), 512x512x512 fp16 rcrr: - default 2/2/2 all-PassThrough: `max_rel = 2.9e-4` - CDE = MultiDAdd: `max_rel = 5.7e-4` - A-op = MultiDAdd: `max_rel = 4.1e-4` - all far below the fp16 tolerance (2e-2); 0 failed measurements. - CI config expands to 16 arch-valid kernels; `default_config.json` → 8896. - `standard` variant `expand_sweep` regression clean. - clang-format-18 (18.1.8) clean on `gemm_multi_abd_ctypes_lib.cpp`. - Serialized A/B perf-parity vs Old-TE (MI300X / gfx942, fp16 rcrr, 16 stems × 5 shapes = 80 rows, interleaved, fair 50/100/flush/rotating both sides): **at parity** — median gap -0.24%, mean -0.66%, 100% within ±15%, 87.5% within ±5% (range [-9.57%, +5.35%]). See the parity comment for details. --- **Related PRs / references (TileEngine → Dispatcher GEMM bridge series):** #8997 (regular GEMM fp16/bf16 all-layout), #9000 (grouped GEMM), #9028 (stream-K), #8887 (fp8/bf8/int8). This PR is a sibling in the same bridge effort tracked across those PRs. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>

ISSUE ID: #8997 ## Motivation The TileEngine → Dispatcher bridge had no path for the **gemm_multi_d** op, which fuses one or more extra D operands into the GEMM epilogue (`E = elementwise_op(A@B, D0, D1, ...)`). This is a real Old-TE capability used for fused bias/residual-style epilogues with no dispatcher equivalent, so this PR adds a complete bridge so the dispatcher can generate, build, and launch multi_d at parity with the legacy Tile Engine version. The capability set matches the Old-TE `gemm_multi_d_instance_builder.py` exactly: `fp16`, the 4-char layouts `{rcrr, rrrr, ccrr, crrr}` (A/B vary, C and D row-major), the element-wise ops `{MultiDAdd, MultiDMultiply, PassThrough}`, and a swept number of D tensors (1 and 2). It follows the registry-bypass bridge pattern used by the grouped (#9000) and stream-K (#9028) bridges. ## Test Plan - Run the CPU-only unit tests (no GPU required): `python3 -m pytest dispatcher/tests/test_multi_d_bridge.py -v` - On-GPU numeric verify over the full capability matrix (fp16 × {rcrr, rrrr, ccrr, crrr} × {MultiDAdd, MultiDMultiply} × {num_d 1, 2} = 16 combos) at M=N=K=1024 against an fp32 reference, gate 2e-2. - Confirm the CI config builds real kernels and the sweep covers all ops × D counts. ## Test Result - CPU-only unit tests pass (10 passed). - On-GPU numeric verify: 16/16 combos pass at M=N=K=1024, worst-case `max_rel = 6.16e-4` (~30x under the 2e-2 gate). Col-major `ccrr` / `crrr` have real on-GPU numeric evidence. - CI config now builds real kernels (was zero); the sweep expands evenly across `{MultiDAdd, MultiDMultiply} × {num_d 1, 2}` per layout. - clang-format (18.1.8) clean on `multi_d_gemm_ctypes_lib.cpp`. - Serialized A/B perf-parity vs Old-TE (MI300X / gfx942, fp16, 4 layouts × 2 ops × num_d=1 = 8 stems × 5 shapes = 40 rows, interleaved, fair 50/100/flush/rotating both sides): **at parity, bridge consistently faster** — median gap +9.44%, 100% within ±15% (range [+3.76%, +14.99%]; positive = bridge faster, from the registry-bypass direct launch avoiding the Old-TE profiler's per-call overhead). num_d=1 is the fair slice since the Old-TE `gemm_multi_d` benchmark is single-D. See the parity comment for details. --- **Related PRs / references (TileEngine → Dispatcher GEMM bridge series):** #8997 (regular GEMM fp16/bf16 all-layout), #9000 (grouped GEMM), #9028 (stream-K), #8887 (fp8/bf8/int8). This PR is a sibling in the same bridge effort tracked across those PRs. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

…idges (#10612) ISSUE ID: #8997 ## Summary This PR combines two sibling TileEngine → Dispatcher bridge ops into a single PR: - **batched GEMM** (previously #9306) - **batched_contraction** (previously #9328) Both follow the same **direct-launch, registry-bypass** pattern (as the stream-K bridge #9028), because their launch ABIs carry variable-length / batch-specific arguments the single-pointer registry backend cannot express. The two ops touch disjoint files except `dispatcher/tests/CMakeLists.txt`, where both GPU-correctness test registrations are kept. ## Motivation The dispatcher had no path for batched GEMM (same GEMM across many independent problems with per-batch strides) or batched_contraction (generalized batched tensor contraction `E[G.., M.., N..] = sum_K A[G.., M.., K..] * B[G.., N.., K..]` with multi-dim G/M/N/K index groups). Both are real Old-TE ops; this bridge lets Python callers generate, build, and launch them at parity with the legacy Tile Engine — without writing C++. ## Test Plan / Result - CPU-only unit tests for both bridges + gemm_utils: **74 passed** (`pytest dispatcher/tests/test_batched_bridge.py dispatcher/tests/test_batched_contraction_bridge.py dispatcher/tests/test_gemm_utils.py`). - Batched GEMM: end-to-end name-parity + correctness across batch counts 1/2/4/8 (`max_rel ~5e-4`), non-packed strides, split-K; full `default_config` codegen 6672 kernels / 0 failures. - Batched contraction: on-GPU verify (gfx950) across dtype × layout × shape × multi-dim × pipeline + D-tensor epilogue, all PASS. - clang-format-18 clean on both ctypes libs. ### Perf parity vs Old-TE - Batched GEMM (MI300X, fp16 rcr, batch=8): at parity / slightly ahead — median gap +4.00%, 100% within ±15%. - Batched contraction (MI350X, fp16 rcr): at parity — median gap -0.95%, 100% within ±15%. ## Scope / known limitations - Batched contraction: `rcr` only, `k_batch==1` only (split-K is a shared Old-TE kernel defect — hard-rejected, never silently-wrong), non-tile-multiple M/N/K rejected by `IsSupportedArguments`. --- Supersedes and closes #9306 (batched GEMM) and #9328 (batched_contraction). **Related PRs (TileEngine → Dispatcher GEMM bridge series):** #8997 (regular GEMM), #9000 (grouped), #9028 (stream-K), #8887 (fp8/bf8/int8), #9305 (multi-ABD), #9307 (preshuffle), #9308 (multi-D), #10439 (block-scale quant, 5 ops). --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Thrupti Raj Lakshmana Gowda <thruptiraj.lakshmanagowda@amd.com>

feat(ck-tile): TE to dispatcher GEMM bridge (fp16/bf16, all layouts) > Re-opened from #8479 with a compliant branch name (users/muozturk/ck-tile/gemm-bridge-all-layout-bf16-fp16). Supersedes #8479. ## Summary This PR routes the **Tile Engine (TE) regular-GEMM sweep through the Dispatcher**, making the Dispatcher the single source of truth for **codegen → build → runtime** while the Tile Engine keeps only the **config search space** and the **benchmark loop**. It is the consolidated, **single-commit** GEMM bridge covering **all four layouts (`rcr`/`rrr`/`crr`/`ccr`)** and **both `fp16` and `bf16`**. It is a clean re-roll of the earlier bridge work (previously split across #8123 + the stacked key/bf16/layouts/parity/example PRs and consolidated in #8261). Those branches accumulated unrelated cross-project commits through repeated `develop` merges; **this branch is a single clean commit off the latest `develop`** containing only the GEMM-bridge files. It supersedes and replaces #8123 / #8261. ## Motivation The Tile Engine historically owned its own codegen/build/runtime for GEMM (`tile_engine/ops/gemm/gemm_universal/`). The consolidation goal is for the **Dispatcher** to own all of that — exactly as it already does for **FMHA** and **Grouped Conv** — so there is one kernel-generation/build/runtime path and the TE shrinks to a config+benchmark frontend. This PR brings regular GEMM in line with that reference binding. ## The binding (mirrors the FMHA/Conv reference, six stages) 1. **Config JSON (TE side)** — the sweep search space lives in `tile_engine/ops/gemm/configs/` (flat op-root layout, matching the `fmha/` and `grouped_conv/` bridges). 2. **Codegen (Dispatcher)** — `dispatcher/codegen/unified_gemm_codegen.py` emits one fully-typed `.hpp` per kernel; `GemmKernelConfig.name` reproduces `KERNEL_NAME` **byte-for-byte** (the thread tying config → kernel → runtime). 3. **Compile to `.so`** — a single static `gemm_ctypes_lib.cpp` is force-included (`-include <kernel.hpp>`); one `.so` per kernel. 4. **Flat `extern "C"` ABI** — `dispatcher_run_gemm(A, B, C, M, N, K, time_ms)` + the kernel-name enumeration entry points. **Host-pointer** memory model (the C lib `hipMalloc`s internally) — the FMHA-forward branch of the reference. 5. **Python ctypes wrapper** — `dispatcher/python/gemm_utils.py` (`GemmDispatcherLib` + `GpuGemmRunner`). 6. **TE driver (3 phases)** — `gemm_full_benchmark.py` (parallel codegen+build → `expand_sweep` → subprocess-isolated benchmark) + the disposable per-kernel worker `run_one_gemm_kernel.py`. ## What's included **Bridge core** - `dispatcher/codegen/unified_gemm_codegen.py` — GEMM codegen, byte-exact naming. - `dispatcher/bindings/ctypes/gemm_ctypes_lib.cpp` — flat C ABI, host-pointer model. - `dispatcher/python/gemm_utils.py` — `GemmKernelConfig`, multi-kernel build (`setup_multiple_gemm_dispatchers`), `expand_sweep`, one-`.so`-per-kernel. - `tile_engine/ops/gemm/gemm_full_benchmark.py` + `run_one_gemm_kernel.py` — 3-phase, multi-GPU, subprocess-isolated driver/worker. **Feature surface (the point of this PR)** - **All four layouts** `rcr`/`rrr`/`crr`/`ccr` (row-major C only — ck_tile rejects column-major C at build) with layout-aware host transpose. - **`fp16` + `bf16`** (bf16 via uint16 byte-encoding; dtype derived from kernel name). - **Trait-derived registry `KernelKey`** — replaces the earlier hard-coded fp16/rcr key so the registry path generalizes across dtype/layout/tile. **Correctness & performance hygiene** - **`--verify`** opt-in fp32 numpy-reference gate (global `max|out-ref|/max|ref|`), `verified`/`max_rel` columns in the CSV; a mismatch counts as a failure. - **Tile Engine AMDGPU `-mllvm` codegen-flag parity** (without these the kernel builds with different occupancy and the timing diverges) and **arch-validated tile filtering** against the real pipeline/scheduler. - **Multi-GPU** fan-out across all visible GPUs (`--devices`, device-pinned `HIP_VISIBLE_DEVICES` workers). **Example & tests** - `dispatcher/examples/gemm/python/12_te_bridge.py` — runnable end-to-end example. - `dispatcher/tests/test_gemm_parity.py`, `test_gemm_utils.py`, and a parity regression harness. **Cleanup** - Removes the legacy standalone `gemm_universal` build path (`gemm_universal_instance_builder.py`, `*_benchmark*.{py,cpp,hpp}`, `gemm_universal/CMakeLists.txt`) and the old `test/ck_tile/gemm_tile_engine/` harness; promotes the sweep configs to the flat op-root `configs/`. ## Design decisions (consistent with the reference) - **Host-pointer memory ownership** (C lib owns device memory) — matches FMHA-forward; the Python runner passes host numpy arrays straight through. - **One `.so` per kernel** — packaging choice; the multi-kernel name ABI is retained (`get_kernel_name_at(0)` reports the single kernel), so the Python enumeration path is unchanged from FMHA/Conv. - **Flat `configs/`** at the op root — matches the `fmha/`/`grouped_conv/` convention; the not-yet-bridged variants keep their per-variant `configs/` dirs, selected by `--variant`. ## Validation (gfx942 / MI300X) - Bridge build + benchmark + `--verify` across **`fp16` and `bf16`** and **all four layouts**, checked against an fp32 numpy reference (`A @ B`). - **Name parity** holds end-to-end: each `.so`'s reported runtime name equals `GemmKernelConfig(...).name`. - bf16 passes under a widened fp16/bf16 tolerance; fp16 within the standard `max_rel` gate. ## Test plan - [ ] `gemm_full_benchmark.py --verify` over `configs/default_ci_config.json` for `fp16` and `bf16`, each of `rcr`/`rrr`/`crr`/`ccr`. - [ ] `unified_gemm_codegen.py` emits a header whose stem == `GemmKernelConfig.name`. - [ ] `setup_multiple_gemm_dispatchers` builds + links each config against `gemm_ctypes_lib.cpp`. - [ ] `pytest dispatcher/tests/test_gemm_parity.py dispatcher/tests/test_gemm_utils.py`. - [ ] `examples/gemm/python/12_te_bridge.py` runs end to end. ## Notes - Single clean commit off the latest `develop`; the diff is **35 files, all under `projects/composablekernel/`** (dispatcher + tile_engine/ops/gemm + test/ck_tile). - **Supersedes #8123 and #8261**, which will be closed. - Stream-K (#8136) and grouped GEMM are separate bridge efforts, not in this PR. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Muhammed Ozturk <muozturk@ctr2-alola-ctrl-01.amd.com>

…ops (#10132) ISSUE ID: #8997 ## Summary Adds **on-device GPU correctness tests** for the merged bridge operators that had **no in-tree GPU test**: grouped (#9000), multi-ABD (#9305), multi-D (#9308), mx (#9329). Closes part of the GPU-test coverage gap (11 of 15 bridge ops previously shipped only CPU-only unit tests). ## What each test does Builds the op's dispatcher `.so`, runs it **on-device** via the op's `Gpu*Runner`, and compares output to an independent **numpy fp32 reference** with a scale-relative tolerance (`max|err| / max|ref|`). Real math — reference is `A@B` per group / `A@B + D0 + D1` / the runner's built-in multi-ABD verify. No trivial always-pass asserts. They **SKIP cleanly** (not fail) when no GPU / hipcc / static-lib is present, so they add **zero cost on a CPU-only runner** and only execute on a GPU runner. ## Verification (MI300X / gfx942) - grouped — PASS (max_rel 7.6e-3 bf16 / 4.8e-4 fp16) - multi-ABD — PASS (6.4e-4) - multi-D — PASS (1.8e-3) - mx — **gfx950-only**, so it arch-gates and SKIPs on gfx942; verify-pending on a gfx950 runner. Registered in `dispatcher/tests/CMakeLists.txt` under label `dispatcher;python;bridge;gpu`. ## Scope Merged ops only (their runners are on develop). The 7 open bridge ops (batched/preshuffle/batched-contraction + 5 block-scale) get their GPU tests in their own PR branches. ## Test plan - [x] grouped / multi-ABD / multi-D PASS on gfx942 - [x] all SKIP cleanly with no GPU (CPU gate stays green) - [ ] mx verified on a gfx950 runner - [ ] CI runs the new `dispatcher_*_gpu_correctness` entries on a GPU runner --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Thrupti Raj Lakshmana Gowda <thruptiraj.lakshmanagowda@amd.com>

The regular-GEMM TE->Dispatcher bridge (#8997 fp16/bf16, #8998 fp8/bf8/int8) was CDNA-only (gfx90a/gfx942/gfx950, MFMA). This enables it on gfx1250 (MI400, RDNA4/WMMA): - gemm_utils: add gfx1250 to _SUPPORTED_ARCHES + wave-combo table; add an OCP fp8/bf8 host codec (float8_e4m3fn / float8_e5m2) selected by arch (gfx950/gfx12xx -> OCP, gfx942 -> FNUZ), since gfx1250's device fp8_t is OCP. - gemm_validation_utils: gfx1250 warp-configuration combos. - arch_specs_generated: gfx1250 fp8/bf8 WMMA warp tile (16x16x64). - gemm_full_benchmark: allow --dtype fp8/bf8. - configs: gfx1250 CI configs (fp16/bf16 warp_tile 16x16x32; fp8 16x16x64). - tests: CPU-only gfx1250/OCP-fp8 codec + arch tests. Validated on MI400/gfx1250 (ROCm 10.1): fp16/bf16 and fp8 kernels build and pass --verify (fp8 max_rel ~3.6e-2 vs fp32 ref, i.e. E4M3 quantization only).

ISSUE ID: #8997 ## Motivation The TileEngine → Dispatcher bridge had no path for the weight-preshuffled GEMM op **gemm_preshuffle**, which pre-permutes the B (weight) tensor into the pipeline's packed layout for higher throughput on weight-heavy GEMMs. This is a real Old-TE capability with no dispatcher equivalent, so this PR adds the bridge so the dispatcher can generate and launch preshuffle GEMM at parity with the legacy Tile Engine version. The capability set matches the Old-TE op exactly: `fp16`, `bf16`, `fp8`, `bf8`, `rcr` layout only (A row-major, B col-major, C row-major), with the B tensor preshuffled into the packed layout. It follows the same bridge scheme as #8997 (regular GEMM), #9000 (grouped), and #9028 (stream-K). ## Test Plan - Run the CPU-only unit tests (no GPU required): `python3 -m pytest dispatcher/tests/test_preshuffle_bridge.py -v` - Codegen + build the `.so` for all 4 dtypes and verify output vs an fp32 NumPy reference at M=N=K=512 on gfx942 / MI300X. - Confirm the preshuffle B-shuffle matches `ck_tile::shuffle_b` byte-for-byte. ## Test Result - CPU-only unit tests pass (8 passed). - Codegen + `.so` build succeed for all 4 dtypes (0 codegen failures). Output verified vs fp32 reference at M=N=K=512: | dtype | max_rel | status | |-------|---------|--------| | fp16 | 5.2e-4 | PASS | | bf16 | 5.5e-3 | PASS | | fp8 | 3.4e-2 | PASS | | bf8 | 6.5e-2 | PASS | - The preshuffle B-shuffle (the #1 correctness risk) is verified byte-for-byte via `ck_tile::shuffle_b`; the permuteN shuffle gives ~1.25 rel error and is deliberately not used, since the bridge does not emit the permuteN pipeline. > **Note (arch support):** the correctness above is on gfx942 (MI300X). fp8/bf8 preshuffle is unsupported on **gfx950** for *both* this bridge and Old-TE — the underlying kernel fails to launch there (`status -2`) — while fp16/bf16 are correct on both arches (see the cross-arch note in the comments). --- **Related PRs / references (TileEngine → Dispatcher GEMM bridge series):** #8997 (regular GEMM fp16/bf16 all-layout), #9000 (grouped GEMM), #9028 (stream-K), #8887 (fp8/bf8/int8). This PR is a sibling in the same bridge effort tracked across those PRs. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Thrupti Raj Lakshmana Gowda <thruptiraj.lakshmanagowda@amd.com>

The regular-GEMM TE->Dispatcher bridge (#8997 fp16/bf16, #8998 fp8/bf8/int8) was CDNA-only (gfx90a/gfx942/gfx950, MFMA). This enables it on gfx1250 (MI400, RDNA4/WMMA): - gemm_utils: add gfx1250 to _SUPPORTED_ARCHES + wave-combo table; add an OCP fp8/bf8 host codec (float8_e4m3fn / float8_e5m2) selected by arch (gfx950/gfx12xx -> OCP, gfx942 -> FNUZ), since gfx1250's device fp8_t is OCP. - gemm_validation_utils: gfx1250 warp-configuration combos. - arch_specs_generated: gfx1250 fp8/bf8 WMMA warp tile (16x16x64). - gemm_full_benchmark: allow --dtype fp8/bf8. - configs: gfx1250 CI configs (fp16/bf16 warp_tile 16x16x32; fp8 16x16x64). - tests: CPU-only gfx1250/OCP-fp8 codec + arch tests. Validated on MI400/gfx1250 (ROCm 10.1): fp16/bf16 and fp8 kernels build and pass --verify (fp8 max_rel ~3.6e-2 vs fp32 ref, i.e. E4M3 quantization only).

…atcher bridge (#13968) ## Summary Adds `mx_gemm_full_benchmark.py`, a JSON-driven sweep driver for microscaling (MX) GEMM through the TE to dispatcher bridge. It is the MX counterpart of `gemm_full_benchmark.py`. The Dispatcher stays the single source of truth for codegen, build and runtime. The Tile Engine keeps only the config search space (the same `mx_gemm/configs/*.json` the Old-TE CMake build reads) and the benchmark loop. Until now MX GEMM had the bridge Python API (#9329, #10916) but no sweep driver, so sweeping MXFP4/MXFP8 from a config JSON meant the Old-TE CMake flow. ## Why MX needs its own driver - The kernel space comes from `MxGemmKernelBuilder` (MX pipelines per arch, no K padding or persistence on gfx1250), not from `gemm_utils.expand_sweep`. - Inputs are packed fp4 or fp8 bytes with one e8m0 scale per 32 K elements, so the worker generates scale tensors and checks against the microscaled numpy reference instead of `A @ B`. - Kernels reject unsupported shapes at run time (rc -2, e.g. K not a multiple of tile_k). The worker reports these as skipped, not failed. ## Design note `expand_mx_sweep()` enumerates with the Old-TE `MxGemmKernelBuilder._get_sampled_kernel_list()`. That is the list the CMake `--list_kernels` step builds, including `--max-instances` sampling. A bridge sweep therefore runs exactly the kernels of an Old-TE build of the same JSON: all 472 kernels of `default_config.json` and `default_config_gfx1250.json` (fp4 and fp8) have byte-identical names. To avoid a second copy of the multi-GPU Phase 3 loop, `gemm_full_benchmark.py` now exposes `unique_built_kernels()`, `benchmark_on_devices()` and `print_summary()`. Its own `main()` uses them, and the MX driver imports them. `_run_batch_on_device()` gained an `extra` payload dict and a `skipped` result state; the regular GEMM worker never emits `skipped`, so its behaviour is unchanged. ## Changes **New** - `tile_engine/ops/gemm/mx_gemm_full_benchmark.py`: driver. Phase 1 expands the JSON and compiles the .so files in parallel, Phase 2 loads problems, Phase 3 runs on all GPUs. Supports `--dtype fp4|fp8`, `--verify`, `--max-instances/--seed`, `--output-dir` (build cache) and `--compile-only`. - `tile_engine/ops/gemm/run_one_mx_gemm_kernel.py`: disposable per-GPU worker. Uses random e8m0 scales (2^-3 to 2^2) and the microscaled reference check. - `mx_gemm/configs/bridge_ci_config_gfx950.json`, `bridge_ci_config_gfx1250.json`: small sweeps covering every MX pipeline of the arch (gfx950: 16 fp4 / 9 fp8 kernels; gfx1250: 13 fp4 / 11 fp8). These are the driver defaults. - `mx_gemm/configs/example_problems.json`: 5 tile-aligned M/N/K problems. **Modified** - `dispatcher/python/mx_gemm_utils.py`: adds `expand_mx_sweep()`. - `tile_engine/ops/gemm/gemm_full_benchmark.py`: shared Phase 3 helpers (no behaviour change). - `dispatcher/tests/test_mx_gemm_bridge.py`: Old-TE name parity, per-arch pipeline coverage, sampling, problem/tile fit, and worker verdicts (exact, wrong, unsupported). - `dispatcher/tests/test_gemm_driver_verify.py`: includes the new driver in the exit-code check; updates the 4-tuple mock. - `dispatcher/docs/mx_gemm.md`: usage section. Related bridge PRs: #8997 (regular GEMM), #9000 (grouped), #9028 (stream-K), #8887 (fp8/bf8/int8), #9329 (MX bridge), #10916 (MX gfx1250 pipelines). ## Capability | Axis | Supported | |---|---| | dtype | MXFP4 (e2m1, packed), MXFP8 (OCP e4m3); e8m0 scales per 32 K; fp16 output | | layout | rcr | | arch | gfx950 (comp_async, comp_async_eight_waves, weight_preshuffle), gfx1250 (+ comp_tdm, comp_tdm_v2) | ## Validation (gfx950, 8 GPUs) | Run | Kernels | Measurements | Result | |---|---|---|---| | CPU tests (`test_mx_gemm_bridge.py`, `test_gemm_driver_verify.py`, `test_gemm_driver_fp32.py`) | - | - | 207 passed | | `bridge_ci_config_gfx950.json` fp4 `--verify` | 16 | 80 | 80 verified, max_rel 0 | | `bridge_ci_config_gfx950.json` fp8 `--verify` | 9 | 45 | 45 verified | | Old-TE `default_ci_config.json` fp4 / fp8 `--verify` | 2 / 2 | 10 / 10 | all verified | | Regular `gemm_full_benchmark.py` fp16 `--verify --max-kernels 4` (refactor check) | 4 | 20 | all verified | | Old-TE `default_config.json` fp4 `--verify` | 194 (all built) | 970 | 928 verified, 4 mismatch, 42 errors (see below) | gfx1250 has CPU coverage only (expansion, names, pipeline coverage); it has not been run on hardware yet. ## Open items (why this is a draft) The full fp4 sweep found two groups of failures that I'm still investigating: 1. **rc=-1 for 14 kernels** (comp_async 256x64 tiles and weight_preshuffle 192x512) on the 2048^3, 4096^3 and 4096x2048x8192 problems. Run alone in a fresh worker, the same kernel/problem passes and verifies, so this looks like interaction between kernels batched in one worker process, not a kernel defect. 2. **Wrong results for comp_async_eight_waves 128x128x128 4x2x1** (max_rel 0.2 to 0.33) at K = 4096 and 8192; K <= 2048 passes. Next step is to check whether the Old-TE build of the same kernel gives the same mismatch (kernel issue) or not (bridge host-side issue). gfx1250 hardware validation is also pending. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>

…atcher bridge (#10931) JIRA ID : AICK-2106 ## Summary This PR routes the gfx1250 (MI400) `gemm_preshuffle` variant through the TE to dispatcher bridge for all five gfx1250 pipelines. The Dispatcher is the single source of truth for codegen -> build -> runtime; the Tile Engine keeps only the config search space and the benchmark loop. The pipeline names follow the gemm_universal TDM/async bridge (#12544). | Pipeline (JSON) | Native pipeline | B layout | Scheduler | Epilogue | Datatypes | | --- | --- | --- | --- | --- | --- | | `preshufflev2` | WeightPreshufflePipelineAGmemBGmemCRegV2 | packed (`shuffle_b_v0`) | `default` | `cshuffle` | fp16, bf16, fp8, bf8 | | `preshuffle_tdm` | WeightPreshufflePipelineAGmemBGmemCRegTDM | packed (`shuffle_b_v0`) | `default` | `cshuffle` | fp16, bf16, fp8, bf8 | | `comp_tdm` | GemmPipelineAgBgCrCompTDMV1 | ordinary | `intrawave` | `tdm` | fp16, bf16, fp8, bf8 | | `comp_tdm_v2` | GemmPipelineAgBgCrCompTDMV2 | ordinary | `intrawave` | `tdm` | fp16, bf16, fp8, bf8 | | `comp_async` | GemmPipelineAgBgCrCompAsync | ordinary | `intrawave` | `cshuffle` | fp16, bf16; fp8/bf8 only with warp_tile_k >= 128 | No core kernel changes (`include/ck_tile/ops`, `test/ck_tile`) are needed. ## Motivation Before this PR, the preshuffle bridge only emitted `preshufflev2`, with wave64 assumptions, so gfx1250 had no preshuffle kernels in the dispatcher or TE. gfx1250 has four more pipelines that can serve the preshuffle op. The gfx1250 TDM/async work in #12544 explicitly left preshuffle out, so this PR adds it. ## Design note - Only the two weight-preshuffle pipelines read packed B. The kernel trait is `Preshuffle = pipeline in PACKED_B_PIPELINES` (`preshufflev2`, `preshuffle_tdm`). The ctypes host path keys its B shuffle off the same `SelectedKernel::Preshuffle` metadata (from #11042), and the TE profiler does the same, so a comp_* kernel always receives ordinary B. - Packed B uses `shuffle_b_v0`, the packing of the native preshuffle V2 example. gfx1250 loads B in 16-byte chunks, and `shuffle_b` interleaves those chunks differently when a lane holds more than one. - I first tried making the comp_* kernels read packed B too (#12621, now closed). A packed-B vs ordinary-B benchmark on gfx1250 showed no benefit (median -0.4% over 198 pairs; comp_async with tile_k 64 was 20-32% slower), so the comp_* pipelines keep ordinary B and need no core change. - comp_tdm/comp_tdm_v2 skip padding validation in their native argument check, so the generated launcher and the dispatcher backend reject partial N/K tiles and split-K at runtime instead of returning wrong results. ## Validity rules One reject helper, `preshuffle_pipeline_reject_reason`, holds the rules. Dispatcher codegen (`codegen_common.py`) and TE validation (`gemm_validation_utils.py`) each have a copy, and a parity test keeps them identical: - On gfx1250, all five pipelines are rcr only and non-persistent. - `preshufflev2`/`preshuffle_tdm` use the `default` scheduler; comp_* use `intrawave`. comp_tdm/comp_tdm_v2 use the `tdm` epilogue; the others use `cshuffle`. - TDM pipelines (`preshuffle_tdm`, `comp_tdm`, `comp_tdm_v2`) require pad_m=pad_n=pad_k=False. - `comp_tdm_v2` requires 4 waves. - `comp_async` requires pad_m=pad_n=pad_k=True, and fp8/bf8 require warp_tile_k >= 128 (same rules as gemm_universal on gfx1250). - `preshufflev2` rejects pad_n/pad_k on gfx1250: `shuffle_b_v0` packs whole N x K tiles, so a partial N tile overflows the packed buffer and a partial K tile packs the wrong elements. pad_m is still allowed. - Feature-qualified targets (for example `gfx1250:xnack-`) and string pad flags (`"false"`, `"True"`) are normalized the same way on both sides. ## Default config and CMake - `tile_engine/ops/gemm/gemm_preshuffle/configs/default_config_gfx1250.json`: rcr, 128x128x{64,128} block, 1x4x1 waves, 16x16 warp tile, all five pipelines, both pad settings, non-persistent. It yields 12/12/10/10 kernels for fp16/bf16/fp8/bf8 (fp8/bf8 have no `comp_async`, since their warp_tile_k is 64). - The TE CMake uses the gfx1250 config and pipeline list only when gfx1250 is the only GPU target. For a mixed list (for example `gfx942;gfx1250`), it skips gfx1250 with a warning and builds the other architectures from the default config. - `gemm_full_benchmark.py --variant gemm_preshuffle --arch gfx1250` picks the gfx1250 config by default. Related bridge PRs: #8997 (GEMM), #9000 (grouped), #9028 (stream-K), #8887 (fp8/bf8/int8), #12544 (gfx1250 TDM/async). ## Validation (gfx1250 / MI400, merged with current develop) | Check | Result | |---|---| | Simple test (`dispatcher/tests/test_preshuffle_gpu_correctness.py --gfx gfx1250`) | fp16: 94 cases pass, max rel err 4.2e-4; bf16: 94 cases pass, max rel err 3.5e-3; 8 unsupported shapes per dtype rejected cleanly | | Comprehensive, dispatcher: all 44 instances of `default_config_gfx1250.json` (12/12/10/10), 9 shapes each | 0 wrong results. Worst rel err: fp16 4.5e-4, bf16 3.4e-3, fp8 4.6e-4, bf8 4.1e-4. Unpadded kernels reject 1000^3, 200x300x136 and 129x257x520 at argument check | | comp_async (pads TTT) on unaligned shapes | 1000^3 runs and verifies; 200x300x136 and 129x257x520 are rejected by the kernel argument check (no wrong output) | | comp_async fp8/bf8 at warp_tile_k 64 and 128 | rejected by `expand_sweep`, `validate_kernel_config` and codegen | | Comprehensive, Tile Engine: `benchmark_gemm_preshuffle_all` for fp16/bf16/fp8/bf8 rcr | 44 binaries (12/12/10/10). `-verify=1`: 92/92 executed runs correct; the 8 comp_async runs on 200x300x136 / 129x257x520 report "Arguments not supported" (same argument check as the dispatcher) | | CMake configure | `GPU_TARGETS=gfx1250` builds all five pipelines; `gfx942` unchanged; `gfx942;gfx1250` skips gfx1250 with a warning | | CPU tests | dispatcher preshuffle / arch filter / LDS / autocorrect / quant tests: 152 passed (680 subtests); TE `test_gemm_preshuffle_gfx1250.py`: 12 passed; `test_kernel_key`: 7/7 | ## Example ```bash cd projects/composablekernel/tile_engine/ops/gemm python gemm_full_benchmark.py --variant gemm_preshuffle --arch gfx1250 --dtype fp16 --layout rcr \ --verify --verify-tol 0.01 --csv gfx1250_fp16.csv ``` --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Thrupti Raj Lakshmana Gowda <thruptiraj.lakshmanagowda@amd.com>

…idge (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). --------- Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com> Co-authored-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com> Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 53fc044c-7bc8-41f2-b5ae-b03bfde3b135

ozturkosu and others added 16 commits June 16, 2026 00:25

Potential fix for pull request finding

a741d14

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Potential fix for pull request finding

1af77db

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Potential fix for pull request finding

fa7aa1b

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

polish comments

29c8cd5

Fix stale checkout ref in libraries PR bot workflow

b45fadb

Merge branch 'develop' into muozturk/gemm-bridge-all-layouts-bf16

2361d88

ozturkosu requested review from a team as code owners July 1, 2026 05:11

github-actions Bot added documentation project: composablekernel github actions labels Jul 1, 2026

ozturkosu self-assigned this Jul 1, 2026

ozturkosu requested review from ThruptiRajLakshmanaGowda and yraparti July 1, 2026 05:12

ozturkosu mentioned this pull request Jul 1, 2026

feat(ck_tile): TE -> Dispatcher GEMM bridge (all layouts, fp16/bf16) #8479

Closed

5 tasks

ozturkosu mentioned this pull request Jul 1, 2026

feat(ck-tile): TE to dispatcher GEMM bridge for fp8/bf8/int8 (all layouts) #8998

Merged

3 tasks done

ozturkosu changed the title ~~feat(ck_tile): TE -> Dispatcher GEMM bridge (all layouts, fp16/bf16)~~ feat(ck-tile): TE to dispatcher GEMM bridge (fp16/bf16, all layouts) Jul 1, 2026

ozturkosu mentioned this pull request Aug 10, 2026

feat(ck-tile): batched GEMM + batched-contraction TE to dispatcher bridges #10612

Merged

ozturkosu mentioned this pull request Oct 2, 2026

feat(ck-tile): gfx1250 preshuffle GEMM with five pipelines in TE/dispatcher bridge #10931

Merged

This was referenced Oct 19, 2026

feat(ck-tile): MX GEMM full benchmark sweep driver for the TE to dispatcher bridge #13968

Merged

feat(ck-tile): comp_tdm_pc producer/consumer GEMM TE to dispatcher bridge (gfx1250) #13975

Merged

ozturkosu mentioned this pull request Nov 3, 2026

feat(ck-tile): adaptive kernel search for the TE to dispatcher GEMM bridge #14938

Open

Navigation Menu

Repository navigation

feat(ck-tile): TE to dispatcher GEMM bridge (fp16/bf16, all layouts) - #8997

feat(ck-tile): TE to dispatcher GEMM bridge (fp16/bf16, all layouts)#8997
ozturkosu merged 25 commits into
developfrom
users/muozturk/ck-tile/gemm-bridge-all-layout-bf16-fp16

ozturkosu commented Jul 1, 2026 •

edited by afagaj

Loading

Uh oh!

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

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

ozturkosu commented Jul 1, 2026 • edited by afagaj Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Motivation

The binding (mirrors the FMHA/Conv reference, six stages)

What's included

Design decisions (consistent with the reference)

Validation (gfx942 / MI300X)

Test plan

Notes

Uh oh!

therock-pr-bot Bot commented Jul 1, 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 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

ozturkosu commented Jul 1, 2026 •

edited by afagaj

Loading

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

edited

Loading

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

edited

Loading