feat(ck-tile): gfx1250-only arch-validity gate + gfx1250 default_config for grouped rowcol/tensor quant [WIP] by ozturkosu · Pull Request #11693 · ROCm/rocm-libraries

ozturkosu · 2026-09-04T02:13:34Z

JIRA ID : AICK-2108
JIRA ID : AICK-2141
JIRA ID : AICK-2160

Summary

Stops a gfx9-oriented dispatcher sweep from emitting warp configurations that cannot work on
gfx1250 (MI400). Three pieces:

A corrected Tile Engine warp whitelist for gfx1250, plus a gfx1250 branch in the
warp-tile validation.
A central, opt-in arch-validity gate (codegen_common.arch_config_supported()) consulted
by the sweep-expansion paths.
default_config_gfx1250.json for grouped rowcolquant / tensorquant.

The gate is partial today — see Known gaps, which are stated up front rather than buried.
This is why the PR is still a draft.

Why

gfx1250 is wave32 with RDNA-style WMMA; gfx90a/gfx942/gfx950 are wave64 with MFMA. A gfx9 MFMA
warp tile such as 32x32x32 compiles cleanly on gfx1250 and returns garbage, and 8-warp maps
build but abort at launch with a missing device symbol. default_config.json was written for
gfx9 and nothing filtered it per-arch.

A full default_config sweep of the two grouped quant ops on gfx1250 produced 11,840 rows:

status	count
PASS	3,760
LAUNCH-ABORT	3,220
FAIL (ran, wrong numbers)	2,908
CLEAN-REJECT	1,952

Both failure buckets have a single cause, and it is enumeration blind to the architecture:

100% of the FAILs were warp tile 32x32x32 — an MFMA shape with no WMMA equivalent
(max_rel 0.95–729). That warp tile passed zero times.
100% of the LAUNCH-ABORTs were 8-warp maps (2x4, 4x2), failing with
device symbol missing from .so. Every 2- and 4-warp map had zero aborts.

What changed

Tile Engine whitelist — tile_engine/ops/gemm/gemm_validation_utils.py.
WARP_SUPPORTED_COMBINATIONS["gfx1250"] goes from nine entries
[[2,4,1],[1,8,1],[8,1,1],[4,2,1],[2,1,1],[1,2,2],[4,1,1],[1,4,1],[2,2,1]] down to the four
legal ones [[1,4,1],[2,1,1],[2,2,1],[4,1,1]]. The five removed are the four 8-warp maps
(launch-abort) and [1,2,2] (warp_k=2, compiles and returns wrong results, GPU-measured
max_rel 1.37). A gfx1250 branch is added to the 8-bit warp-tile validation: only 16x16xK
with K in (64, 128) is accepted, since there is no 32x32 WMMA instruction.

Central gate — codegen_common.arch_config_supported(), consulted by the grouped
rowcolquant / tensorquant codegen, unified_contraction_multi_abd,
batched_contraction_utils.expand_sweep, gemm_utils.expand_sweep, and
codegen_common.iter_quant_axes. It is opt-in: ARCH_VALIDITY_RULES = {"gfx1250": {...}},
and any arch without a row returns True on the first statement.

default_config_gfx1250.json for both ops — warp tile 16x16x{64,128}, the four legal warp
maps, block tiles 64..256, persistent pinned false. 736 kernels per (op, dtype).

Also: _cshuffle_store_ok was defined twice in gemm_utils.py with byte-identical bodies; the
later copy was the live one and the earlier "shared helpers" copy was dead. Removed the duplicate.

Gate hardening — three follow-up fixes made on this branch after the first review pass:

arch_config_supported() now normalizes the arch string with the shared
normalize_gfx_arch() before looking anything up. Previously it did
ARCH_VALIDITY_RULES.get(arch or "") on the raw string, so a real device name such as
gfx1250:xnack- found no row and returned True immediately — the gate was inert on the most
common real-world spelling of the only architecture it gates.
The warps-per-block cap is now arch-wide rather than keyed to ("compv3", "intrawave")
alone. The measured failure was cannot find symbol — no launchable kernel entry emitted for
an 8-warp block at all — which is a wave32 block-size property of the target, not of the
pipeline scheduled into it, and the Tile Engine whitelist in this same PR already asserts it
unconditionally. A pair-keyed cap left the mem pipeline admitting exactly the maps measured
to abort. Pair keys remain available for genuinely pipeline-specific evidence.
gemm_utils.py imported the gate behind a bare except Exception whose handler substituted a
constant-True stub. Any defect inside codegen_common — not just the missing-path case the
fallback exists for — silently disabled the gate at import time. Narrowed to ImportError and
made loud with a RuntimeWarning.

None of this changes gfx9, structurally: no gfx9 arch has a row in ARCH_VALIDITY_RULES.
And none of it changes what this PR emits: enumerating the shipped default_config_gfx1250.json
axes through the gate yields the identical 16 accepted combinations before and after, because
those configs are compv3/intrawave only.

Known gaps

These are real and unresolved in the current head. They are the reason the effectiveness claim
above is scoped to "partial". Two gaps previously listed here (the un-normalized arch lookup and
the pipeline-keyed warps-per-block cap) have since been fixed on this branch; see
Gate hardening above.

Two disagreeing warp tables ship in the same PR.
dispatcher/codegen/arch_specs_generated.py (~line 46) carries the nine-entry gfx1250
list, including [1,2,2] and all four 8-warp maps. To be precise about attribution: that
generated list is inherited unchanged from develop — this PR's diff to
arch_specs_generated.py touches only the fp8/bf8 warp-tile rows, and the Tile Engine table
is the only warp list this PR narrows. What this PR does newly do is write the same
nine-entry list into arch_specs.json, which had no gfx1250 block at all before, so the PR
enshrines as source-of-truth a list it did not originate.

This matters because arch_specs_generated.py is what the central gate actually reads:
arch_config_supported() resolves warp_map through _get_arch_data()["warp_combos"]. So
the gate's warp-map rule does not reject [1,2,2] on gfx1250. The 8-warp maps are now
rejected, but by the warps-per-block cap rather than by the warp-map rule.

Narrowing the generated table is a follow-up. It would be emission-neutral for the two
configs this PR ships — enumerating default_config_gfx1250.json through the gate already
yields exactly [1,4,1] [2,1,1] [2,2,1] [4,1,1], since warp_k is pinned to 1 there and the
cap removes the 8-warp maps — but it is a two-file change (arch_specs.json + regeneration)
whose blast radius extends to every consumer of warp_configs, and regenerating it would
detach the validation numbers below from the SHA they were measured at. It is deliberately not
folded in here.
[1,2,2] at max_rel 1.37 is quoted from a prior sweep, not from the campaign below.
The entry's presence in the tables is verifiable in-tree; the measurement itself should be
traced to a specific run before it is relied on in review.
History is still stacked on feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043 even though the base has been retargeted to develop.
Several commits belong to feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043 rather than to this change. This PR should land after
feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043, or be rebuilt before review — and since feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043 will most likely land as a squash,
rebuilt rather than rebased.

gfx9 must not shrink — and does not

Emitted kernels from the raw default_config.json expansion:

arch	develop	this branch
gfx90a	35,496	35,496
gfx942	35,496	35,496
gfx950	35,496	35,496
gfx1250	35,496	768

An earlier revision of this work did shrink gfx9, by applying WARP_SUPPORTED_COMBINATIONS
as if it were a hardware closure. It is a curated whitelist, so the loss was exactly proportional
to table size (gfx90a 3 rows → −23,832; gfx942 7 rows → −9,144; gfx950 9 rows → −6,624). The gate
is now opt-in per arch, which makes that structurally impossible rather than merely absent.

Further evidence: 790,272 full-signature combinations per gfx9 arch, 0 rejected;
batched_contraction 37,248 → 37,248; multi_abd 11,616 → 11,616; the five quant families
identical; gemm_utils unchanged on all three gfx9 arches.

Validation — gfx1250 / MI400, ROCm 10.1

Built from a clean checkout of this branch (git status --porcelain empty), not a patched tree.

2,944 kernels × 12 shapes = 35,328 rows: 21,024 PASS / 14,304 CLEAN-REJECT / 0 FAIL /
0 LAUNCH-ABORT / 0 FAILED-TO-COMPILE. Worst global max_rel 4.36e-4 against a 5e-2
tolerance; lowest non-zero fraction 0.999992. Shapes include non-square, differing quant-axis
extents, tile-aligned and unaligned, an operand in [256, 448], and M % 4 != 0 cases.

No regression: against the 3,760 PASS rows of the pre-change baseline, 3,712 still PASS and
0 went PASS → non-PASS. The 48 not emitted are 8-warp maps, deliberately removed — in that same
baseline 3,220 rows launch-aborted and 100% of them were 8-warp.

Test suites versus a clean develop worktree, same command: 16 failed / 975 passed on the host
suite, 32 failed / 983 passed on the device suite — failure sets byte-identical to develop,
zero new failures. The 8 fewer skips are gfx1250 tests that now actually run.

Correctness used a global metric, max|diff| / (max|ref| + 1e-6). A per-element relative
error is the wrong default here: it explodes near zero and once reported max_rel 27 on
numerically fine output.

What is NOT validated

The gfx9 no-op is enumeration-level evidence only. No gfx90a / gfx942 / gfx950 device was
available. The argument is strong (opt-in table; 790k combinations rejected zero) but it is not
a device run.
The remaining table gap (Known gaps 1) is not covered by any test. The measured gfx1250 results
above were produced through the Tile Engine path with the bare gfx1250 arch string, which is
exactly the configuration in which the gate is effective. The gfx1250:xnack- spelling and
the mem pipeline now have host-test coverage but no device run; a [1,2,2] request routed
through the generated table is still admitted by the warp-map rule.
A real kernel defect is diagnosed but not fixed: rowcolquant zeroes row M−1 when
M % 4 != 0. gemm_quant_kernel.hpp builds the AQ/BQ broadcast views with no padding for the
quant tail, unlike A/B/C. Output is ~99.7% correct, so an all-zero guard and a 99%-non-zero
threshold both miss it; only a reference comparison catches it. feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043 makes it a loud
rejection rather than a silent wrong answer; the kernel fix needs gfx9 hardware and belongs in
its own change.
M = 1021 escapes that rule on 14 of 16 (warp map, tile_m) cells — unexplained.
rcr layout only; no split-K; no performance or parity numbers.

Related: #11043, #11042, #10934.

… gfx1250 (MI400) Re-applies the gfx1250 enablement of PR #11043 on top of current develop. The original branch predates #10439, which merged the bridges themselves into develop, so only the gfx1250 delta is carried over here. - codegen_common.py: add ROWCOL_TENSOR_QUANT_DEFAULT_TILE_GFX12 and the rowcol_tensor_quant_default_tile(gfx_arch) selector. The stock default tile is the gfx9 MFMA 32x32x16 fragment, which does not exist on gfx12 WMMA hardware: the kernel compiles but produces all-zero output. gfx12 needs the FlatMM 8-bit tile 16x64x256 / warp 1x4x1 / warp_tile 16x16x128. - grouped_gemm_{rowcolquant,tensorquant}_utils.py: _default_config() selects the tile through that helper. gfx942/gfx950 kernel names are unchanged. - grouped_gemm_{rowcolquant,tensorquant}_ctypes_lib.cpp: add gfx1250 to kSupportedArchs, which otherwise rejects the .so in dispatcher_initialize with "compile-time GFX_ARCH 'gfx1250' is not a supported architecture". - test_{rowcolquant,tensorquant}_gpu_correctness.py: add gfx1250 to _SUPPORTED_ARCHES, and default --gfx to the detected device instead of a hardcoded gfx950 (which built a gfx950 image and aborted on any other GPU). Validated on gfx1250 / MI400 (ROCm 10.1): both harnesses PASS fp8 + bf8.

…silent failures Addresses review feedback on #11043 plus two defects found while validating the enablement on gfx1250 hardware. All three share a failure mode: the code was wrong in a way that reported success. 1. Arch gate silently skipped on supported devices. The GPU correctness tests gated on `_GFX_ARCH in _SUPPORTED_ARCHES`, an exact string match. rocm_agent_enumerator may report feature suffixes, so a supported device enumerating as "gfx942:sramecc+:xnack-" caused the tests to SKIP. A skip is not a failure, so CI reported green while running none of the GPU coverage. Both test modules now normalize the arch (strip at the first ':') and compare the whole base token, matching the prefix tolerance the C++ bridges already had in is_supported_arch(). Token comparison rather than startswith, so a hypothetical "gfx9421" is not accepted on the strength of the "gfx942" prefix. Pinned by a new CPU-only test module. 2. RowColQuant silently zeroed row M-1 of C when M % 4 != 0 on gfx1250. MakeAQBlockWindow() builds the AQ view as a broadcast make_naive_tensor_view({M, N}, {1, 0}) and then creates the tile window with no padding transform, unlike the A/B/C windows which pad from kPadM/kPadN/kPadK. The M axis therefore has an unguarded tail. TensorQuant is correct at the same M because it passes scalar scales; the BQ view has the same unguarded tail but is unreachable because pad_n=false already forces N % 64 == 0. The root-cause fix belongs in the shared ck_tile quant kernel header, which is also compiled for gfx942/gfx950 -- untested here. Rather than change tile distribution for architectures that were not measured, the bridge now rejects these shapes with an explicit message. The guard is scoped to gfx12 targets via the compile-time GFX_ARCH and to RowColQuant only, and its comment records the condition for removing it. 3. Undocumented argument constraints returned a bare -2. N % TileN is now pre-checked in both bridges, derived from the kernel's own kPadN/TileN constants rather than hardcoded. The -2 message enumerates the known constraints, including that K % 16 == 0 is required even with pad_k=true (pad_k covers the K-loop tail, not the global-load vector width). Constraints documented on the ABI comment and in both run() docstrings, and the Python RuntimeError is now self-contained instead of deferring to stderr. K % 16 is reported but not hard-enforced: it is checked inside ck_tile's IsSupportedArgument() and has not been shown to be general across dtypes and tiles, so enforcing it could reject shapes the kernel would accept. Verified on gfx1250 / MI400: both GPU correctness harnesses 2/2 at max_rel_err=0.0005; every M % 4 == 0 shape unchanged at 2.2e-4 to 3.9e-4; every M % 4 != 0 shape now a clean rejection instead of a wrong answer; TensorQuant still correct at those same M. Host-side suite 59 -> 95 passing, 0 failures. clang-format-18 clean. Known gap, unchanged: the codegen's _default_config() is still arch-agnostic, so a CMake-built lib on gfx1250 would compile the gfx9 tile. Only the Python runtime path is arch-aware. Tracked as a follow-up.

…p expansion Sweep expansion had no single place that answered "is this (arch, warp map, warp tile, dtype, pipeline) combination valid?", so each operator re-derived its own partial answer -- and three of them had no arch input at all. A gfx1250 default_config sweep of grouped_gemm_rowcolquant + grouped_gemm_ tensorquant (11,840 rows) shows the cost: 2,908 rows returned wrong numbers and 100% of them were warp_tile 32x32x32, a wave64 MFMA shape that does not exist on wave32 WMMA hardware (0 of 3,760 16x16x64 rows were wrong); another 3,220 rows aborted at launch and 100% of them were 8-warp blocks. Both are enumeration defects, not kernel defects. codegen_common.arch_config_supported() is now the one authoritative gate, sitting next to the existing valid_wave_configs/valid_warp_configs helpers that already own the arch tables from arch_specs.json. It applies three data-driven rules: the Old-TE warp-map table, the per-arch warp-tile table (enforced only where that table is a hardware closure, i.e. the wave32 WMMA families -- the CDNA rows are curated whitelists and stay advisory), and a per-arch warps-per-block cap table. No rule is expressed as a gfx string test; adding an arch is a row in arch_specs.json. Wired into the expansions that previously bypassed it: the grouped rowcolquant/tensorquant _build_specs(), unified_contraction_multi_abd build_specs(), batched_contraction_utils.expand_sweep(), and codegen_common iter_quant_axes() (which covers the five remaining quant families). gemm_utils.expand_sweep() now delegates its inline gfx1250 8-warp rule to the same helper; its warp-map gate deliberately stays on the Old-TE table, whose contract is byte-parity with Old-TE rather than arch_specs.json. Also: add gfx1250 to arch_specs.json. It was present in the generated module but absent from the JSON source, so any regeneration would have deleted it. Its 8-bit warp-tile row gains 16x16x128, the GPU-validated gfx12 shape. Also: remove the second, byte-identical definition of _cshuffle_store_ok in gemm_utils.py. Python keeps the last definition, so that copy was live and the canonical one in the shared-helpers block was dead code. Emitted-kernel regression gate, grouped rowcolquant and tensorquant, fed by the Old-TE instance builder's own default_config expansion: gfx942 592 -> 592 (0 removed, 0 added) for both ops gfx950 600 -> 600 (0 removed, 0 added) for both ops gfx1250 1184 -> 464 for both ops; every one of the 720 removed kernels is warp_tile 32x32x32 or an 8-warp block, and nothing else is removed. Exhaustive over the GEMM default_config axis space (12,150 points x 5 dtypes x 3 pipelines x 2 schedulers), the new rules reject 0 configurations on gfx942 and gfx950.

…uant configs/default_config.json for grouped_gemm_rowcolquant and grouped_gemm_ tensorquant describes gfx9 MFMA hardware. Expanded for gfx1250 it offers 45 warp-tile combinations where the hardware has one, and a GPU sweep of the result (11,840 rows) returned 2,908 silently-wrong answers -- 2,908/2,908 of them warp_tile 32x32x32, a wave64 MFMA shape that does not exist on wave32 WMMA -- plus 3,220 launch aborts, 3,220/3,220 of them 8-warp blocks. Adds default_config_gfx1250.json per op, same schema, sized to what gfx1250 can actually run. Expansion: 184 tile configs x 4 trait combinations = 736 kernels per (op, dtype). Axes, and why each is what it is (all statements GPU-measured on gfx1250): warp tile 16x16x{64,128}. 32x32xK is removed: there is no 32x32 WMMA. 16x16x128 is included on evidence, not on the arch table -- V_WMMA_*_16x16x128 kernels were built and run and agree with a CPU fp32 reference to max_rel 4e-4, identically to 16x16x64. warp map the 1x{1,2,4} / 2x{1,2,4} / 4x{1,2,4} product, of which the divisibility and arch rules keep 1x4x1, 2x1x1, 2x2x1, 4x1x1. Every >4-warp map aborts at launch and [1,2,2] returns wrong results, so both are removed from the arch table below rather than worked around in the config. block tile 64..256 step 64 on all three axes; 29 of the 64 combinations survive divisibility against the warp map, 46 once 16x16x128 is allowed. pipeline / pinned to compv3 / intrawave / cshuffle, which is not a scheduler / collapse this config can undo: the codegen emits no other epilogue pipeline or epilogue for these ops, and compv3+interwave is accepted by codegen and then fails to compile on gfx1250 ("no member named 'operator()'" in GemmPipelineAgBgCrCompV3). pad pad_m and pad_k are opened; all eight pad combinations were verified to run correctly. pad_n stays at the shipped runtime default (false). persistent left false. persistent=true runs correctly at 128x128x128 but at 64x256x128 / 2x1x1 it hangs the GPU hard enough that the process is unkillable, so it is not safe to sweep by default. Two arch tables in gemm_validation_utils.py disagreed with the hardware and are corrected so the config cannot express an invalid kernel: WARP_SUPPORTED_COMBINATIONS["gfx1250"] listed four 8-warp maps and [1,2,2]. All five are unusable (launch abort / wrong results). The remaining four are exactly what the central arch-validity gate in codegen_common already accepts, so the two now agree. _validate_fp8_mfma_warp_tile_k applied the gfx942 MFMA K-block table to gfx1250, which accepts 32x32x32 (the wrong-answer shape) and rejects 16x16x128 (a valid one). gfx1250 now has its own WMMA branch. No behaviour change on gfx9: both edits are inside gfx1250-only branches.

…ers/muozturk/ck/gfx1250-combined-gate

… into users/muozturk/ck/gfx1250-combined-gate

…nstruction The previous revision reduced gfx9 coverage. Under a raw expansion of the grouped rowcolquant/tensorquant default_config ranges (tile 64..256 step 64, warp_m/n {4,2,1}, warp_k {1}, warp_tile_m {4,16,32}, warp_tile_n {16,32,64}, warp_tile_k {8,16,32,64,128}, divisibility-valid, fp8+bf8, rcr) fed straight into _build_specs, it deleted 9,144 of 35,496 gfx942 candidates and 6,624 of 35,496 on gfx950. That is real, it is not a measurement artifact, and it was missed because the accepted regression gate used a different denominator -- the Old-TE instance builder's own tile expansion -- under which the same code is byte-identical on gfx9 (592 -> 592, 600 -> 600). Cause: rule 1, the warp-map rule, ran on every arch. WARP_SUPPORTED_COMBINATIONS is a curated Old-TE whitelist rather than the set of warp maps the hardware can run -- gfx942 has no [4,2,1] / [2,4,1] / [4,4,1] row -- so applying it as a closure deletes working kernels. It is also what produced the 3,312 8-warp gfx942 rejections observed on a direct probe of the predicate; those did not come from the warps-per-block cap, which was already gfx1250-only. The warp-tile rule was already confined to gfx1250 in effect, via a family closure (ARCH_FAMILY_MAP -> "rdna*" => the table is exhaustive). Correct, but the wrong shape for this requirement: it makes the gfx9 no-op a property of the current table contents rather than a guarantee, and gfx1100/gfx1200/gfx1201 would have been swept in by the same family test with no evidence behind them. Replaced with an explicit per-arch opt-in table, ARCH_VALIDITY_RULES. An arch is gated if and only if it has a row; arch_config_supported() returns True on its first statement otherwise, so no rule -- present or future -- can reach an ungated arch. Only gfx1250 has a row. The family-closure helpers (_arch_family, _is_wave32_arch, _warp_tile_table_is_exhaustive), the standalone PIPELINE_MAX_WARPS_PER_BLOCK table and the now-unused ARCH_FAMILY_MAP plumbing in _get_arch_data() are all removed; each row now carries its own rule flags. Proof under BOTH denominators, before = origin/develop, after = this branch: raw range expansion, per op (rowcolquant and tensorquant identical) gfx942 35,496 -> 35,496 name sets byte-identical (cmp) gfx950 35,496 -> 35,496 name sets byte-identical (cmp) gfx1250 35,496 -> 768 unchanged by this commit Old-TE instance-builder expansion, per op gfx942 592 -> 592 0 removed, 0 added gfx950 600 -> 600 0 removed, 0 added gfx1250 1,184 -> 464 720 removed, unchanged by this commit exhaustive predicate probe, 14,580 points per arch gfx90a / gfx942 / gfx950 / gfx1100 / gfx1200 / gfx1201 / unknown: 0 rejected gfx1250: 11,316 rejected The gfx1250 result is untouched: the 720 removed kernels are still exactly 360 warp_tile 32x32x32, 128 both, 232 8-warp, 0 for any other reason, and 0 of the 464 survivors is 32x32x32 or 8-warp. Replaying the gate over the 11,840-row baseline sweep still gives FAIL 2,908 -> 0 and LAUNCH-ABORT 3,220 -> 0 with 3,712 of 3,760 PASS rows retained (the 48 lost are the documented 8-warp rowcolquant track-B rows). Every other call site is a no-op on gfx9 for the same structural reason, measured per op with byte-identical name sets on gfx90a/gfx942/gfx950: unified_contraction_multi_abd 133,952 -> 133,952 (gfx1250 -> 3,168), batched_contraction_utils.expand_sweep 18,624 -> 18,624 (gfx1250 -> 4,352), codegen_common.iter_quant_axes 50,232 -> 50,232 (gfx1250 -> 1,440), gemm_utils.expand_sweep 320 -> 320 on both shipped CI configs. Tests: test_warp_map_gate_uses_the_arch_table asserted the removed gfx942 behaviour and is gone. Replaced by test_only_gfx1250_is_gated (the opt-in table has exactly one key), test_gfx9_is_a_constant_true_no_op (exhaustive over the axis space on gfx90a/gfx942/gfx950) and test_ungated_wave32_archs_are_also_- untouched. Suite unchanged otherwise: 16 failed / 927 passed, identical FAILED set to origin/develop, zero new failures. (cherry picked from commit 96d7e01)

therock-pr-bot · 2026-09-04T02:17:55Z

✅ 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?

Copilot

🟡 Changes recommended

The central arch_config_supported() gate should normalize feature-suffixed arch strings (and a couple of Python error messages currently hard-code TileN=64), otherwise gfx1250 gating/diagnostics can be incorrect in real-world suffix-reporting environments.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

This PR adds gfx1250-specific (MI400 / RDNA4 wave32 WMMA) safeguards and defaults to prevent sweep expansion/codegen from emitting kernel configurations that are known to be invalid on that architecture, while keeping gfx9-class arches effectively unchanged.

Changes:

Introduces a centralized, opt-in arch_config_supported() gate (currently only active for gfx1250) and wires it into multiple sweep/codegen expansion paths.
Adds gfx1250-tuned default config JSON sweeps for grouped RowColQuant and TensorQuant Tile Engine ops.
Improves gfx1250 enablement and diagnostics across Python GPU tests and C++/Python dispatcher bridges (including suffix-tolerant arch detection and clearer constraint reporting).

File summaries

File	Description
projects/composablekernel/tile_engine/ops/gemm/grouped_gemm_quant/grouped_gemm_tensorquant/configs/default_config_gfx1250.json	Adds a gfx1250-focused sweep config for grouped tensorquant.
projects/composablekernel/tile_engine/ops/gemm/grouped_gemm_quant/grouped_gemm_rowcolquant/configs/default_config_gfx1250.json	Adds a gfx1250-focused sweep config for grouped rowcolquant.
projects/composablekernel/tile_engine/ops/gemm/gemm_validation_utils.py	Narrows gfx1250 warp-map combos and adds WMMA-aware fp8 warp-tile validation.
projects/composablekernel/dispatcher/tests/test_tensorquant_gpu_correctness.py	Makes arch gating suffix-tolerant and removes hardcoded `--gfx` default.
projects/composablekernel/dispatcher/tests/test_rowcolquant_gpu_correctness.py	Same suffix-tolerant arch gating and `--gfx` default behavior update.
projects/composablekernel/dispatcher/tests/test_gpu_test_arch_gate.py	Adds CPU unit tests for the Python arch-gate predicate behavior.
projects/composablekernel/dispatcher/tests/test_codegen_common.py	Adds unit tests pinning the central arch-validity gate’s opt-in behavior and gfx1250 rules.
projects/composablekernel/dispatcher/python/grouped_gemm_tensorquant_utils.py	Routes default tile selection through an arch-aware selector; improves error context.
projects/composablekernel/dispatcher/python/grouped_gemm_rowcolquant_utils.py	Same arch-aware default tile routing; improves error context (plus gfx12 M%4 constraint explanation).
projects/composablekernel/dispatcher/python/gemm_utils.py	Removes duplicate helper definition and routes sweep expansion through the central arch-validity gate.
projects/composablekernel/dispatcher/python/batched_contraction_utils.py	Adds optional `arch` parameter to sweep expansion and applies the central arch gate.
projects/composablekernel/dispatcher/codegen/unified_grouped_gemm_tensorquant_codegen.py	Adds `--arch` support and applies the central arch-validity gate during spec generation.
projects/composablekernel/dispatcher/codegen/unified_grouped_gemm_rowcolquant_codegen.py	Same `--arch` support and central gating for rowcolquant spec generation.
projects/composablekernel/dispatcher/codegen/unified_contraction_multi_abd_codegen.py	Adds arch-aware gating to contraction_multi_abd spec expansion.
projects/composablekernel/dispatcher/codegen/codegen_common.py	Defines `ARCH_VALIDITY_RULES`, `arch_config_supported()`, and uses it in shared quant axis iteration; adds gfx12 default tile selector.
projects/composablekernel/dispatcher/codegen/arch_specs.json	Adds/updates gfx1250 arch spec data (warp size, warp configs, warp tile combos).
projects/composablekernel/dispatcher/codegen/arch_specs_generated.py	Regenerates generated arch tables to include gfx1250 fp8/bf8 warp-tile combos (64/128).
projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_tensorquant_ctypes_lib.cpp	Adds gfx1250 support and improves constraint checks/messages (e.g., N%TileN).
projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_rowcolquant_ctypes_lib.cpp	Adds gfx1250 support and improves constraint checks/messages (including gfx12-scoped M%4 guard).

Review details

Files reviewed: 19/19 changed files
Comments generated: 3
Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

+    rules = ARCH_VALIDITY_RULES.get(arch or "")
+    if not rules:
+        return True
+
+    data = _get_arch_data()


+                f"for kernel {self.kernel_name} at M={M} N={N} K={K}. "
+                f"(-1 = rejected by the bridge, -2 = rejected by the kernel, "
+                f"-3 = launch threw.) Shape constraints: K % 16 == 0 (required even "
+                f"with pad_k=True); N % 64 == 0 when pad_n=False; and on gfx12 targets "


+                f"for kernel {self.kernel_name} at M={M} N={N} K={K}. "
+                f"(-1 = rejected by the bridge, -2 = rejected by the kernel, "
+                f"-3 = launch threw.) Shape constraints: K % 16 == 0 (required even "
+                f"with pad_k=True); N % 64 == 0 when pad_n=False. "


…250 tile match The parent branch renamed ROWCOL_TENSOR_QUANT_DEFAULT_TILE_GFX12 to ..._GFX1250 and replaced the `"gfx12" in gfx_arch` substring test in rowcol_tensor_quant_default_tile() with an exact match on the normalized target, because gfx1200/gfx1201 are WMMA parts whose 8-bit warp fragment is 16x16x16 rather than gfx1250's 16x16x64 / 16x16x128 -- a family match there compiles cleanly and returns garbage. This branch still carried the old name and the substring test, so the two halves of the same stack disagreed about which parts get the gfx1250 tile. Adopt the parent's version verbatim: * normalize_gfx_arch() is the single source of truth for stripping feature suffixes ("gfx1250:xnack-" -> "gfx1250"); the two grouped GPU-correctness tests now import it from codegen_common instead of each defining a private copy. * rowcol_tensor_quant_default_tile() matches gfx1250 exactly. * The OCP-FP8 hipcc define stays deliberately family-wide: gfx1200 and gfx1201 do use OCP FP8, so the two sites need opposite treatment. Tests: dispatcher/tests/test_codegen_common.py 131 passed, 7 subtests.

…rp cap arch-wide Two ways the gate could be present and inert. 1. arch_config_supported() looked the target up exactly as handed in: rules = ARCH_VALIDITY_RULES.get(arch or "") if not rules: return True Callers pass whatever rocm_agent_enumerator or gcnArchName reported, and on a real part that is routinely "gfx1250:xnack-". That string has no row, so the function returned True before any rule could run -- the gate was a no-op on the most common real-world spelling of the only architecture it gates. It now normalizes with normalize_gfx_arch() before the lookup, which also keeps the ARCH_VALIDITY_RULES, WARP_SUPPORTED_COMBINATIONS and WARP_TILE_SUPPORTED_COMBINATIONS lookups consistent -- all three are keyed by the bare target. 2. The warps-per-block cap was keyed {(pipeline, scheduler): cap} and only registered for ("compv3", "intrawave"), so every other pipeline -- mem in particular -- admitted exactly the 8-warp maps that were measured to abort. The measured failure is "cannot find symbol": no launchable kernel entry was emitted for an 8-warp block at all, which is a wave32 block-size property of the target rather than a property of the pipeline scheduled into it. The Tile Engine whitelist for gfx1250 in this same change already asserts it unconditionally. The cap is now registered arch-wide under a None key; pair keys remain available for genuinely pipeline-specific evidence, and the lookup falls back None-ward. It also no longer requires pipeline/scheduler to be supplied: a caller that omits its trait pair must not thereby escape a cap that describes what the target can launch. gfx9 is untouched, by construction rather than by inspection: no gfx9 arch has a row in ARCH_VALIDITY_RULES, so no rule -- present or future -- can reach it. test_gfx9_is_a_constant_true_no_op covers that exhaustively over the sweep axis space, including the mem pipeline. Emission is unchanged for the two configs this PR ships. Enumerating the grouped rowcolquant/tensorquant default_config_gfx1250.json axes through the gate yields the identical 16 accepted combinations before and after -- the four warp maps [1,4,1] [2,1,1] [2,2,1] [4,1,1] x 2 dtypes x warp_tile_k in {64,128} -- because those configs are compv3/intrawave only, where the pair-keyed cap already applied. The change closes the hole for every other caller. Tests: 137 passed, 37 subtests (test_codegen_common.py, test_gpu_test_arch_gate.py). Full non-GPU dispatcher suite: 990 passed, identical 16-failure set to the parent commit (all require hipcc or a device).

…lure gemm_utils imports the central arch-validity predicate behind a bare `except Exception`, whose handler substitutes a stub that returns True unconditionally. The fallback exists for one narrow reason -- keeping expand_sweep importable when dispatcher/codegen is not on sys.path -- but as written any failure inside codegen_common (a syntax error, a bad table edit, a missing name) also lands there, replacing the safety predicate with a constant-True stub, silently, at import time. That is the worst possible failure mode for a gate: it does not stop working loudly, it stops working invisibly. Catch ImportError only, so a defect inside codegen_common surfaces as the exception it is, and warn on RuntimeWarning when the fallback does engage so an ungated run is at least visible in the log. Tests: gemm_utils imports cleanly and binds the real codegen_common.arch_config_supported; full non-GPU dispatcher suite unchanged (990 passed, same 16 pre-existing failures).

therock-pr-bot · 2026-09-09T22:34:23Z

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

… the data The comment on the narrowed gfx1250 warp-map list said an 8-warp block "has no launchable kernel entry" and that every such map aborts, 3,220/3,220. Recounting the 14,208-row sweep archive, that is not what was measured, and the two failure modes were confounded: 8-warp + 16x16x64 (legal tile): 2168 abort, 424 reject, 192 pass, 0 FAIL 8-warp + 32x32x32 (illegal tile): 1184 abort, 240 reject, 112 FAIL <=4-warp + 32x32x32: 1206 reject, 3114 FAIL All 3,226 wrong-result rows use 32x32x32, and no other warp tile produces one. 8-warp maps paired with a legal tile never returned a wrong answer; they abort, cleanly reject, or pass, and the split moves between runs. So the warp tile is the silent-corruption case and the warp map is the dependability case. Both stay rejected; the comments now say which is which. The previously quoted 3,220 and 2,908 figures also came from a different, smaller sweep than the archive they were attributed to. Comment text is byte-identical to the same correction on #11043, so the two branches do not diverge here.

ozturkosu added 7 commits September 1, 2026 18:54

Merge branch 'users/muozturk/ck/central-arch-validity-filter' into us…

7fb733d

…ers/muozturk/ck/gfx1250-combined-gate

Merge branch 'users/muozturk/ck/gfx1250-grouped-quant-default-config'…

9f2425c

… into users/muozturk/ck/gfx1250-combined-gate

ozturkosu changed the base branch from users/muozturk/ck/rowcolquant-tensorquant-gfx1250-enable to develop September 4, 2026 02:14

github-actions Bot added the project: composablekernel label Sep 4, 2026

assistant-librarian Bot added the organization: ROCm label Sep 4, 2026

ozturkosu self-assigned this Sep 4, 2026

ozturkosu requested a lite review from Copilot September 4, 2026 03:01

Copilot started reviewing on behalf of ozturkosu September 4, 2026 03:02 View session

Copilot AI reviewed Sep 4, 2026

View reviewed changes

ozturkosu changed the title ~~feat(ck-tile): gfx1250-only arch-validity gate + gfx1250 default_config for grouped rowcol/tensor quant~~ feat(ck-tile): gfx1250-only arch-validity gate + gfx1250 default_config for grouped rowcol/tensor quant [WIP] Sep 4, 2026

ozturkosu added 3 commits September 9, 2026 18:23

Navigation Menu

feat(ck-tile): gfx1250-only arch-validity gate + gfx1250 default_config for grouped rowcol/tensor quant [WIP] - #11693

feat(ck-tile): gfx1250-only arch-validity gate + gfx1250 default_config for grouped rowcol/tensor quant [WIP]#11693
ozturkosu wants to merge 11 commits into
developfrom
users/muozturk/ck/gfx1250-arch-validity-and-config

ozturkosu commented Sep 4, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 4, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

therock-pr-bot Bot commented Sep 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

ozturkosu commented Sep 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Why

What changed

Known gaps

gfx9 must not shrink — and does not

Validation — gfx1250 / MI400, ROCm 10.1

What is NOT validated

Uh oh!

therock-pr-bot Bot commented Sep 4, 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!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

therock-pr-bot Bot commented Sep 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

ozturkosu commented Sep 4, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 4, 2026 •

edited

Loading