Repository navigation
feat(ck-tile): grouped rowcolquant/tensorquant GEMM bridges on gfx1250 [Experimental- Dont Review] - #11904
Closed
ozturkosu wants to merge 27 commits into
Closed
feat(ck-tile): grouped rowcolquant/tensorquant GEMM bridges on gfx1250 [Experimental- Dont Review]#11904ozturkosu wants to merge 27 commits into
ozturkosu wants to merge 27 commits into
Conversation
… 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.
The RowColQuant/TensorQuant default-tile selector matched the whole gfx12
family. gfx1200/gfx1201 are WMMA parts too, but their 8-bit warp fragment is
16x16x16, not gfx1250's 16x16x64/16x16x128, so they would have inherited a
tile that compiles cleanly and returns garbage. Nothing upstream stops that:
the runtime helpers have no arch allowlist. Match gfx1250 exactly instead.
The OCP-FP8 hipcc define stays family-wide on purpose -- every gfx12xx part
uses OCP FP8 encoding -- and is now written as an explicit family test with a
comment, so the two sites cannot be mistaken for the same rule.
Add a shared normalize_gfx_arch() in codegen_common as the single source of
truth for stripping feature suffixes ("gfx1250:xnack-"), use it at the arch
detection boundary so suffixes never reach --offload-arch/-DGFX_ARCH (the C++
side prefix-matches, so the bare target still matches a suffixed device), and
drop the duplicated local copies in the two GPU correctness tests.
The C++ derives its N-divisibility bound from SelectedKernel::TileN, but the
Python docstring and RuntimeError hardcoded 64. Export dispatcher_get_tile_n()
and dispatcher_get_pad_n() from both ctypes libs and use the real values, so
the diagnostics stay true if the compiled tile changes.
Tests cover suffix normalization, exact-vs-family arch matching for the tile
selector, and the intentional gfx12-family OCP-FP8 behaviour.
… arch The grouped RowColQuant/TensorQuant .so cache was keyed on (kernel name, arch), which says nothing about the exported symbol set. A caller with a persistent output_dir predating the dispatcher_get_tile_n()/dispatcher_get_pad_n() exports was handed that stale artifact, and the runner then died at attribute lookup with a bare "undefined symbol". Default usage is tempfile.mkdtemp() so CI never saw it, but the validation harnesses use persistent directories precisely to get cache hits. Fold an _SO_ABI revision into the artifact filename so a pre-ABI .so simply cannot be selected -- rejecting the stale artifact rather than masking it with a symbol-lookup fallback, which would hide the staleness instead of ending it. Reproduced on gfx1250 against a real pre-ABI library built from the parent commit's ctypes source: it raises "undefined symbol: dispatcher_get_tile_n" under the old name and is ignored under the new one. normalize_gfx_arch() was applied only to the detection branch, so a caller passing gfx_arch="gfx1250:xnack-" sent the suffixed string straight into --offload-arch, -DGFX_ARCH and the cache name. Normalize once at each entry boundary instead: the setup function, the compile helper, and _default_config (reached from the public default_fp8_config/default_bf8_config). The GPU correctness tests derived the host fp8 encoding from the module-level detected arch rather than the arch being compiled for, so with --gfx pointing at a different target the host bytes and the kernel disagreed and decoded to NaN/Inf on device. Thread the arch through _fp8_ml_dtype/_encode_fp8/_decode_fp8 and _make_inputs. The gfx12-family test inside _fp8_uses_ocp is correct as-is (gfx1200/gfx1201 use OCP FP8 too) and is unchanged. Register test_codegen_common.py and test_gpu_test_arch_gate.py with CTest; both existed but nothing ran them. Validated on gfx1250: 4/4 kernels compile with both exports present, worst global max_rel 4.22e-4 across rowcolquant/tensorquant x fp8/bf8 over five shapes each, M%4!=0 cleanly rejected for RowColQuant and correct for TensorQuant, reported tile_n=64 matches the compiled kernel, and a suffixed target builds and runs end to end. Dispatcher suite 212 passed / 0 failed on device.
…dation The sweep validity tables were written for gfx9 and let two families of configuration through on gfx1250, both of which are worse than a build error. WARP_SUPPORTED_COMBINATIONS listed nine gfx1250 warp maps. Four of them are 8-warp blocks -- [2,4,1] [4,2,1] [1,8,1] [8,1,1] -- and gfx1250 is wave32, so a block of more than four warps has no launchable kernel entry: 3,220 of 3,220 such rows aborted at launch in a GPU sweep. A fifth, [1,2,2], compiles and then returns wrong results (GPU-measured max_rel 1.37). Only [1,4,1] [2,1,1] [2,2,1] [4,1,1] are legal. The 8-bit warp-tile check had no gfx1250 case, so it fell through to the gfx9 MFMA expectations. gfx1250 is RDNA-style WMMA: the only 8-bit fragments are 16x16x64 and 16x16x128, and there is no 32x32 WMMA at all, so a 32x32xK tile compiles and returns garbage -- 2,908 of 2,908 wrong-result rows in that same sweep were 32x32x32. The reverse was also true: 16x16x128, which is the tile this bridge actually ships, was being rejected. Measured on the real predicate, gfx1250 fp8 rowcol/tensor quant: 32x32x32 ACCEPT -> REJECT, 16x16x128 REJECT -> ACCEPT, 16x16x64 unchanged. gfx942 and gfx950 are bit-for-bit unchanged across the same probe, as is the per-arch entry count of every other WARP_SUPPORTED_COMBINATIONS row. Content is identical to the same fix on the stacked gate branch, so that branch rebases without a semantic conflict.
… invalid tile
Listing gfx1250 in kSupportedArchs makes the runtime arch check accept the target,
but only the Python helper picks an arch-correct tile. The CMake/codegen path does
not: unified_grouped_gemm_{rowcolquant,tensorquant}_codegen.py::_default_config()
takes no arch argument and hard-codes the gfx9 tile. A CMake-built gfx1250 library
would therefore initialize cleanly and then execute a fragment that does not exist
on the device, which compiles and returns garbage rather than failing.
gfx1250 is not reachable through CMake today because DESIRED_TARGETS omits it, but
that is an accident of omission rather than a decision, and it stops protecting the
moment someone adds the target. Make the bridge itself fail closed instead.
The kernel's tile is already a compile-time property of the generated header, so
this is a static_assert rather than a runtime check: an invalid gfx1250 library
cannot be produced at all. The rules mirror the Python validator -- on gfx1250 the
only 8-bit WMMA fragments are 16x16x64 and 16x16x128, there is no 32x32 WMMA, and
a block of more than four warps has no launchable kernel entry on wave32.
The test is exact-gfx1250, not gfx12-family: gfx1200/gfx1201 have a 16x16x16 8-bit
fragment and are not in kSupportedArchs. Prefix matching accepts the suffixed form
("gfx1250:xnack-") that hipDeviceProp_t reports.
Verified against a standalone mock of the generated header, 11 cases: legal gfx1250
tiles and both warp-map shapes compile; 32x32x32, 32x32x16, 16x16x32 and the 8-warp
maps are refused; gfx942, gfx950 and an undefined GFX_ARCH are entirely unaffected.
… data The comment on the gfx1250 warp-map list, and the static_assert message that echoed it, both said an 8-warp block "has no launchable kernel entry" and cited 3,220/3,220 aborted rows. Recounting the 14,208-row sweep archive that the claim rests on, that is not what the data shows: LAUNCH-ABORT 3352 CLEAN-REJECT 664 FAIL 112 <- wrong answers, not aborts PASS 192 <- correct, all 16x16x64, 192 distinct kernels So 8-warp maps are not uniformly broken. They are nondeterministic: re-running the same surface flips hundreds of rows between abort, pass and fail. That is a stronger reason to reject them, not a weaker one -- a configuration that is sometimes correct can survive a clean validation run and reach users, which a configuration that always aborts cannot. Rejecting them therefore stands; only the stated justification changes. The numbers cited previously also came from a different, smaller sweep than the archive they were attributed to. This diverges from the identical comment on the stacked gate branch, which carries the same overstatement and should receive the same correction.
The previous correction fixed one overstatement and introduced another: it said the 8-warp maps "returned wrong results 112 times". Cross-tabulating warp map against warp tile on the same 14,208-row archive shows those 112 rows are not attributable to the warp map at all. 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 in the archive use 32x32x32, and no other warp tile produces one. The two variables were confounded: every "8-warp wrong answer" is a 32x32x32 wrong answer that happens to also be 8-warp. So the two rejections have different justifications, and the comments now say so: 32x32x32 is the sole source of silent wrong answers -- highest severity. 8-warp maps are unreliable, not incorrect: mostly abort, sometimes pass, and the split moves between runs. Still rejected, but for dependability. Both remain rejected; only the reasoning is now attributable to the data.
Both guards added earlier advertised protection they did not fully provide. The static_assert capped warps-per-block at four, but [1,2,2] has a product of four and passed it -- and [1,2,2] is the one warp map this PR documents as compiling and then returning wrong results (max_rel 1.37). The corrupting case was the one getting through. warp_k is now required to be 1, as its own rule rather than a side effect of the cap. WarpTileN was never asserted at all, so 16x32xK also passed; it is now pinned to 16. The validator guards were bypassed by a suffixed architecture name. Both WARP_SUPPORTED_COMBINATIONS.get() and the gfx1250 warp-tile branch compared the raw gpu_target, and an unrecognized key falls through permissively -- so a device reporting "gfx1250:xnack-" disabled the entire new restriction and landed in the branch that accepts 32x32x32, the sole measured source of wrong answers. The file already computed a normalized base_gpu_target twenty lines away; both sites now normalize the same way. Verified against the standalone mock, 9 cases: [1,2,2] and 16x32x128 are now refused, every previously accepted gfx1250 tile still compiles, suffixed gfx1250 still resolves, and gfx942/gfx950 remain unaffected including for [1,2,2].
Narrowing WARP_SUPPORTED_COMBINATIONS["gfx1250"] from nine warp maps to four was a regression. That table is shared by seven instance builders, and one of the maps removed, 4x2x1, is plain GEMM's default gfx1250 WMMA configuration -- issue #11161 reported it returning wrong results, PR #11175 fixed it, and the fix was verified correct on device. Measured on the predicate: validate_warp_configuration(4, 2, 1, "gfx1250") develop ACCEPT after my change REJECT <- deleted another operator's working default The table is now byte-identical to develop again. The restriction itself is still needed, but only here: these two grouped quant bridges do not emit a loadable kernel entry for >4-warp blocks on gfx1250, so those kernels abort at launch with "cannot find symbol" -- 3,220 rows in a GPU sweep. That is a property of this bridge's code generation, not of the hardware, which is precisely why it does not belong in a hardware capability table. It now lives in validate_gemm_rowcol_tensor_quant(), alongside the warp-tile rule that was already scoped to the four quant validators. warp_k > 1 keeps its own rule rather than riding on the warps-per-block cap: [1,2,2] has a product of 4 and would otherwise pass. Verified: plain GEMM's 4x2x1 default is accepted again; this operator still rejects 4x2x1, 2x4x1, 4x4x1 and 1x2x2; 32x32x32 is still rejected; the legal 16x16x128 tile is still accepted; and across gfx90a, gfx942 and gfx950 there are zero decision changes.
…fx12 test Two findings from an independent review of an external code review. The Python comment justifying the warp-map rule cited 3,220 abort rows while the C++ assertion message beside it cited 2,168 for what reads as the same population. Both numbers are real, but they come from different sweeps: 3,220 is from an 11,840-row run, while the C++ message quotes the 14,208-row archive. Recounting that archive, the 4,320 8-warp rows split as 3,352 aborts overall and 2,168 when paired with an otherwise legal tile. The comment now quotes both figures from the one archive and records the earlier mixup, since this is the third time numbers from these two sweeps have been merged in a comment. Separately, the M%4 guard called is_gfx12_arch(std::string(GFX_ARCH)) -- building a std::string on every invocation to test a value fixed at compile time, in a file that already had a constexpr arch predicate a hundred lines above. Replaced with a constexpr kCompiledForGfx12, so on gfx9 the branch folds away entirely. The family-versus-exact distinction is preserved and now explicit: kCompiledForGfx12 covers gfx1200/gfx1201/gfx1250 and gates the M%4 guard, because that defect is an unpadded AQ view in a shared kernel header rather than a gfx1250 silicon property; kCompiledForGfx1250 stays exact and gates the warp tile assertions, because the WMMA fragment shapes genuinely differ within the family. Verified across seven arch strings including suffixed forms.
ozturkosu
changed the title
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
`normalize_gfx_arch` in dispatcher/codegen/codegen_common.py claimed to be
"the single source of truth for that rule; do not re-implement it", while
gemm_validation_utils.py open-codes `.split(":")[0]` three times. A comment
the adjacent code contradicts is worse than no comment: the next reader
either believes it and looks for an import that cannot exist, or stops
believing the comments in this area at all.
Two ways to fix it. Move the helper somewhere both trees can import, or
state the constraint honestly. Moving it is not available: the dependency
direction is dispatcher -> tile_engine (dispatcher/python/gemm_utils.py
imports gemm_validation_utils, not the other way round), so the shared home
would have to be inside tile_engine -- the tree that is being deprecated and
whose disappearance this bridge work is explicitly preparing for. Parking
new shared infrastructure there buys one copy today and a migration later.
So: keep two copies, say so, and make the duplication testable rather than
incidental.
- codegen_common.normalize_gfx_arch: scope the claim to the dispatcher tree,
and record why tile_engine cannot import it.
- gemm_validation_utils: add a module-private `_base_gfx_arch` and route the
pre-existing inline split in validate_lds_capacity through it, so the rule
has one definition per tree instead of one per call site.
- test_codegen_common.py: assert the two helpers agree on bare, single-suffix,
multi-suffix, empty and malformed inputs. Without this the pair is a comment
promising something nothing checks, which is the defect being fixed.
Tests: dispatcher/tests/test_codegen_common.py 82 -> 86 passed
(15 subtests), 0 failed.
Reviewers spotted `if gpu_target == "gfx950":` sitting a few lines below a
normalized gfx1250 test inside the same function. That is not a cosmetic
inconsistency -- it is a live bug, and it predates this branch.
_validate_fp8_mfma_warp_tile_k encodes ISA fragment shapes. gfx950 doubles
the fp8 K-block relative to gfx90a/gfx942 (MFMA_F32_16x16x256_F8 and
MFMA_F32_32x32x128_F8). A raw `== "gfx950"` means "gfx950:sramecc+:xnack-"
misses and falls through to the gfx942 branch, so the validator rejects
gfx950's correct warp_tile_k and accepts gfx942's wrong one. Suffixed targets
are not hypothetical here: projects/composablekernel/CMakeLists.txt sets
GPU_TARGETS to "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+" for
ASAN builds, and that string reaches --gpu_target unmodified.
Fixed by normalizing once at function entry into `base_gpu_target` and using
it for every arch test in the body, rather than adding a third normalized
comparison beside the raw one. Same treatment for the operator-scoped gfx1250
rule in validate_gemm_rowcol_tensor_quant. `gpu_target` itself is left alone
so diagnostics still echo what the caller passed.
Measured over {gfx90a, gfx942, gfx950, gfx1201, gfx1250} x {bare, suffixed} x
{fp8, bf8, fp16, bf16} x warp_tile_m {16,32,64} x warp_tile_k
{16,32,64,128,256}, 24 decisions change versus develop and all 24 are on
suffixed names:
- gfx950:sramecc+:xnack- : 12 changed. 16x16x128 and 32x32x64 now accepted
(correct for gfx950), 16x16x64 and 32x32x32 now rejected (those are the
gfx942 shapes).
- gfx1250:xnack- : 12 changed, matching bare gfx1250 -- the branch this
branch added.
- every bare target, and suffixed gfx90a/gfx942/gfx1201 : 0 changed.
New test file test_gemm_validation_arch_normalization.py pins it: suffixed
must equal bare for every arch in the matrix, gfx950's K-block must be the one
accepted under a suffix, gfx1250 must refuse the nonexistent 32x32 WMMA
fragment either way, and the error text must still show the caller's string.
Registered with CTest as dispatcher_test_gemm_validation_arch_normalization;
tile_engine has no CTest registration of its own, and the bridge/Old-TE
instance-set parity that depends on this module is a dispatcher property.
Tests: new file 4 passed / 308 subtests, 0 failed.
The previous revision normalized the WARP_SUPPORTED_COMBINATIONS lookup in
validate_warp_configuration for every architecture. Measured against develop
over warp_m/n/k in {1,2,4,8}: bare names give 0 changed decisions, suffixed
names give 45, all True -> False, 15 each for gfx90a/gfx942/gfx950. The
question the review asked is whether those 45 are a fix or a regression.
They are a regression, and the evidence is in the tree.
dispatcher/codegen/arch_specs.json describes itself as the single JSON source
of truth for arch data and generates both arch_specs_generated.py and the C++
runtime filter. It lists for gfx942:
[1,1,1] [1,2,1] [1,4,1] [2,1,1] [2,1,2] [2,2,1] [4,1,1]
and for gfx950 those plus [8,2,1] and [4,4,1]. The table in
gemm_validation_utils lists three maps for each. It is a stale hand-maintained
subset, not a description of the hardware. So 1x1x1 -- the example raised in
review -- is declared valid for gfx942 by the repo's own spec, and the
normalized lookup rejects it. Activating this table for suffixed names does
not stop invalid configurations from being generated; it stops valid ones.
That is a silent coverage loss on the ASAN configure in
projects/composablekernel/CMakeLists.txt, which sets GPU_TARGETS to
"gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+".
Second, independent reason. dispatcher/python/gemm_utils.py::_warp_config_supported
mirrors this function on purpose -- its comment says so, including "Unknown
arch => permissive (matches Old-TE's log-and-allow behavior)" -- so that the
bridge and Old-TE emit identical instance sets. It raw-gets too. Making one
side strict for suffixed names and not the other diverges the sets, which is
the precise failure that comment exists to prevent.
Option (b) from the review, then: restore develop's raw lookup and confine
normalization to gfx1250. gfx1250 is the one target where the normalization is
load-bearing and where nothing is lost -- arch_specs.json does not describe it
at all, so this table is its only listing, and there is no shipped gfx1250
behaviour for the other six operators to regress against.
Result versus develop, same grid, bare and suffixed:
gfx90a, gfx942, gfx950, gfx1201, gfx1250 (bare) : 0 changed
gfx90a:*, gfx942:*, gfx950:*, gfx1201:xnack- (suffix): 0 changed
gfx1250:xnack- : 55 changed
Everything outside the target of this branch is byte-identical to develop.
Widening the table to match arch_specs.json is the real fix for gfx9, but it
moves what seven instance builders generate on their primary targets and needs
its own PR with its own build evidence. Deliberately not done here.
Six tests pin this in test_gemm_validation_arch_normalization.py, including
one that asserts the table is a STRICT SUBSET of arch_specs.json -- written to
start failing the day the two agree, because at that point the argument above
no longer holds and the raw lookup should be revisited. Verified as a real
negative test: reinstating the unconditional normalization fails 4 of the 10.
Tests: test_gemm_validation_arch_normalization.py 4 -> 10 passed
(320 subtests), 0 failed.
_default_config() in unified_grouped_gemm_{rowcolquant,tensorquant}_codegen.py
took no architecture and returned the gfx9 MFMA tile unconditionally, while
bindings/ctypes/CMakeLists.txt derives GFX_ARCH from CMAKE_HIP_ARCHITECTURES[0].
Configure for gfx1250 and the two disagree, the compile-time guard fires, and
the target fails to build.
Three options were on the table. Leave the hard failure. Skip the target with
a status message, the way mx_gemm does twelve targets above in the same file.
Or wire the architecture in, using the selector rowcol_tensor_quant_default_tile(arch)
that already exists in codegen_common and was simply never called.
Wiring it in, and the static_assert stays.
Skipping is right for mx_gemm because mx_gemm is a single-architecture
operator: there is no correct gfx942 build of it to produce, so declining to
build one loses nothing. These two bridges are the opposite case -- gfx1250 is
a target they are meant to work on -- so a skip would leave the user with a
target that silently does not exist and no indication that it should have.
Meanwhile wiring the arch in, on its own, makes the build succeed on the
strength of an off-tree sweep, so the assertion stays as the floor: it is now
unreachable through the supported path and fires only if something upstream of
it goes wrong.
- _default_config(gfx_arch="") and generate_kernels(..., gfx_arch="").
Default "" reproduces the previous output byte for byte; gfx_arch is
ignored when an explicit --config/--config-json supplies tile_configs,
because such a config already carries its own tile.
- New --gfx-arch CLI flag on both codegens, honoured by --list-names too.
- CMake: strip feature suffixes from the derived GFX_ARCH (ASAN configures
put "gfx942:sramecc+:xnack-" in CMAKE_HIP_ARCHITECTURES, and GFX_ARCH is
compared against kSupportedArchs and the tile guards, which want the bare
target), and tell the user to pass --gfx-arch in the message printed when
no generated header is found.
Verified: no flag gives 128x128x64_2x2x1_32x32x16 as before; --gfx-arch
gfx1250 gives 16x64x256_1x4x1_16x16x128; --gfx-arch gfx1250:xnack- gives the
same as bare gfx1250.
Also makes the drift guard arch-aware, per the related finding.
TestDefaultConfigAlignment compared the runtime and codegen defaults at one
architecture, so it could not see a divergence that only exists on another --
which is precisely the divergence this change removes. Both cases are now
parametrized over gfx942/gfx950/gfx1250 plus a suffixed spelling of two of
them. A parametrized equality test can pass for the wrong reason (if
_default_config stopped honouring gfx_arch, both sides would fall back to the
gfx9 tile together and every case would still pass), so a companion test
asserts the gfx1250 and gfx942 default kernel names are disjoint.
Confirmed as a real negative test: pinning the codegen back to the gfx9 tile
fails 5 of the 11 cases in that class, on gfx1250 and gfx1250:xnack- only.
Tests: test_grouped_gemm_rowcolquant_utils.py + test_grouped_gemm_tensorquant_utils.py
108 -> 128 passed, 0 failed.
The block of static_asserts in the two ctypes libs had no test. It is the one
gfx1250 guard that outlives Tile-Engine deprecation -- the Python validator
disappears with tile_engine, and the runtime tile selector only covers the
Python build path -- so it is the guard that most needs to survive an unrelated
refactor, and nothing in dispatcher/tests or its CMakeLists so much as named it.
The guard is pure C++: a constexpr prefix test on GFX_ARCH plus static_asserts
over SelectedKernel's compile-time constants. It needs no device, no ROCm and
no ck_tile headers. So the test extracts the block from the shipped .cpp
between BEGIN/END markers, pairs it with a mock SelectedKernel, and asks the
host C++ compiler whether the result compiles.
Extraction, not transcription. A hand-copied guard in the test keeps passing
forever after someone deletes the original -- which is the failure mode being
guarded against, so a copy would test the wrong artifact. The markers are the
contract; test_guard_block_is_present_in_both_sources fails if they go.
Coverage, for both operators:
- four device-validated gfx1250 configurations must compile,
- twelve must not: 32x32x32, 16x16x32, the four 8-warp maps, [1,2,2], and
five single-rule probes,
- a suffixed "gfx1250:xnack-" is still guarded,
- and every rejected configuration must still compile for gfx942/gfx950,
because the guard is gfx1250-scoped and 32x32x32 is the correct gfx942 fp8
tile.
The single-rule probes matter more than they look. The asserts overlap:
32x32x32 trips the M, N and K rules at once, so a suite built only from
realistic tiles still passes after any one assert is deleted -- I confirmed
that before adding them. 32x16x64, 16x32x64, 16x16x256, 16x16x64 with 8x1x1,
and 16x16x64 with 1x1x2 are each caught by exactly one rule, which makes every
rule individually load-bearing. Verified: deleting the warp_tile_m assert from
the real source now fails the suite, and did not before.
test_guard_is_what_rejects_the_bad_configs is the other control. It strips the
static_asserts from the extracted block and requires all twelve to compile.
Without it, a syntax error in the extraction or a missing symbol would fail
every case and the suite would be measuring nothing while looking green.
Registered as dispatcher_test_gfx1250_compile_guard; skips itself if no host
C++ compiler is on PATH.
Tests: new file 6 passed / 132 subtests, 0 failed.
Correction to an earlier commit on this branch: e8ff28c quotes the
rowcolquant+tensorquant utils suites as "108 -> 128 passed". The measured
figures are 83 -> 101. The delta is right, the endpoints were not.
Review is right: kCompiledForGfx12 cannot differ from kCompiledForGfx1250.
gfx1200 and gfx1201 are the only other gfx12xx parts, neither appears in
kSupportedArchs, so dispatcher_initialize() returns -1 on them before
dispatcher_run_gemm -- and the M%4 guard the constant feeds is inside
dispatcher_run_gemm. The old comment presented a family/exact distinction as a
careful piece of reasoning without saying it currently distinguishes nothing.
Of the two offered resolutions -- delete it and use the exact predicate, or
keep it and say the generality is deliberate -- keeping it is the safer one,
and the safety is concrete rather than hypothetical. Two blocks above,
kSupportedArchs documents enabling a target as "a one-line addition here plus a
CMake arch entry". With the exact predicate, that one line applied to a gfx12xx
part would silently switch the M%4 guard off for the target it was just enabled
on, reintroducing the silent wrong answer the guard exists to prevent. Deleting
the constant trades a comment that overclaims for a trap.
So the constant stays and the comment now states outright that it buys nothing
today and why it is there anyway. A static_assert pins
kCompiledForGfx1250 => kCompiledForGfx12, so the claim cannot rot into a
falsehood the way the previous wording could.
clang-format-18 -style=file: clean.
Tests: test_gfx1250_compile_guard.py 6 passed / 132 subtests, 0 failed --
the new assert is inside the extracted block, so it is compiled by that test
too. Also fixed a false positive it exposed: the negative control matched the
bare word "static_assert", which now occurs in the block's prose, so it matches
"static_assert(" instead.
_SO_ABI was a literal 2 in two files, with nothing forcing a bump and nothing noticing when one was missed. The review asked for at minimum a test that the two agree, and for something self-maintaining if one was available. One is. The token exists because setup_multiple_*_dispatchers() reuses a cached .so on a filename match, so the filename has to describe the exported C ABI. A hand-maintained integer describes it only while everyone remembers, and the symptom of forgetting is an "undefined symbol" in a cache directory on a different machine, weeks later. A test that the two constants agree does not help with that: it catches the two files disagreeing, which is the lesser failure, and stays green when both are equally stale. So derive the token from the thing it is supposed to describe. _ctypes_abi_token() scans the export signatures in bindings/ctypes/grouped_gemm_*_ctypes_lib.cpp -- return type, name, parameter types -- and hashes the sorted set. Add, remove, rename or re-type an exported function and the cache key moves on its own. Signatures rather than the whole file: keying on file content would invalidate every cached artifact for a comment fix, which trains people to distrust the cache. Signatures are the ABI surface and nothing else. The Python path already compiles this exact .cpp (_CTYPES_LIB_SRC is passed to hipcc a few hundred lines down), so the source is guaranteed present wherever the token is needed. If it is unreadable, or the scan finds no exports, the helper logs and returns a fixed fallback -- degrading to the old hand-bumped behaviour rather than silently keying every build identically. Verified by hand before writing the tests: changing dispatcher_get_tile_n()'s return type from int to long moves the token ecce4e2b0e -> 6f68feb723; adding an export moves it to f0e3414b2a; reverting returns it to ecce4e2b0e. test_so_abi_token.py covers all of it -- token derived not hardcoded, stable across calls, moves for an added export / changed return type / changed parameter type, does NOT move for a comment edit, the current eight-export set is pinned so a silently non-matching regex cannot freeze the token, both modules use the same rule, and the unreadable-source fallback works. Mutations are applied to a temporary copy; the checked-in sources are never written to, because an interrupted test that edits the working tree leaves it corrupt. Registered as dispatcher_test_so_abi_token. Tests: new file 9 passed / 16 subtests, 0 failed; the two utils suites stay at 101 passed, 0 failed.
…oken Self-inflicted, and caught by the suite: cbdadea turned _SO_ABI from an integer into a derived string but left TestSoCacheAbiKey::test_abi_is_versioned asserting `isinstance(..., int) and >= 2` in both utils test files. Two failures. The commit message on cbdadea reports "101 passed, 0 failed" for those suites; that was the state before this hunk, not after. Correct figures are at the bottom of this message. The assertion is rewritten rather than deleted -- the property it was protecting (the cache filename carries something that distinguishes ABI revisions) still matters. It now checks the token is a non-empty alphanumeric string, safe to put in a filename, and not the read-failure fallback, with a pointer to test_so_abi_token.py where the derivation itself is tested. The companion test_pre_abi_artifact_is_not_reused needed no change: it already interpolates UTILS._SO_ABI rather than a literal. Tests: test_grouped_gemm_rowcolquant_utils.py + test_grouped_gemm_tensorquant_utils.py + test_so_abi_token.py 108 passed / 2 failed -> 110 passed / 0 failed.
The derivation read the export signatures out of the ctypes source with a regex at import time. The regex could not read a multi-token or namespaced return type: int dispatcher_get_tile_n() -> parsed const char* dispatcher_get_kernel_name() -> parsed unsigned long long dispatcher_x() -> [] std::size_t dispatcher_y() -> [] const unsigned char* dispatcher_z() -> [] so adding such an export dropped it, left the token unchanged, and the stale .so was reused anyway -- the exact failure the derivation existed to prevent, now silent. The companion test compared parsed names against a hardcoded eight-name set, so the dropped export left that set unchanged and the test still passed. Hardening the regex was the other option and is what one reviewer wanted. Rejected for where the parse lives, not for how good it is. Any parser of C++ declarations will have an edge it cannot read; on the runtime cache-key path that edge is silent by construction, because the only signal is "the token did not move", which is indistinguishable from "the ABI did not change". The same weakness in a test is a red build. Two further arguments the alternative does not answer: the derivation read a source file at import time and fell back to a fixed string when it could not, so an install shipping the Python without the C++ sources keys every build identically -- silently, and only in that environment; and the token is needed whether or not it is derived, so keeping a literal is strictly less machinery in the place that cannot report its own failure. So _SO_ABI is a literal again. The parse moves to tests/test_so_abi_token.py with a prefix pattern that reads any return type, comment stripping (the sources document their exports in declaration-shaped prose), and -- the part that matters -- a report of every dispatcher_* name in the export block that no signature accounted for. A parser weakness is now an assertion failure naming the symbol it could not read. The test also holds a frozen snapshot of each operator's signatures, checks the source against it, and derives the required literal from the snapshot, so changing the ABI without changing the literal fails, and updating the snapshot without updating the literal fails too. test_so_abi_token.py: 8 tests -> 11 tests, 21 subtests, all passing. The _SO_ABI_FALLBACK assertion in the two grouped-utils suites is updated; the fallback no longer exists, and the test now asserts the withdrawn names stay withdrawn.
…sing gfx9 --gfx-arch was accepted verbatim. It selects the built-in default tile and nothing else, so it never reaches --offload-arch and hipcc never sees it: a typo does not fail, it falls through to the gfx9 MFMA 32x32x16 tile. Generate with "--gfx-arch gfx1205" and compile the result for gfx1250 and you get a kernel that builds cleanly and returns garbage -- the exact failure this branch exists to close, entered through the front door. Validated at the CLI boundary rather than inside rowcol_tensor_quant_default_tile(). Making the tile selector itself raise was the tempting option, but it would delete a decision the panel already accepted: test_codegen_common.py asserts that gfx1200/gfx1201 fall back to the MFMA tile rather than inheriting gfx1250's, and that test documents why the gfx1250 match is exact and not a gfx12 family match. Raising there would replace a reviewed behaviour as a side effect of adding argument checking. The CLI is also the only entry point where an unknown target is genuinely silent -- every other path eventually hands the value to a compiler, which rejects it -- so that is where the check belongs and where its justification is honest. The accepted set is the same four as gemm_utils._SUPPORTED_ARCHES, pinned by a test, so the bridge does not advertise different support on its codegen and its runner. Not taken from arch_specs.json: that file omits gfx1250 entirely and lists gfx908/gfx1100/gfx1200/gfx1201, for which neither default tile is correct. Suffixed targets are normalized before the check, so the name a real device reports is usable. Behaviour change, deliberate and narrow: "--gfx-arch gfx1200" (and gfx908, gfx1100, gfx1201) previously produced the gfx9 tile and now exits 1. No automated caller passes the flag -- it appears only in a CMake diagnostic aimed at a human -- so this reaches manual invocations only, and only ones that were already getting a tile no one had validated for that target. tests/test_codegen_common.py: 87 -> 93 tests, all passing. CLI verified: gfx1250 and gfx1250:xnack- list kernels and exit 0; gfx1205 and gfx1200 print the supported set and exit 1.
The comment defending the unnormalized WARP_SUPPORTED_COMBINATIONS lookup gave
two reasons. The second one is not true.
It claimed that dispatcher/python/gemm_utils.py mirrors this function with a raw
.get, so normalizing here and not there would diverge the two instance sets. But
the mirror cannot observe the difference. _warp_config_supported() has exactly one
call site, expand_sweep(), and expand_sweep() sets `arch = _resolve_arch(arch)`
before it -- and _resolve_arch() raises ValueError for any name outside
("gfx90a", "gfx942", "gfx950", "gfx1250"). A suffixed target never reaches the
mirror's .get; it fails earlier. Whatever this side does with suffixed names, the
two sets cannot diverge on them.
Deleted rather than reworded. A reason that does not hold is worse than no reason:
it makes the decision look doubly supported and invites the next reader to
maintain a constraint that does not exist. The first reason carries the argument
by itself and was re-verified against arch_specs.json at this head -- gfx942 has
seven warp_configs, gfx950 nine, this table three -- so normalizing a suffixed
gfx942 would reject valid configurations, not invalid ones.
Comment-only. No test asserted the removed text.
_resolve_arch() checked its input against _SUPPORTED_ARCHES and only the callers
downstream of it normalized, so a target carrying feature flags was rejected
outright:
gfx1250 -> gfx1250
gfx1250:xnack- -> ValueError: Unsupported GPU architecture
gfx950:sramecc+:xnack- -> ValueError: Unsupported GPU architecture
This branch spent several commits pushing normalization to the right boundaries
and never reached this one, which is the entry point every other path goes
through -- expand_sweep() calls it before anything else.
Measured delta, _resolve_arch() over 15 inputs, before vs after:
gfx90a / gfx942 / gfx950 / gfx1250 unchanged (pass through)
gfx90a:xnack+ ValueError -> gfx90a CHANGED
gfx942:sramecc+:xnack- ValueError -> gfx942 CHANGED
gfx950:sramecc+:xnack- ValueError -> gfx950 CHANGED
gfx1250:xnack- ValueError -> gfx1250 CHANGED
gfx908 / gfx908:xnack+ / gfx1200 ValueError (unchanged)
gfx942x / gfx ValueError (unchanged)
"" / None unchanged (autodetect)
Four of fifteen change, all in the same direction: previously-raising becomes
succeeding with the base name. Nothing that used to succeed changes, and nothing
that used to raise for being unsupported starts passing -- gfx908:xnack+ still
raises, because normalizing feeds the gate, it does not bypass it. The error text
now shows both the input and the normalized form.
One correction to the finding as reported. "A real device reports the suffixed
form" is true of hipDeviceProp_t::gcnArchName and of rocminfo's ISA line, but
_get_arch()'s rocminfo fallback reads the agent "Name:" field, which is bare
(verified: this host prints "Name: gfx90a" alongside
"amdgcn-amd-amdhsa--gfx90a:sramecc+:xnack-"). So autodetection on such a host was
not broken. What was actually unusable is the explicit path -- and there is a
concrete source of suffixed names for it inside this repository:
projects/composablekernel/CMakeLists.txt sets GPU_TARGETS to
"gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+", so a caller copying a
target from the build configuration was rejected. amd-smi's
TARGET_GRAPHICS_VERSION is captured as \S+ and can also carry a suffix.
The suffix is dropped rather than carried into --offload-arch. That is the same
target string every already-working path used, since a suffixed name could not
get through at all; a caller needing a specific xnack setting must still pass it,
as before.
Scoped to gemm_utils.py. batched_gemm_utils.py has the identical _resolve_arch
and the identical defect -- it is a different operator, not touched by this PR,
and folding it in would widen the blast radius the panel is already asking to
narrow. Reported, not fixed here.
tests/test_gemm_utils.py: 34 -> 41 tests, 31 subtests, all passing. The
normalize_gfx_arch fallback (codegen_common is not importable at gemm_utils
import time) is pinned to the canonical helper on both branches.
Contributor
Author
Assessment: should
|
| commit | also touches |
|---|---|
6bbbcb4887 |
— (validator only) |
678fb564b4 |
— (validator only) |
5599a9f75a |
— (validator only, comment) |
fd4a1b3c22 |
both ctypes libs |
13a0f7dfcf |
both ctypes libs |
ce25d60ffb |
both ctypes libs |
6795101c13 |
rowcolquant ctypes lib |
cb87d4311e |
tests/CMakeLists.txt, new test file |
ba1dd96747 |
new test file |
d490a72058 |
codegen_common.py, test_codegen_common.py |
Seven of the ten are mixed. Splitting means separating hunks inside seven commits and reassembling two
histories — and the per-commit rationale is the thing three reviewers have been reading. That rationale
does not survive being cut in half.
The remainder would not stand alone in either direction. The Python validator and the C++
static_assert block in the two ctypes libraries encode the same rule, deliberately and redundantly:
the Python one disappears with Tile-Engine, so the C++ one is the survivor (32f66c6e9d, 99c7b1e183,
and its test 56974c97d9). Landing them in separate PRs puts develop in a window where one half of a
guard exists and the other does not, for a rule whose entire purpose is that a wrong tile compiles
cleanly and returns garbage. fd4a1b3c22 and 13a0f7dfcf exist precisely because the two wordings had
drifted apart once already.
Test coverage splits unevenly. test_gemm_validation_arch_normalization.py (10 tests, 320 subtests)
is entirely about this file and would move cleanly. But test_codegen_common.py::TestNormalizeGfxArch
pins tile_engine's _base_gfx_arch to codegen_common.normalize_gfx_arch across the two trees, and
normalize_gfx_arch itself was introduced by d490a72058 — a commit that lands in both halves. One of
the two PRs would carry a drift-guard whose other side is not in it.
The evidence does not split. The gfx1250 rules in this file rest on a 14,208-row sweep of these two
operators. Presented as a standalone validation-utils PR, the natural question is "what evidence
supports rejecting 32x32 on gfx1250 for AQuant?" and the honest answer lives in the operator PR.
But the concern behind the request is right
The reviewers are not wrong that something here has shipping-architecture blast radius. It is one narrow
thing, and it is worth naming precisely rather than splitting a file over: cb87d4311e normalizes the
target in _validate_fp8_mfma_warp_tile_k, which flips 8 decisions for a suffixed gfx950 across
four operators (16/64 and 32/32 accept->reject, 16/128 and 32/64 reject->accept, for both fp8 and bf8;
measured at f0af3afd34 over 128 probes). Bare gfx950 is unaffected. The change is correct — a suffixed
gfx950 was previously being handed gfx942's K-blocks — but this PR has no gfx950 device evidence.
Proposal
Gate rather than split: a gfx950 device run covering those 8 flipped decisions before this merges.
That addresses the actual risk, costs no history surgery, and keeps the guard's two halves in one
changeset. If reviewers still want the split after that, it should be done by re-authoring the two
histories deliberately, not by git-splitting these commits.
Happy to be overruled — this is a repository-structure call, not a technical one.
The _validate_arch docstring justified normalize-before-validate by claiming
amd-smi and rocminfo "report the target with its feature flags on a real
device" and that the suffixed form was "the only name such a device ever
gives". Both are false, measured on a gfx90a device: amd-smi
TARGET_GRAPHICS_VERSION, rocm_agent_enumerator and the rocminfo agent Name:
line all return a bare gfx90a. _get_arch cannot even emit a suffixed name --
rocminfo's only suffixed string is the ISA line
amdgcn-amd-amdhsa--gfx90a:sramecc+:xnack-, which fails the startswith("gfx")
check after the split.
The ordering change is still correct; the reason is different. Suffixed names
reach the function from callers: this repo's CMakeLists.txt ASAN branch sets
CK_GPU_TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+", and
hipDeviceProp_t::gcnArchName returns gfx90a:sramecc+:xnack- (measured). Either
can be copied into --gfx-arch and arrives here unmodified.
Also drop the false "gfx1250 -> ValueError" row from the before/after table in
both the docstring and the mirrored comment in tests/test_gemm_utils.py.
gfx1250 is in _SUPPORTED_ARCHES and always passed, as
test_bare_supported_arches_are_unchanged asserts directly below that comment.
Comment-only; no behaviour change.
… names normalize_gfx_arch's docstring named rocm_agent_enumerator as a source of feature-suffixed targets. It is not: on a gfx90a device it prints a bare gfx90a, as do amd-smi TARGET_GRAPHICS_VERSION and the rocminfo agent Name: line. Only the hipDeviceProp_t::gcnArchName half of that sentence was right -- it returns gfx90a:sramecc+:xnack- on the same device. Name the sources that actually produce suffixes: gcnArchName, and build configuration such as this repo's CMakeLists.txt ASAN branch setting CK_GPU_TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+". This defect is pre-existing -- introduced in 36c88c1, not by the recent commits on this branch. Comment-only; normalize_gfx_arch is unchanged.
Both entries pointed at one _COMMON_EXPORTS object. Changing one operator's ABI by editing that object moved BOTH tokens and still passed every test in this file: the untouched operator's .so cache was invalidated for nothing, and the diff did not show which operator actually changed. Each entry is now its own concrete tuple. Written as (*_EXPORTS_TEMPLATE,) and not tuple(_EXPORTS_TEMPLATE) because tuple() on a tuple returns the same object back, which leaves the entries aliased exactly as before -- verified by identity check. The token value does NOT move: token_for derives from tuple contents, which are unchanged, so both operators stay at re8d662d9 and no _SO_ABI literal needs updating. The two tokens are equal rather than distinct, which is correct -- the operators export the same C ABI, and the .so filename separates their caches through cfg.name. Appending an export to one entry now moves only that entry's token, confirmed by running token_for on both. tests/test_so_abi_token.py: 11 passed, 21 subtests, unchanged from before.
Contributor
Author
|
No need this anymore!! |
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.
Draft. Successor to #11043, opened so that review findings can be applied and re-reviewed against a
stable base rather than accumulating on a branch already under review.
Why a second PR
#11043 went through several review rounds. The code findings were addressed, but three reviewers
independently converged on objections that were structural rather than local, and applying them in place
would have made an already large diff harder to follow. This branch carries #11043's history up to its
tip (
6795101c13) and applies the rest as separate, reviewable commits.What this enables
grouped_gemm_rowcolquantandgrouped_gemm_tensorquanton gfx1250 / MI400. gfx1250 is wave32 withRDNA-style WMMA; gfx9 is wave64 with MFMA, and a gfx9 warp tile compiles cleanly on gfx1250 and returns
garbage rather than erroring. Every guard here exists because of that property, at three layers:
c2ccdc24c1,e8ff28c080) —rowcol_tensor_quant_default_tile(gfx_arch)selects theFlatMM 8-bit tile 16x64x256 / warp 1x4x1 / warp_tile 16x16x128 for gfx1250 instead of the gfx9 MFMA
32x32x16 tile. The match is exact, not a
gfx12family match: gfx1200/gfx1201 have a 16x16x16 8-bitfragment and would take a tile no one has validated for them.
6bbbcb4887,678fb564b4,cb87d4311e,ba1dd96747) — rejects gfx1250-invalid warpmaps and warp tiles before codegen.
32f66c6e9d,99c7b1e183,ce25d60ffb) —static_asserton the tile,and gfx1250 added to
kSupportedArchsso the.sois not rejected atdispatcher_initialize.Scope is wider than the title
Two things in this PR reach past the two named operators. Both are deliberate; neither is implied by the
title.
1. It changes AQuant, BQuant and ABQuant on gfx1250. The gfx1250 branch added to
_validate_fp8_mfma_warp_tile_katcb87d4311esits in a helper shared byvalidate_gemm_aquant,validate_gemm_bquant,validate_gemm_abquantandvalidate_gemm_rowcol_tensor_quant— five operatorscounting rowcolquant and tensorquant separately. The gfx1250 rule (only 16x16x{64,128} is a legal 8-bit
WMMA warp tile) is a property of the silicon, so it is correct for all of them, but only rowcolquant and
tensorquant have gfx1250 evidence behind it. The operator-specific rule — the >4-warp block restriction,
which is a property of this bridge's codegen and not of the hardware — was deliberately kept out of the
shared helper and out of
WARP_SUPPORTED_COMBINATIONSat678fb564b4, because plain GEMM's gfx1250default is an 8-warp 4x2x1 map that #11161 / #11175 verified correct on device.
2. It changes decisions on gfx950, a shipping architecture.
_validate_fp8_mfma_warp_tile_kcomparedgpu_target == "gfx950"raw;cb87d4311enormalizes the target first. Measured atf0af3afd34across128 probes (8 arch spellings x fp8/bf8 x warp_tile_m 16/32 x warp_tile_k 16/32/64/128), 16 decisions
differ from develop:
gfx90a,gfx942,gfx950,gfx1250(bare, gfx9)gfx950:sramecc+:xnack-gfx1250,gfx1250:xnack-A suffixed gfx950 previously fell through to the gfx90a/gfx942 branch and was told to use K-blocks 32/64
instead of 64/128 — the correct tile rejected, the wrong one accepted. The fix is right, but it flips
eight decisions on a shipping part for four operators, and this PR has no gfx950 device evidence.
Bare
gfx950is unaffected; the change bites only when the target carries feature flags, which thisrepository's own
CMakeLists.txtdoes (GPU_TARGETS=gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+).Findings applied on this branch
cb87d4311e,d490a72058). Measured delta atf0af3afd34: see the table above.ba1dd96747).WARP_SUPPORTED_COMBINATIONSis ahand-maintained subset, not a hardware description —
arch_specs.jsonlists 7 warp maps for gfx942and 9 for gfx950 where the table lists 3 — so normalizing it wholesale would reject valid
configurations, not invalid ones. Normalization is scoped to gfx1250, which
arch_specs.jsondoesnot describe at all. A second justification for this was found to be false on inspection and
deleted at
5599a9f75a: it claimed thegemm_utils.pymirror could observe a divergence, butexpand_sweep()calls_resolve_arch()first, so a suffixed target never reaches the mirror.e8ff28c080). The codegen default config took no architecture, soCMake's injected target produced a hard compile failure.
--gfx-archnow selects the tile.56974c97d9). It is the only guard that survivesTile-Engine deprecation and had none. 217 lines, registered in CTest.
normalize_gfx_archsingle-source-of-truth claim (d490a72058). Narrowed to "for thedispatcher tree", which is what it is:
tile_engine/cannot import it (dependency direction isdispatcher -> tile_engine, and tile_engine is on the deprecation path), so it keeps
_base_gfx_archand the two are pinned by
test_codegen_common.py::TestNormalizeGfxArch.d490a72058). Now multi-arch..socache key vs the ctypes export set (c17151fd58,949d3a8a14). A.sopredating thedispatcher_get_tile_n()/dispatcher_get_pad_n()exports was a name match in a persistent cachedir and died at attribute lookup.
_SO_ABIis folded into the filename. It was briefly derived atimport time by regex-scanning the ctypes source; that is withdrawn at
949d3a8a14, because theregex could not read a multi-token or namespaced return type, dropped those exports, and left the
token unchanged — reproducing the exact stale-cache failure, silently. The parse now lives in
tests/test_so_abi_token.py, reports anydispatcher_*name it could not account for, and derivesthe required literal from a frozen snapshot.
--gfx-archwas unvalidated (47fdbd53cc). It only selects a tile, so it never reaches--offload-archand a typo was never rejected — it fell through to the gfx9 MFMA tile. Now checkedagainst the same four architectures as
gemm_utils._SUPPORTED_ARCHES, pinned by a test.Behaviour change:
--gfx-arch gfx908|gfx1100|gfx1200|gfx1201previously produced the gfx9 tile andnow exits 1. No automated caller passes the flag.
_resolve_archraised on the name the hardware reports (f0af3afd34). It validated beforenormalizing, so
gfx1250:xnack-andgfx950:sramecc+:xnack-raisedValueError. Measured deltaover 15 inputs: 4 change, all previously-raising -> now resolving to the base name; nothing that
succeeded changes and no unsupported target starts passing.
batched_gemm_utils.pycarries theidentical defect and is not fixed here — different operator, and widening this diff is the
opposite of what review asked for.
Open review question
Two reviewers have asked for
tile_engine/ops/gemm/gemm_validation_utils.pyto be split into its own PR,on the grounds that its blast radius is five operators and the gfx950 change lands there. Not done
unilaterally; see the assessment comment on this PR. The short version: 10 of the 21 commits touch that
file and 7 of those 10 also touch the ctypes libraries,
codegen_common.pyor the tests, so the split isa hunk-level rewrite of history rather than a commit move, and the remaining PR would not stand alone —
its C++
static_assertand the Python validator encode the same rule and are currently pinned to eachother.
What is NOT validated
0 FAIL / 1,472 clean-reject / 0 launch-abort, worst global relative error 4.36e-4 against a 0.05
tolerance, against an independent CPU fp32 reference — was measured before this branch existed. I could
not locate the commit it was taken at, in feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043's body, in feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043's comments, or in this branch's
history, so its provenance is unpinned and it should not be treated as evidence for this head.
compare against on this architecture.
test_gfx1250_compile_guard.pydoes compile it -- it extracts the block verbatim from the shipped.cppand hands it to the host C++ compiler -- but against a mockSelectedKernel, not against CK'sheaders under hipcc.
launch; 2,168 of those when paired with an otherwise legal tile; all 3,226 wrong-result rows using
32x32x32) was recounted and reconciled to a single archive at
6795101c13, afterfd4a1b3c22foundthat earlier comments had merged figures from two different sweeps. Those figures are pinned to those
commits; they were not re-measured on this head.
Supersedes: #11043 (kept open until this one is reviewed).