Repository navigation
Conversation
The merged multi-D GEMM bridge (#9308) ships only an MFMA CI config (warp_tile 32x32x16), which is valid on gfx942/gfx950 but reports as unsupported (status -2/-1) on gfx1250. gfx1250 has no MFMA units; its RDNA4 WMMA fp16/bf16 warp tile is 16x16x32. Add a gfx1250 multi_d CI config that pins the WMMA warp tile (16x16x32) while keeping the multi_d_config D-fusion block, extend the multi_d bridge unit test to assert the gfx1250 config uses WMMA (not MFMA) and retains the multi_d block, and add gfx1250 to the driver's --arch help. Depends on #10904 (shared gfx1250 dispatcher enablement: _SUPPORTED_ARCHES, wave-combo fallbacks, arch_specs_generated.py gfx1250 warp tiles). Validated on an MI400 (gfx1250, ROCm 10.1) node: built the dispatcher static lib for gfx1250 and ran the multi_d sweep with the new CI config (fp16, rcrr, --verify) across all 4 default problems (1024^3, 2048^3, 4096^3, 512x512x8192). 80 measurements verify against the fp32 numpy reference (max_rel ~5e-4), spanning compv3/compv4/mem pipelines, both MultiDAdd/MultiDMultiply ops, and both D-counts, at 26-69 TFLOPs (median 31). Remaining status -1 kernels are unsupported trait combos gracefully surfaced by the bridge, not correctness failures.
Contributor
Author
gfx1250 (MI400) GEMM bridge enablement — status across the stacked PR series
All enablement PRs are stacked on #10904 (shared gfx1250 arch enablement), Python/JSON-only, each with a CPU test. Validated on MI400 / gfx1250, ROCm 10.1. Remaining C++ follow-ups: stream-K cross-workgroup coherency handshake (CDNA atomics → RDNA4/wave32); mx gfx1250 scale-preshuffle (ck_tile already ships |
This was referenced Aug 18, 2026
Contributor
Author
|
Superseded by #10921, which combines the shared gfx1250 enablement (gemm_universal) with the grouped, multi-D, and multi-ABD variant configs+tests into a single PR against develop. Branch retained (not deleted) as the base of the remaining partial PRs. |
ozturkosu
added a commit
that referenced
this pull request
Aug 31, 2026
…rouped + multi-D (#10921) JIRA ID : AICK-2098 JIRA ID : AICK-2099 JIRA ID : AICK-2100 JIRA ID : AICK-2137 ## Summary Enables the **regular (gemm_universal)**, **grouped**, **multi-D**, and **multi-ABD** GEMM Tile-Engine → Dispatcher bridges on **gfx1250 (MI400, RDNA4 / WMMA)**, for fp16/bf16 (+ fp8 on the regular path). This is the consolidation of the previously-stacked PRs #10904 / #10914 / #10915 / #10917 into one. Python + JSON only (no C/C++) — no clang-format gate. The two *partial* variants (**stream-K #10919**, **microscaling mx #10916**) are **not** included — they need C++/kernel follow-ups (see below) and remain separate. ## Do the merged bridges (Universal/grouped/multi-D) support gfx1250? — No. As merged, none of these bridges run on gfx1250. They share a CDNA-only arch gate (`_SUPPORTED_ARCHES = gfx90a/gfx942/gfx950`) and most pinned the MFMA `32x32x16` warp tile in their CI config, which is invalid on gfx1250's RDNA4 WMMA engine (fp16/bf16 WMMA is `16x16x32`, fp8 is `16x16x64`). ## Changes **Shared enablement (was #10904)** - `dispatcher/python/gemm_utils.py`: `gfx1250` in `_SUPPORTED_ARCHES` + gfx1250 wave combos; OCP fp8/bf8 codec (`float8_e4m3fn`/`float8_e5m2`) selected by arch (gfx950/gfx12xx → OCP, gfx942 → FNUZ). - `tile_engine/ops/gemm/gemm_validation_utils.py`: gfx1250 warp-configuration combos. - `dispatcher/codegen/arch_specs_generated.py`: gfx1250 fp8/bf8 WMMA `[[16,16,64]]` (fp16/bf16 `[[16,16,32]]` already present). - `tile_engine/ops/gemm/gemm_full_benchmark.py`: `SUPPORTED_DTYPES` += fp8, bf8. - `tile_engine/ops/gemm/configs/`: gfx1250 regular CI configs (fp16/bf16 16x16x32; fp8 16x16x64). **Variant configs + tests** - grouped: `grouped_gemm/configs/default_ci_config_gfx1250.json`. - multi-D: `gemm_multi_d/configs/default_ci_config_gfx1250.json` + `--arch` help. - multi-ABD: `gemm_multi_abd/configs/default_ci_config_gfx1250.json` (fp16-only, upstream design). - CPU-only tests across `dispatcher/tests/` (gfx1250 arch, OCP fp8 codec, grouped/multi-D/multi-ABD configs). 63 tests pass. ## Validation (MI400 / gfx1250, ROCm 10.1) | Bridge | Result | |---|---| | gemm_universal | 640/640 fp16+bf16 VERIFY, 0 fail; fp8 up to ~625 TFLOP/s (max_rel ~3.8e-2 = E4M3 quant) | | grouped | 64/64 fp16 OK (~178–217 TFLOP/s) | | multi-D | 80 measurements VERIFY (max_rel ~5e-4) | | multi-ABD | 32/32 fp16 VERIFY (fp16-only by upstream design) | ## Notes - `arch_specs_generated.py` is auto-generated from `arch_specs.json`, which has **no gfx1250 entry** on develop (the gfx1250 fp16/bf16 rows were hand-added to the generated file). This PR follows that precedent; a follow-up should add a full gfx1250 stanza to `arch_specs.json`. - fp8 `--verify` needs a looser tol (~0.05–0.15) than the fp16 default (2e-2), since the reference is fp32 and fp8 carries E4M3 quantization. - **Not included (separate, partial):** stream-K (#10919 — kernel hangs on RDNA4: cross-workgroup K-reduction coherency relies on CDNA atomics) and microscaling mx (#10916 — GPU `.so` build `static_assert`s gfx950-only via `preShuffleScaleBuffer_gfx950`; ck_tile already ships `preShuffleScaleBuffer_gfx1250`). Both need C++ follow-ups. Supersedes #10904, #10914, #10915, #10917. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
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
Enables the merged multi-D GEMM bridge (#9308) on gfx1250 (MI400).
Stacked on #10904 (
users/muozturk/ck/gemm-universal-gfx1250-enable), which supplies the shared gfx1250 dispatcher enablement. This PR adds only the multi_d-specific piece on top.Does the merged multi_d bridge support gfx1250? No (as merged).
Two reasons the merged #9308 does not run on gfx1250:
gemm_utils._SUPPORTED_ARCHES, the wave-combo fallback tables, andarch_specs_generated.pydid not include gfx1250. This is fixed by the base PR feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) #10904.gemm_multi_d/configs/default_ci_config.json) pinswarp_tile = 32x32x16, an MFMA shape. gfx1250 has no MFMA units; it runs the RDNA4 WMMA path whose fp16/bf16 warp tile is 16x16x32. On gfx1250 the MFMA kernels report as unsupported (status -2/-1) and never run.Changes
gemm_multi_d/configs/default_ci_config_gfx1250.json(new): mirrors the multi_d CI config but pins the WMMA warp tile (16x16x32) and keeps themulti_d_configD-fusion block (elementwise_ops x num_d_tensors).dispatcher/tests/test_multi_d_bridge.py: newTestMultiDGfx1250Configasserting the gfx1250 config exists, uses WMMA (16x16x32, not MFMA 32x32x16), and retains the multi_d block. The pre-existingtest_configs_are_valid_sweepsalso now covers the new config.gemm_multi_d_full_benchmark.py: add gfx1250 to the--archhelp text.No C/C++ changes, so no clang-format needed. The merged multi_d bridge is fp16-only by design, so this enablement targets fp16.
MI400 validation (gfx1250, ROCm 10.1)
Built the dispatcher static lib for
-DGPU_TARGETS=gfx1250and ran:max_rel ~5e-4), evenly across all 4 default problems (1024^3, 2048^3, 4096^3, 512x512x8192) — 20 verified rows per problem.compv3/compv4/mempipelines, bothMultiDAdd/MultiDMultiply, and bothnum_d_tensors1/2.status -1kernels are unsupported trait combos gracefully surfaced by the bridge (not correctness failures); every problem still has passing kernels.Notes
arch_specs_generated.pyis auto-generated fromarch_specs.json, which has no gfx1250 entry. Per the precedent in feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) #10904, the generated file is hand-edited for gfx1250 until the source spec is updated.Test plan
pytest dispatcher/tests/test_multi_d_bridge.py(13 passed)--verify: 80/80 measured kernels correct