Repository navigation
Conversation
The grouped GEMM bridge (#9000) shares the arch gate and WMMA warp tiles with the regular bridge, so it runs on gfx1250 once gfx1250 is enabled (see the base branch / #10904). Its CI config already uses the gfx1250-valid WMMA warp tile 16x16x32, so no grouped-specific codegen change is needed. - add an explicit gfx1250 grouped CI config (default_ci_config_gfx1250.json) - add CPU-only tests: grouped multi-problem plumbing, gfx1250 arch acceptance, and that the grouped gfx1250 CI config uses the WMMA 16x16x32 warp tile Validated on MI400/gfx1250 (ROCm 10.1): grouped_gemm_full_benchmark over the CI config, fp16, 64/64 kernels OK (~178-217 TFLOP/s), 0 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 grouped GEMM bridge (merged #9000) on gfx1250 (MI400, RDNA4/WMMA). Stacked on #10904 (the shared gfx1250 enablement); Python + JSON only.
Does the merged grouped bridge support gfx1250? — Not on its own.
As merged, #9000 shares the regular bridge's arch gate (
_SUPPORTED_ARCHES = gfx90a/gfx942/gfx950) and WMMA/warp-tile tables, so on gfx1250 it fails the same way (Unsupported GPU architecture 'gfx1250'). Once the shared gfx1250 enablement (#10904:_SUPPORTED_ARCHES, wave tables, arch_specs) is in place, grouped runs with no grouped-specific codegen change — its CI config already uses the gfx1250-valid WMMA warp tile16x16x32(unlike the regular CI config, which used the CDNA MFMA32x32x16).Changes
tile_engine/ops/gemm/grouped_gemm/configs/default_ci_config_gfx1250.json— explicit gfx1250 grouped CI config (WMMA 16x16x32), for parity with the regular gfx1250 config.dispatcher/tests/test_gemm_utils.py— CPU-only tests: grouped multi-problem roundtrip/flops, gfx1250 arch acceptance, and that the grouped gfx1250 CI config uses the WMMA warp tile.Validation (MI400 / gfx1250, ROCm 10.1)
grouped_gemm_full_benchmark.py grouped_gemm/configs/default_ci_config.json --arch gfx1250 --dtype fp16: 64/64 kernels OK, 0 failures, ~178–217 TFLOP/s across the trait sweep.Notes
--arch gfx942; pass--arch gfx1250on MI400 (a follow-up could switch it to rocminfo auto-detect).arch_specs_generated.pyis auto-generated fromarch_specs.json, which has no gfx1250 entry on develop (the gfx1250 rows were hand-added to the generated file); this stack follows that precedent.Sibling PRs: #10904 (gemm_universal + shared arch, base of this stack).