[CK_TILE] Add Tile Engine -> Dispatcher bridge for GEMM#8123
[CK_TILE] Add Tile Engine -> Dispatcher bridge for GEMM#8123ozturkosu wants to merge 15 commits into
Conversation
Introduce the shared-config bridge that lets Tile Engine drive the GEMM dispatcher the same way FMHA and grouped conv already do: one config dataclass owned by the dispatcher, imported by Tile Engine, with no translator between two vocabularies. - dispatcher/python/gemm_utils.py: GemmKernelConfig (the common contract; .name mirrors the codegen KERNEL_NAME byte-for-byte), GemmProblem, GemmDispatcherLib, GpuGemmRunner, setup_multiple_gemm_dispatchers (codegen + hipcc -> .so paths, CPU-only/parallel), and expand_sweep. - dispatcher/bindings/ctypes/gemm_ctypes_lib.cpp: add the indexed multi-kernel ABI dispatcher_get_kernel_name_at(index, buf, size); legacy single-kernel dispatcher_get_kernel_name retained. - tile_engine/ops/gemm/gemm_full_benchmark.py: 3-phase TE driver (compile -> load problems -> benchmark) that generates no binaries. - tile_engine/ops/gemm/run_one_gemm_kernel.py: disposable GPU worker for subprocess fault isolation. Scope: regular GEMM, fp16, rcr (Phase 1). Name parity verified end-to-end (config.name == generated .hpp stem == runtime registry name). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
There was a problem hiding this comment.
Pull request overview
Adds a Tile Engine ↔ Dispatcher bridge for GEMM, aligning GEMM with the existing FMHA/grouped-conv “shared config owned by dispatcher” pattern so Tile Engine can generate/search configs while the dispatcher handles codegen/compile/runtime dispatch.
Changes:
- Introduces
dispatcher/python/gemm_utils.pyproviding the sharedGemmKernelConfigcontract plus helpers to expand TE sweeps, build per-config.solibraries, and run kernels via a thin ctypes wrapper. - Extends the GEMM ctypes library ABI to support enumerating kernel names by index (
dispatcher_get_kernel_name_at) while retaining the legacy single-kernel name API. - Adds Tile Engine-side benchmark driver + isolated subprocess worker for running GEMM kernels safely during sweeps.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 3 comments.
| File | Description |
|---|---|
| projects/composablekernel/tile_engine/ops/gemm/run_one_gemm_kernel.py | New isolated worker process that loads a compiled GEMM .so and runs one/batched kernels, returning JSON results. |
| projects/composablekernel/tile_engine/ops/gemm/gemm_full_benchmark.py | New 3-phase benchmark driver (compile → load problems → benchmark via subprocess isolation) that consumes the dispatcher bridge. |
| projects/composablekernel/dispatcher/python/gemm_utils.py | New dispatcher-owned GEMM bridge module defining GemmKernelConfig, sweep expansion, parallel build to .so, and runtime wrappers. |
| projects/composablekernel/dispatcher/bindings/ctypes/gemm_ctypes_lib.cpp | Adds indexed multi-kernel name enumeration ABI while keeping the legacy single-kernel name function. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| except Exception as e: | ||
| print(f" Batch error: {e}") | ||
| try: | ||
| if proc and proc.poll() is None: | ||
| proc.kill() | ||
| except Exception: | ||
| pass |
| parser.add_argument("--arch", default="gfx942") | ||
| parser.add_argument("--dtype", default="fp16") | ||
| parser.add_argument("--layout", default="rcr") |
There was a problem hiding this comment.
Fixed in 35da60d9ff. Added module-level SUPPORTED_DTYPES/SUPPORTED_LAYOUTS and wired them into argparse choices=, so a value the Phase-1 runner can't honor now fails fast instead of silently benchmarking the wrong thing:
error: argument --dtype: invalid choice: 'bf16' (choose from 'fp16')
On this anchor PR the supported sets are fp16 / rcr. Propagated up the stack, widening each set to the branch's real capability: bf16 added in #8190, and rrr/crr/ccr in #8191 (column-major C stays out — ck_tile rejects it at build).
| for i in unique: | ||
| c = configs[i] | ||
| codegen_args.append( |
Phase 2 parity gate — all 4 test-plan items executed on gfx942 (MI300X), fp16/rcrEverything below was produced through the bridge ( 1. Numeric parity vs numpy fp32 reference — PASS
max_rel ~3–4e-04 = fp16 accumulation tolerance. 2. Performance medians (≥12 runs, +3 warmup) — STABLE
CV ≤ 2.9% → reproducible. 3. Top-K fastest kernels — coherent & reproducibleSwept 48 configs → 32 unique kernels × 4 shapes. Top-1 is a
Top-1 identical in both runs everywhere. Lower overlap on smaller shapes is boundary churn among near-ties (top-5 span < 1.5% TFLOPS). Old-TE head-to-head: the legacy TE GEMM flow emits only 4. fp16/rcr sweep pass rate — explained
Pass rate on shapes compatible with the swept kernels' pad settings: 100%. VerdictPhase-2 parity gate met: correct numerics, stable perf, coherent/reproducible top-K, fully-explained sweep. All rejections are the dispatcher faithfully surfacing real kernel constraints (no-pad divisibility, codegen geometry) shared with the legacy TE flow. Ready to route stream_k / grouped through the same bridge (Phase 3). |
Status update — FMHA-gap closure (measured vs reference PR #5260)Two of the generality gaps called out when benchmarking this bridge against the
Both validated on gfx942 (MI300X):
Suggested land order: #8123 → #8187 → #8190. Remaining toward full parity with the reference: layouts beyond |
Update — third generality gap closed (layouts)Following the gap-closure table above, the rcr-only gap is now also addressed:
Validated on gfx942 (MI300X), numeric vs CPU ref: fp16 rcr/rrr/ccr/crr max_rel 5.8e-4; bf16 rrr 7.8e-3 (dtype × layout compose). Column-major C is rejected at build by ck_tile's universal GEMM (row-major C only) — a kernel limitation, not a host issue. Full stack land order: #8123 → #8187 → #8190 → #8191. All three FMHA generality gaps (hard-coded key, fp16-only, rcr-only) are now closed. |
Update — final gap closed (runnable example); stack completeThe last item from the FMHA-parity evaluation ( Full FMHA-parity stack (each PR stacked on the previous):
Validation for #8193 (gfx942/MI300X, M=N=K=512): 6/6 pass — fp16 All architectural generality gaps and the example gap are closed. The only remaining parity item is merge (#4 in the rubric), which needs human review — these are all still draft. |
setup_multiple_gemm_dispatchers built its per-config codegen args without the config's variant, and _generate_single_kernel_subprocess hard-coded --variants standard. A GemmKernelConfig(variant='preshuffle'/'multi_d') passed to the bridge would therefore be code-generated as a standard kernel, while its name (and the hpp_glob_pattern derived from it) still carried the variant suffix -- so the lookup could never match the emitted header. Pass the variant through; the subprocess default stays "standard" so all existing callers are unaffected. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Update — variant now threaded through the bridge codegen path (Copilot review item)Addresses the automated-review flag that Fix (commit
Backward-compatible: the subprocess helper is shared, pre-existing develop Fix committed here and propagated up the stack (#8187 → #8190 → #8191 → #8193). |
…dler If subprocess.Popen() itself raises, the generic except handler referenced proc before it was bound, masking the real error with an UnboundLocalError. Initialize proc = None before the try; the handler already guards with `if proc and proc.poll() is None`. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Update — fixed UnboundLocalError risk in the GEMM benchmark (Copilot review item)Addresses the flag on Fix (commit |
Copilot review item: the driver exposed unrestricted --dtype/--layout while the Phase-1 worker hard-codes fp16 inputs and an rcr (column-major B) host transpose. Passing e.g. --dtype bf16 would codegen bf16 kernels but feed them fp16 data, silently benchmarking the wrong thing. Add SUPPORTED_DTYPES/ SUPPORTED_LAYOUTS and wire them into argparse choices so a mismatch fails fast. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
How to use the GEMM bridgeThree entry points, all driven by the single shared Prereq: build the dispatcher first so A. CLI sweep — the benchmark drivercd projects/composablekernel/tile_engine/ops/gemm
python3 gemm_full_benchmark.py gemm_universal/configs/default_config.json \
--arch gfx942 --dtype fp16 --layout rcr --csv gemm_results.csv
# optional: --problems shapes.json --workers 8 --batch-size 20 --max-kernels 50Three phases: compile (CPU, parallel — returns
B. Programmatic — the
|
There was a problem hiding this comment.
it will be good to include example config files to sweep. And later include the json used in nightly tests.
| ) | ||
|
|
||
| # ======================================================================== | ||
| # Phase 3: Benchmark via subprocess (serial GPU, batched) |
There was a problem hiding this comment.
please see if you can launch on multiple GPU devices in parallel. the serial gpu design in Conv might be an inferior design compared to fmha launcher.
|
the code is still not organized as per the gemm variant. Also please add a README once it is done. |
Parity check is done. Look at #8193 (comment) However there is a validity issue for (rrr, crc, ccr) maybe we can land this PR now as only rcr and then focus on remaning format. |
Addresses the three review items on the TE->Dispatcher GEMM bridge driver, scoped to this foundation PR's fp16/rcr surface (bf16/layouts follow in the #8190/#8191 stack): 1. Example configs to sweep - gemm_full_benchmark.py defaults to the selected variant's configs/default_ci_config.json (small CI sweep) when no config is passed, and to configs/example_problems.json when --problems is omitted; configs/default_config.json remains the full sweep. - New gemm_universal/configs/example_problems.json (square / rectangular / large M,N,K). Nightly-test JSON drops into the same configs/ dir -- no driver change needed. 2. Multi-GPU launch in parallel (supersedes grouped_conv's serial-GPU design) - Phase 3 fans the (kernel x problem) work across every visible GPU: one worker thread per device pulls batches from a shared queue and spawns a disposable subprocess pinned with HIP_VISIBLE_DEVICES, so an N-GPU box runs ~Nx faster while keeping per-batch fault isolation. - Devices auto-detected (HIP_VISIBLE_DEVICES, then rocm-smi/amd-smi); override with --devices (count, explicit ids, or all). 3. Variant organization + README + deprecation note - --variant selects the per-variant configs/ directory. - New README "Dispatcher Bridge Workflow" section: scripts, per-variant config layout, run examples, multi-GPU explanation, supported surface (fp16/rcr here), and a deprecation note for the legacy *_instance_builder.py generators. Driver --dtype/--layout choices stay fp16/rcr to match this PR's dispatcher host path; run_one_gemm_kernel.py (fp16 host gen) is unchanged.
Layout × awkward-size sweep: old-TE ground truth + root causeFollowing up on the "full layout × dtype gfx942 sweep failed for Ground truth (old TE, develop branch)
Rejections are the kernel's own raw strings, e.g. Conclusion: this is not a bridge bug. Old TE rejects 257×129×512 on every layout using the identical kernel and Root causefp16 vector width = 16 B / 2 B = 8 elements. Which dim is contiguous depends on layout (3-char = A-major, B-major, C-major):
Checks fire in order A → B → C, so the reported dim is the first violator:
Every cell in the table matches this rule. Fix options
ReproductionBuilt from a |
…duler The codegen arch filter hard-coded pipeline=compv4 / scheduler=intrawave when validating tile geometry, ignoring each config's actual traits. Since compv4 has the strictest MFMA constraints, tiles legal under mem/compv3 were wrongly rejected -- collapsing the generated fp16/rcr set from ~1520 to 512 kernels (compv3 and mem each decimated ~5x; compv4 roughly preserved). Thread the trait's real pipeline/scheduler into _is_tile_arch_valid at both call sites; the tile pre-filter now keeps a tile if it is legal under any configured pipeline/scheduler, with the precise per-trait check deferred to _get_configs_for_variant. Verified on the 6144 fp16/rcr config population: emitted kernels 512 -> 1520 (compv3 464, compv4 128, mem 928), and a previously-rejected compv3 64x64x192 config now generates a header end-to-end. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Codegen arch filter: validate tiles against the real pipeline/scheduler (commit 4d14777)BugWhile investigating why the bridge built far fewer fp16/rcr kernels than old Tile-Engine (522
Evidence (fp16/rcr, gfx942)Two validators disagreed:
Generated-header counts on disk matched the bug signature —
FixThread the trait's real Verification
Residual gap1520 vs old-TE's 2385 is expected: old TE applies no gfx942 arch filter, so it builds whatever its instance builder enumerates (some of which may be arch-invalid on gfx942). The bridge's |
…artifact, not a speedup) The sweep's >=20% "bridge faster" cells (all compv4/intrawave/1024^3) are NOT a bridge speedup. Proven on MI300X: the device kernel is byte-identical (same ck_tile::kentry<1,GemmKernel<...>> symbol, rocprof), and through any uniform harness it runs at the same speed on both sides. Ruled out empirically: kernel, compiler/flags (4 toolchain rebuilds incl. clang++-HIP with old-TE flags all give ~189), all bench knobs (warmup/repeat/flush/rotating/timer), allocation/placement (DeviceMem, 4GB decoys), and stale timing headers (byte-identical across trees). rocprof confirms the slowdown is real device time (13.78us vs 11.34us): old TE's *standalone benchmark binary* runs the identical kernel ~18-20% slower purely due to that process's GPU clock/execution state (+8% stall cycles under PMC, plus ~13% lower sustained SCLK). Fix: - ab_same_harness.py: apples-to-apples A/B that builds the old-TE kernel into a .so and runs BOTH it and the bridge .so through the SAME worker. Gap collapses to ~+/-0.5% at 1024^3 (was +20..+24% vs the standalone binary). Proof in ab_same_harness.out. - diagnose.md sec.4 rewritten: the prior hipcc-vs-clang++ toolchain theory is disproven and replaced with the evidence above. - generated_tile_backend.hpp: correct the misleading comment that claimed matching bench knobs makes bridge-vs-old-TE "apples-to-apples". Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The >=20% "bridge faster" cells are a benchmark-harness artifact, not a speedup (commit db55f7e)Followed up on the suspicious cells in the fp16/rcr sweep (all 16 are Ruled out (each tested on-GPU)
What it actually isA clean standalone harness measures the old-TE kernel at 189-194; old TE's standalone benchmark binary measures the same kernel at 156. PMC counters: old-TE kernel = 366k cycles vs 339k (+8% stall cycles); non-PMC wall ratio (1.22) > cycle ratio (1.08) -> remaining ~13% is lower sustained SCLK in that process. Shape-selective because large shapes hit the power cap and clocks converge; 1024^3 has headroom so the two processes' DPM governors diverge most. Fix: measure both kernels through the SAME harness (
|
| shape | bridge | oldTE | gap% |
|---|---|---|---|
| 512^3 | 38.77 | 38.78 | -0.01 |
| 1024^3 | 189.19 | 189.41 | -0.12 |
| 2048^3 | 295.59 | 297.10 | -0.51 |
| 4096^3 | 369.48 | 369.85 | -0.10 |
So bridge and old TE run the same kernel at the same speed. Also corrected the misleading "matching bench knobs => apples-to-apples" comment in generated_tile_backend.hpp. Full evidence in diagnose.md §4.
Resolve CMakeLists.txt conflict in tile_engine/ops/gemm: keep develop's expanded op list and new add_subdirectory entries while preserving this branch's retirement of legacy gemm_universal (dropped from both budget foreach loops and from add_subdirectory).
Remove dispatcher/parity_diag/regression/diagnose.md from the PR; the content now lives on Confluence (MLSE) as a child page under the fp16/rcr A/B sweep report: https://amd.atlassian.net/wiki/spaces/MLSE/pages/1737132108
- ab_same_harness.py: derive ROOT from __file__ and take old-TE header dir from OLD_TE_GEN env var (was hard-coded dev paths); drop unused statistics import - generated_tile_backend.hpp: make env_bool case-insensitive (handles False/Off) and align its comment - gemm_full_benchmark.py: clarify in --devices help and resolve_devices docstring that a bare digit is a count; a single id needs the comma form (5,) Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…ridge # Conflicts: # projects/composablekernel/dispatcher/python/gemm_utils.py
The Stream-K bridge (#8136) was branched at #8123's first commit, so it lacked all subsequent regular-GEMM bridge improvements (arch-validated tile filtering, the develop merge + legacy gemm_universal retirement, benchmark-param/--verify work on the shared driver, README). Merge the current #8123 HEAD to pick those up; the Stream-K-specific analogues that live in the duplicated driver/worker/ctypes lib are ported in follow-up commits. Sole conflict: dispatcher/python/gemm_utils.py variant threading. Kept the Stream-K routing (_ctypes_source_name -> streamk_gemm_ctypes_lib.cpp, .name _streamk suffix, variant through codegen_args/expand_sweep) and adopted #8123's explanatory comment.
The Stream-K bridge keeps its own driver, worker and ctypes lib, so the regular-GEMM bridge improvements that landed on #8123 after this branch forked did not arrive via the merge. Port the Stream-K-specific analogues: - streamk_gemm_ctypes_lib.cpp: benchmark knobs now default to old-TE's warmup=50/repeat=100 (was 3/10 -- a cold, un-ramped clock, the root of #8123's spurious "perf gap") and are env-overridable via CK_TILE_BENCH_WARMUP/REPEAT/FLUSH/ROTATING. Unlike the regular path, rotating_count defaults to 1: the Atomic preprocess re-zeros only the original C buffer, so rotating C would corrupt the accumulation. - streamk_gemm_full_benchmark.py: fan the (kernel x problem) work across every visible GPU (device-pinned HIP_VISIBLE_DEVICES workers, --devices, device CSV column), add the --verify/--verify-tol fp32-reference gate, and constrain --dtype/--layout to the supported fp16/rcr surface. Also fixes a latent proc-unbound error in the batch handler. - run_one_streamk_gemm_kernel.py: add the fp32 numpy reference check (global max|out-ref|/max|ref|, verified/max_rel) behind --verify. - README: document the Stream-K bridge driver/worker, flags, _streamk name suffix, fp16 Atomic tolerance, and the rotating_count divergence. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
GEMM Parity — Old TE engine vs Dispatcher bridge (PR #8123)Head-to-head performance + validity comparison between the old Tile-Engine Scope — the buildable-common fp16/rcr set: every kernel that Shapes (M×N×K) — every dim is a multiple of 768 = lcm(64,128,192,256) so every Matched knobs (both engines) — warmup=50, repeat=100, flush_cache=true, Artifacts: Coverage
The 130 no-compare instances are exactly the 26 Validity — both engines, independent
Functional parity is clean. Every validity failure on both engines is the Performance parity — overall (gap = (bridge − oldTE) / oldTE)Over the 5830 comparable instances:
The distribution is centered on parity (median ≈ 0) but has a wide, two-sided Per shape
Dispersion is not just a small-shape launch-overhead artifact: the large, Where the divergence concentrates
10 largest deficits / surpluses (full list in CSV)
Bottom line
|
The bridge compiled GEMM kernels with only -mllvm -enable-noalias-to-md-conversion=0, omitting the AMDGPU inlining/codegen flags Tile Engine passes (-amdgpu-early-inline-all=true, -amdgpu-function-calls=false, --lsr-drop-solution=1, -enable-post-misched=0, -fno-offload-uniform-block). Those flags change register allocation and thus occupancy. Persistent GEMM kernels size their grid by occupancy (UniversalGemmKernel::MaxOccupancyGridSize = #CUs x occupancy), so the mismatch produced large perf gaps vs Tile Engine on persistent tiles (up to ~-73% on some mem-pipeline 256x256x64 / 192x64x128 kernels) while non-persistent kernels (fixed problem-sized grid) were unaffected. Matching the flags restores parity: on the full persistent buildable-common fp16/rcr subset (596 kernels x 5 shapes), the bridge-vs-Tile-Engine gap collapses to a median ~-2% with no kernel beyond +/-20% (was -51.7% .. +203.8%). Numerically unchanged (max_rel ~3.5e-4). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Perf parity fix — match Tile Engine's AMDGPU codegen flags (
|
| kernel | old-TE | bridge (before) | bridge (after fix) |
|---|---|---|---|
mem_default …True_256x256x64_2x2x1_32x32x16 |
255.4 | 123.7 (−51.6%) | 262.2 (+2.7%) |
mem_cshuffle …True_256x256x64_1x4x1_32x32x16 |
267.6 | 72.7 (−72.8%) | 266.8 (−0.3%) |
Fix
d04623df41 adds the same -mllvm flag set to the bridge compile in gemm_utils.py. Output numerically unchanged (max_rel ~3.5e-4).
Confirmation — full persistent subset (fixed bridge vs current-develop Tile Engine)
596 persistent fp16/rcr kernels × 5 shapes, matched knobs (warmup 50 / repeat 100 / flush / rotating 1000):
- median gap −2.1%, 98.6% within ±10%, 0 kernels beyond ±20% (was −51.7% … +203.8%)
mempipeline specifically: median −2.2%, range −10.2 / +8.6%, 0 beyond ±20%
The residual ~±10% is ordinary codegen dispersion present on both engines, not a bridge effect.
Confirmation complete — full persistent subset (596 kernels × 5 shapes)The post-fix A/B run finished (fixed bridge vs Tile Engine, both built from current
The |
…ridge (#8261) Brings #8261 current with #8123's latest work: arch-validated tile filtering, matching Tile Engine AMDGPU -mllvm codegen flags, the --verify fp32-reference gate, perf-gap diagnosis correction, Copilot nits, and the develop catch-up. Conflict resolution (tile_engine/ops/gemm): - Adopted #8123's flat op-root configs/ layout for gemm_universal (gemm_universal/ removed; matches the fmha/grouped_conv bridge convention and the on-disk rename). - Kept #8261's superset support surface: SUPPORTED_DTYPES=(fp16,bf16), SUPPORTED_LAYOUTS=(rcr,rrr,crr,ccr) in driver + README, since #8261 already wires bf16 and all four layouts through the dispatcher host path. - Took #8123's --verify driver/worker wiring and device-count semantics. - example_problems.json: identical content, kept at configs/.
Superseded by #8479 — work moved thereClosing this PR. The GEMM bridge work has been moved to #8479 This branch's history had accumulated unrelated cross-project commits through All functionality from this PR is preserved in #8479: the TE → Dispatcher regular Please review #8479 instead. |
Re-stacks the grouped_gemm bridge on top of the all-layouts GEMM bridge (muozturk/gemm-bridge-all-layouts-bf16, #8479) instead of the now-closed #8123 base. Squashes the prior grouped-branch history (was 5e665a3.. d075b90 on the old base) into the net grouped delta replayed onto #8479. Contents: - codegen: grouped variant + launch generator (arch_filter.py, unified_gemm_codegen.py), plus the standalone 02_grouped_gemm_driver.cpp and README sections (absorbed from the closed #8075). - bridge: multi-problem registry-bypass ctypes ABI (grouped_gemm_ctypes_lib.cpp), GpuGroupedGemmRunner + dtype/layout codecs (gemm_utils.py), TE driver/worker harness, GROUPED_GEMM_BRIDGE.md. Kept both #8479's GEMM_KEY_* registry descriptors and grouped's ALayout/BLayout/CLayout single-include exports. Coverage: {rcr,rrr,ccr,crr} x {fp16,bf16,fp8,bf8}; validated at Old-TE parity on MI300X/gfx942.
…8997) > 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>
Summary
Adds the Tile Engine → Dispatcher bridge for GEMM, following the exact pattern FMHA and grouped conv already use: a single shared config dataclass owned by the dispatcher and imported by Tile Engine. There is no translator between two vocabularies — both sides share the one object whose
.nameis the runtime registry key.What's in this PR
dispatcher/python/gemm_utils.py(new) — the bridge:GemmKernelConfig— the shared contract..namereproduces the codegenKERNEL_NAMEbyte-for-byte. The warp/wave mapping (TE "warp" = wave count →wave_*; TE "warp_tile" = MFMA shape →warp_tile_*) lives in one place so it cannot drift.setup_multiple_gemm_dispatchers()— codegen + hipcc →.sopaths; CPU-only, parallel, returns paths (no GPU). Reusesctypes_utilsleaf helpers so codegen/compile have a single source of truth.GemmProblem,GemmResult,GemmDispatcherLib,GpuGemmRunner,expand_sweep.dispatcher/bindings/ctypes/gemm_ctypes_lib.cpp(upgrade) — adds the indexed multi-kernel ABIdispatcher_get_kernel_name_at(index, buf, size); the legacy single-kerneldispatcher_get_kernel_nameis retained for backward compatibility.tile_engine/ops/gemm/gemm_full_benchmark.py(new) — 3-phase TE driver (compile → load problems → benchmark). Generates no binaries.tile_engine/ops/gemm/run_one_gemm_kernel.py(new) — disposable GPU worker for subprocess fault isolation (one bad kernel kills one worker, not the sweep).Scope
Regular GEMM, fp16, rcr (Phase 1). stream_k / grouped_gemm / extra dtypes route through this same bridge in a later phase.
Verification
GemmKernelConfig.name== generated.hppstem == runtime registry name, confirmed on validatedexpand_sweepconfigs (6144 valid from the default config)..so; the newdispatcher_get_kernel_name_atsymbol is exported and the registry reports the expected name.Test plan
Next Steps