Skip to content

feat(ck-tile): enable grouped AQuant/ABQuant/BQuant GEMM bridges on gfx1250 (MI400) - #10927

Open
ozturkosu wants to merge 4 commits into
developfrom
users/muozturk/ck/grouped-quant-gfx1250-enable
Open

feat(ck-tile): enable grouped AQuant/ABQuant/BQuant GEMM bridges on gfx1250 (MI400)#10927
ozturkosu wants to merge 4 commits into
developfrom
users/muozturk/ck/grouped-quant-gfx1250-enable

Conversation

@ozturkosu

@ozturkosu ozturkosu commented Aug 18, 2026

Copy link
Copy Markdown
Contributor

Summary

Enables three merged grouped quant GEMM TE->dispatcher bridges on gfx1250 (MI400):

Stacked on the shared gfx1250 foundation in #10921 (base branch users/muozturk/ck/gemm-bridges-gfx1250-enable).

Requires C++? YES (small)

The 3 ctypes-lib runtime arch allow-lists rejected gfx1250 and had to be extended — that is the only C++ change. Each of grouped_gemm_{aquant,abquant,bquant}_ctypes_lib.cpp had a runtime gate of the form if(arch != gfx950 && arch != gfx942 && arch != gfx90a) return -1;; added && ... != gfx1250 and updated the (supported: ...) message. Each file is clang-format-18 clean.

Correct gfx1250 config — FlatMM warp_tile_k=128 (on-GPU determined)

The original revision of this PR pinned the gfx12 WMMA tile (warp_tile_k=16) on the theory that gfx1250 needs WMMA. On-GPU validation on MI400/gfx1250 (fp32-dequant reference) proved that wrong: warp_tile_k=16 returns all-zeros for every grouped-quant variant. The empirically-correct gfx1250 config for the fp8/bf8 quant path is the FlatMM tile, warp_tile_k=128 — the same tile the gfx9 fp8/bf8 FlatMM decode uses.

variant dtype gfx1250 config warp_tile on-GPU status max_rel_err
aquant fp8 16x16x128 PASS ~3e-4
aquant bf8 16x16x128 PASS ~1e-4
abquant (compv3) fp8 16x16x128 PASS ~3e-4
abquant (compv3) bf8 16x16x128 PASS ~3e-4
bquant fp8 16x16x128 (= stock) PASS ~5e-4
bquant bf8 16x16x128 (= stock) PASS ~5e-4
a/bquant fp8i4 / bf8i4 deferred (does not compile on gfx1250)
  • aquant / abquant: gfx9 stock uses warp_tile_k=32 (plain MFMA), which is WRONG on gfx1250 (aquant max_rel=1.0; abquant all-zero). The gfx1250 helpers now pin warp_tile_k=128.
  • bquant: the stock fp8/bf8 config is already warp_tile_k=128 and is correct on gfx1250 as-is, so the gfx1250 helper simply delegates to the stock config.
  • fp8i4 / bf8i4 (packed-int4 weights): do not compile on gfx1250 at any warp_tile_k (no gfx12 instruction for the packed-int4 quant path). The C++ arch-gate stays open, but on-GPU correctness is deferred pending kernel work. No gfx1250 config helper is shipped for i4 (it would emit a non-buildable kernel).

Changes

  • C++ (arch gates only): add gfx1250 to the runtime allow-list + (supported: ...) string in grouped_gemm_{aquant,abquant,bquant}_ctypes_lib.cpp; clang-format-18 applied.
  • Python utils: gfx1250 helpers pin the FlatMM tile warp_tile_k=128:
    • aquant: default_{fp8,bf8}_config_gfx1250
    • bquant: default_{fp8,bf8}_config_gfx1250 (delegate to stock)
    • abquant: default_{fp8,bf8}_compv3_config_gfx1250 (compv3, transpose_c=false)
    • i4 gfx1250 helpers removed (unbuildable on gfx1250).
  • CPU tests: each TestGfx1250Configs now asserts warp_tile_k=128 / 16x16x128 names and the correct pipeline (aquant mem; bquant/abquant compv3; abquant not eightwaves/transpose_c).

Validation

  • CPU + codegen: pytest dispatcher/tests/test_grouped_gemm_{aquant,bquant,abquant}_utils.py176/176 pass.
  • On-GPU (MI400/gfx1250, a07-1, HIP_VISIBLE_DEVICES=0): all 6 fp8/bf8 configs PASS against an fp32-dequant reference (see table + validation comment). Uses the real corrected helper functions.

Dependencies / follow-ups

Note: no Co-Authored-By trailer per team policy.

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).
Fold the grouped (#9000), multi-D (#9308), and multi-ABD (#9305) variant
enablement onto the shared gfx1250 foundation so gemm_universal + these three
"full" variants ship in one PR:

- grouped:   gfx1250 CI config (WMMA 16x16x32) + CPU tests
- multi-D:   gfx1250 CI config (WMMA 16x16x32) + CPU tests + --arch help
- multi-ABD: gfx1250 CI config (WMMA 16x16x32, fp16-only) + CPU tests

Each merged variant pinned the CDNA MFMA 32x32x16 tile in its CI config, which
does not run on gfx1250 WMMA; the new configs use 16x16x32. No variant codegen
change was needed beyond the shared arch enablement.

Validated on MI400/gfx1250 (ROCm 10.1): grouped 64/64 fp16 OK; multi-D 80
measurements VERIFY (max_rel ~5e-4); multi-ABD 32/32 fp16 VERIFY. 63 CPU tests pass.
…fx1250 (MI400)

Enables the three merged grouped quant GEMM TE->dispatcher bridges on
gfx1250, stacked on the shared gfx1250 foundation (PR #10921).

C++ (small, required): the three ctypes-lib runtime arch allow-lists rejected
gfx1250; add gfx1250 to each gate and to the "(supported: ...)" message.
clang-format-18 applied.

Python: add gfx1250 default-config helpers that pin warp_tile_m/n/k=16 so the
codegen selects the gfx12 WMMA instruction (warp_tile_k=32/128 select the
gfx9-only MFMA/FlatMM paths, which silently return zeros on gfx12). ABQuant
uses the standard CompV3 path (not the gfx950-native eightwaves/FlatMM path).

Tests: CPU unit tests assert the gfx1250 configs use WMMA warp_tile_k=16,
propagate gfx_arch, and (abquant) use compv3 rather than eightwaves.

CPU + codegen validated (176/176 quant utils tests pass); on-GPU validation
pending gfx1250 node recovery.
@ozturkosu ozturkosu self-assigned this Aug 18, 2026
@ozturkosu
ozturkosu marked this pull request as ready for review August 18, 2026 07:19
@ozturkosu
ozturkosu requested a review from a team as a code owner August 18, 2026 07:19
@ozturkosu

Copy link
Copy Markdown
Contributor Author

On-GPU validation on MI400/gfx1250 (a07-1)

Node: heliosr-1b114-a07-1 (MI400 / gfx1250, ck-emin container, ROCm)
Date: 2026-08-18
Device: pinned HIP_VISIBLE_DEVICES=0 (device index 1 on this node is a known HW fault, excluded)
Method: fresh standalone composable_kernel clone + overlay of this branch's changed CK files. For each grouped-quant variant, built one gfx1250 fp8 kernel from its default_fp8_*_config_gfx1250() helper via setup_multiple_*_dispatchers(..., gfx_arch="gfx1250"). That step hipcc --offload-arch=gfx1250-compiles the changed grouped_gemm_*quant_ctypes_lib.cpp together with the gfx12 WMMA-codegen kernel header — so a successful .so proves both the C++ ctypes arch-gate and the gfx12 WMMA codegen (warp_tile 16x16x16).

Build (C++ arch-gate + WMMA codegen) — primary goal

Variant gfx1250 config .so builds
aquant grouped_gemm_aquant_fp8_rcr_mem_cshuffle_intrawave_16x64x256_1x4x1_16x16x16_aqg1x1x128
abquant grouped_gemm_abquant_fp8_rcr_compv3_permute_n_intrawave_128x128x128_1x4x1_16x16x16_...
bquant grouped_gemm_bquant_fp8_rcr_compv3_cshuffle_intrawave_16x64x256_1x4x1_16x16x16_qg1x1x128

All three arch gates accept gfx1250 and the gfx12 WMMA (16x16x16) kernels compile cleanly.

On-GPU run

Each .so loads and executes on gpu0 with rc=0 (no crash/hang) and returns a finite output buffer. However, the output was all-zero even for trivial inputs (A=B=all-ones fp8, unit A-scale, both float8_e4m3fnuz and float8_e4m3 codecs), where a real A@B over K=512 should be non-zero. This points to a test-harness / kernel-registration gap in my lightweight standalone runner (fp8/scale representation or a missing example-.so registration step), not something I can attribute to a kernel defect with confidence.

On-GPU numerical run needs a proper reference harness — deferred. The build result above (C++ gate + WMMA codegen OK for all three variants) is the load-bearing validation for this PR, which changed the ctypes arch gates.

No code changes made.

@ozturkosu
ozturkosu changed the base branch from users/muozturk/ck/gemm-bridges-gfx1250-enable to develop August 18, 2026 07:43
@therock-pr-bot

therock-pr-bot Bot commented Aug 18, 2026

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR 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.

🙋 Wish to Override Policy?

@ozturkosu

Copy link
Copy Markdown
Contributor Author

This is PR 3 of 3 in the recommended review order (review after #10921 and #10928; config correction in progress).

Review order & dependencies — gfx1250 (MI400) GEMM-bridge enablement series

These PRs enable the merged CK GEMM bridges on gfx1250 (MI400, RDNA4/WMMA). They all target develop, and each currently includes the shared gfx1250 enablement (gemm_utils._SUPPORTED_ARCHES + wave tables, arch_specs_generated.py gfx1250 rows, gemm_full_benchmark.py dtypes, OCP fp8 codec), so they overlap on those shared files. Please review/merge in this order:

1️⃣ #10921 — foundation: gemm_universal + grouped + multi-D + multi-ABD.
Review & merge first. Carries the shared gfx1250 enablement. Fully GPU-validated on MI400/gfx1250 (single-checkout run): universal 80/0, grouped 64/0, multi-D 256/0, multi-ABD 48/16 VERIFY (the 16 are the non-divisible 257³ default shape correctly rejected by pad=false kernels).

2️⃣ #10928 — batched GEMM + batched contraction.
Review after #10921. GPU-validated (batched_gemm 32/32 verify ~5e-4; batched_contraction bridge 1/1 + default 8/8 verified; the 4 non-builds are CompV3/CompV4+interwave pipeline incompatibilities, not gfx1250 bugs). Adds the batched variants' own arch tuples + gfx1250 WMMA 16x16x32 configs.

3️⃣ #10927 — grouped AQuant/ABQuant/BQuant.
Review last. Needs a small C++ change (3 ctypes-lib runtime arch gates). ⚠️ Config correction in progress: the initial gfx1250 configs (warp_tile 16x16x16) produced all-zero output; on-GPU diagnosis showed the stock FlatMM/MX config (16x16x128) is what runs correctly on gfx1250 (bquant fp8 verified, max_rel ~5e-4). The PR is being updated to the empirically-verified configs — please hold review until the corrected validation comment is posted.

Merge mechanics

After #10921 merges to develop, rebase #10928 and #10927 onto develop; their diffs then collapse to just the variant-specific files (the shared files land via #10921), giving a clean focused review for each.

Separate, not in this order (need C++/kernel follow-ups, kept as drafts): #10916 (mx — GPU blocked by a gfx950-only static_assert), #10919 (stream-K — kernel hangs on RDNA4 wave32).

…le_k=128

On-GPU validation on MI400/gfx1250 (fp32-dequant reference) shows the previous
warp_tile_k=16 (WMMA) gfx1250 overrides return all-zeros for every grouped-quant
variant. The correct gfx1250 config for fp8/bf8 is the FlatMM tile
(warp_tile_k=128):

  - aquant  fp8/bf8: 16x16x16 -> all-zero; stock 16x16x32 -> WRONG (max_rel 1.0);
            16x16x128 -> PASS (max_rel ~3e-4 / ~1e-4)
  - abquant compv3 fp8/bf8: 16x16x16 and 16x16x32 -> all-zero;
            16x16x128 -> PASS (max_rel ~3e-4)
  - bquant  fp8/bf8: stock is already 16x16x128 -> PASS (max_rel ~5e-4);
            16x16x16 -> all-zero. gfx1250 helper now returns the stock config.

Changes:
  - aquant/abquant gfx1250 helpers now pin warp_tile_k=128 (FlatMM).
  - bquant gfx1250 helpers delegate to the stock fp8/bf8 config.
  - Remove fp8i4/bf8i4 gfx1250 helpers: the packed-int4 quant path does not
    compile on gfx1250 at any warp_tile_k (no gfx12 instruction). The C++ ctypes
    arch-gate stays open for i4, but on-GPU correctness is deferred.
  - CPU tests updated to assert warp_tile_k=128 / 16x16x128 names.

C++ ctypes arch-gates (add gfx1250) are unchanged and remain correct.
@ozturkosu

Copy link
Copy Markdown
Contributor Author

Corrected gfx1250 config + on-GPU validation

This supersedes my earlier all-zero comment. That comment concluded the all-zero output was a test-harness/scale gap in a lightweight standalone runner. That conclusion was wrong: using a proper fp32-dequant reference harness on MI400/gfx1250 shows the all-zeros were caused by the config (warp_tile_k=16, gfx12 WMMA), not the harness. The correct gfx1250 config for the fp8/bf8 grouped-quant path is the FlatMM tile, warp_tile_k=128.

Node: heliosr-1b114-a07-1 (MI400 / gfx1250, ck-emin container)
Device: HIP_VISIBLE_DEVICES=0 (device index 1 is a known HW fault, excluded)
Reference: fp32 dequant reference — bquant via dispatcher/tests/test_bquant_gpu_correctness.py (sets proper e8m0/float scales); aquant/abquant via the in-tree examples 14/15 fp32 dequant references. All runs use the real default_*_config_gfx1250 helper functions from this branch.

Per-variant on-GPU result (corrected)

variant dtype config warp_tile pipeline PASS/FAIL max_rel_err
aquant fp8 16x16x128 mem PASS 0.0003
aquant bf8 16x16x128 mem PASS 0.0001
abquant fp8 16x16x128 compv3 PASS 0.0003
abquant bf8 16x16x128 compv3 PASS 0.0003
bquant fp8 16x16x128 compv3 PASS 0.0005
bquant bf8 16x16x128 compv3 PASS 0.0005
a/bquant fp8i4 / bf8i4 deferred build fails on gfx1250

What was wrong before, and why 128 is correct

Direct A/B comparison on gfx1250 for each variant (same shapes, fp32 reference):

variant warp_tile_k=16 (old PR) stock (=32) =128 (FlatMM)
aquant fp8 all-zero WRONG (max_rel 1.0, half nonzero) PASS 3e-4
abquant fp8 all-zero all-zero PASS 3e-4
bquant fp8 all-zero (stock is already 128) PASS 5e-4
  • bquant stock already uses warp_tile_k=128, so it is correct on gfx1250 with no distinct config — the gfx1250 helper now delegates to the stock config.
  • aquant / abquant stock uses warp_tile_k=32 (plain MFMA) which is wrong/zero on gfx12; the gfx1250 helpers now pin warp_tile_k=128.
  • fp8i4 / bf8i4 do not compile on gfx1250 at any warp_tile_k (no gfx12 packed-int4 quant instruction) → arch-gate left open, on-GPU correctness deferred; no unbuildable config helper shipped.

Changes pushed (ad8a03201f)

  • aquant/abquant gfx1250 helpers: warp_tile_k 16 → 128.
  • bquant gfx1250 helpers: delegate to stock fp8/bf8 (128).
  • removed i4 gfx1250 helpers.
  • CPU tests updated to assert warp_tile_k=128 / 16x16x128; 176/176 quant utils tests pass.
  • C++ ctypes arch-gates unchanged (correct as-is), clang-format-18 clean.

@therock-pr-bot

Copy link
Copy Markdown

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

ozturkosu added a commit that referenced this pull request Aug 18, 2026
Enables the 5 block-scale quant GEMM bridges (aquant/abquant/bquant/
rowcolquant/tensor_quant) on gfx1250:

- quant_bridge_common.hpp: accept gfx1250 in the shared runtime arch check
  (+ updated the unsupported-arch message).
- aquant/abquant: add gfx1250 to _SUPPORTED_ARCHS (bquant/rowcolquant/
  tensor_quant already accept gfx12 via _uses_ocp_fp8).
- All 5 arch-derived warp_tile_k selectors (_warp_tile_k_for /
  fp8_warp_tile_k_for_arch): treat gfx1250 like gfx950 -> 128 (FlatMM) for
  the 8-bit-float quant path. warp_tile_k=32 is the gfx9 MFMA path and
  silently outputs zeros on gfx1250 (same trap found in the grouped-quant
  work #10927); fp8/bf8 quant needs the FlatMM 16x16x128 tile on gfx1250.

OCP fp8 for gfx12 was already wired in the utils build path. The CI configs
already sweep warp_tile_k in [16,32,64,128], so 128 is exercised.
@ozturkosu
ozturkosu requested review from illsilin and yraparti August 18, 2026 08:33
@ozturkosu

Copy link
Copy Markdown
Contributor Author

📋 Merge order for the gfx1250 (MI400) bridge-enablement series

This is the coordination note for landing the gfx1250 enablement PRs. There are two independent tracks plus follow-ups.

Track A — core GEMM bridges (must be ordered; they share files)

1️⃣ #10921 — FIRST. Foundation: gemm_universal + grouped + multi-D + multi-ABD plus the shared gfx1250 enablement (gemm_utils.py _SUPPORTED_ARCHES/wave tables, arch_specs_generated.py, gemm_full_benchmark.py dtypes, OCP fp8 codec, gemm_validation_utils.py). Fully GPU-validated on MI400/gfx1250. Everything in this track depends on it.

2️⃣ #10928 — batched GEMM + batched contraction. After #10921. GPU-validated.
3️⃣ #10927 — grouped AQuant/ABQuant/BQuant. After #10921. GPU-validated (fp8/bf8; i4 deferred).

#10928 and #10927 currently target develop and still contain #10921's shared files (they were split off the combined branch). After #10921 merges, rebase #10928 and #10927 onto develop — their diffs then collapse to just their variant-specific files for a clean review. Do not merge #10928/#10927 before #10921 (it would double-land the shared enablement and force conflicts).

Track B — block-scale quant bridges (independent)

#10439 — block-scale AQuant/ABQuant/BQuant/RowColQuant/TensorQuant.Independent of Track A (verified: 0 file overlap with #10921). It introduces its own bridges and its gfx1250 enablement is self-contained (quant_bridge_common.hpp + gemm_*quant_utils.py). Can merge on its own timeline after review — no ordering constraint vs #10921. All 5 ops GPU-validated on gfx1250 (fp8/bf8/i4; MX sub-variants gfx950-only by design).

Follow-ups — not ready (need C++/kernel work)

TL;DR

#10921  ──► then rebase+merge #10928 and #10927
#10439  ──► independent, merge after review anytime
#10916, #10919  ──► after their C++/kernel follow-ups

ozturkosu added a commit that referenced this pull request Aug 18, 2026
…250 (MI400)

Enables the #10010 grouped rowcolquant + tensorquant quant bridges on gfx1250:
- grouped_gemm_{rowcolquant,tensorquant}_ctypes_lib.cpp: accept gfx1250 in the
  runtime arch gate (+ message). clang-format-18 clean.
- {rowcolquant,tensorquant}_utils.py default_fp8/bf8_config: gfx12 branch using
  the FlatMM 8-bit tile (16x64x256, warp 1x4x1, warp_tile 16x16x128). The stock
  configs use MFMA 32x32x16 / CI 16x16x32, which zero-out on gfx1250 WMMA (same
  trap as #10439/#10927); the FlatMM warp_tile_k=128 path is what runs there.
- test_{rowcolquant,tensorquant}_gpu_correctness.py: add gfx1250 to the harness
  _SUPPORTED_ARCHES so the on-device correctness tests run on MI400.

Validated on MI400 / gfx1250 (a07-1, ROCm 10.1, gpu0): both correctness
harnesses PASS 4/4 (fp8 + bf8, vs fp32 dequant reference).

NOTE: local commit (not pushed). #10010 is @tlakshma's open PR and also has a
pre-existing A@B^T correctness bug independent of this gfx1250 work.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants