Conversation
ozturkosu
changed the base branch from
users/muozturk/ck/blockscale_all_bridges
to
develop
August 19, 2026 22:15
✅ 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. |
Re-applies the gfx1250 (MI400) enablement of PR #11042 on top of current develop, now that PR #10439 has merged the block-scale quant bridges themselves (fb1cd33). Only the arch-enablement delta is kept; the stale copy of the bridges carried by the old branch is dropped. - aquant / abquant / tensor_quant: add gfx1250 to _SUPPORTED_ARCHS so _validate_arch()/_detect_gpu_arch() accept an MI400 host instead of raising. (bquant / rowcolquant have no arch allow-list.) - All five arch-derived warp_tile_k selectors (_warp_tile_k_for / fp8_warp_tile_k_for_arch): treat gfx12* like gfx950 -> 128 for the 8-bit-float quant path. warp_tile_k=32 is the gfx9 MFMA value and silently outputs zeros on gfx1250. Already on develop, so intentionally NOT re-applied here: - quant_bridge_common.hpp validate_supported_arch() already accepts gfx1250. - test_bquant_gpu_correctness.py already auto-detects the arch and SKIPs the MX (H3) variants off gfx950 (the b0c8ca5 fix landed with #10439).
ozturkosu
force-pushed
the
users/muozturk/ck/blockscale-quant-gfx1250-enable
branch
from
September 1, 2026 22:40
b0c8ca5 to
d5e1a78
Compare
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
It enables a new GPU-architecture path where correctness hinges on hardware-specific warp-tile behavior and should be confirmed with on-device validation.
Pull request overview
Enables the block-scale quant GEMM Tile-Engine → dispatcher bridges to run on gfx1250 (MI400) by widening runtime arch gates and making fp8/bf8 warp_tile_k defaults treat gfx12-class hardware like gfx950, avoiding the known WMMA/MFMA silent all-zero behavior.
Changes:
- Add
gfx1250to the supported-arch gates for the relevant bridges. - Adjust fp8/bf8
warp_tile_kselection logic to return the “128” path on gfx12-class architectures (including gfx1250). - Update the shared C++ runtime arch gate to accept gfx1250.
File summaries
| File | Description |
|---|---|
| projects/composablekernel/dispatcher/bindings/ctypes/quant_bridge_common.hpp | Extends runtime supported-arch validation to include gfx1250. |
| projects/composablekernel/dispatcher/python/gemm_tensor_quant_utils.py | Adds gfx1250 to supported archs and treats gfx12 as the 128 warp_tile_k case for fp8/bf8. |
| projects/composablekernel/dispatcher/python/gemm_rowcolquant_utils.py | Treats gfx12 as the 128 warp_tile_k case for fp8/bf8 defaults. |
| projects/composablekernel/dispatcher/python/gemm_aquant_utils.py | Adds gfx1250 to supported archs and treats gfx12 as the 128 warp_tile_k case. |
| projects/composablekernel/dispatcher/python/gemm_abquant_utils.py | Adds gfx1250 to supported archs and treats gfx12 as the 128 warp_tile_k case. |
| projects/composablekernel/dispatcher/python/gemm_bquant_utils.py | Treats gfx12 as the 128 warp_tile_k case. |
Review details
- Files reviewed: 5/5 changed files
- Comments generated: 5
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
ThruptiRajLakshmanaGowda
requested changes
Sep 3, 2026
ThruptiRajLakshmanaGowda
left a comment
ThruptiRajLakshmanaGowda
left a comment
Contributor
There was a problem hiding this comment.
Hey Folks, is this an urgent change? Can we wait till freeze to merge this PR?
…add tests The gfx1250 enablement widened the 8-bit warp_tile_k selectors in the five non-grouped block-scale quant bridges with a family-wide "gfx12" substring test. That is wrong for gfx1200/gfx1201: only gfx1250 has a 16x16x128 8-bit WMMA fragment, the others expose 16x16x16. A K=128 warp tile still compiles on those parts and then silently returns garbage -- the same failure mode as K=128 on gfx942 -- so this is a correctness gate, not a naming detail. Narrow the five warp_tile_k selectors (aquant, abquant, bquant, rowcolquant, tensor_quant) to an exact gfx1250 match via a small per-module _is_gfx1250() helper that strips the feature suffix reported by real agents (gfx1250:xnack-). The OCP fp8 encoding predicates and compile defines are left family-wide, which is correct: every gfx12xx part uses OCP e4m3/e5m2, so narrowing those would break fp8 on gfx1200/gfx1201. Both categories are now commented at every site so neither can be "tidied" into the other. Also refresh the five selector docstrings, which still described the 128 case as gfx950-only, and add a CPU-only test suite pinning both categories: gfx1250 (bare and suffixed) positive, gfx1200/gfx1201 negative for warp_tile_k but positive for OCP fp8, and gfx90a/gfx942/gfx950 outputs unchanged. Note: _is_gfx1250 is intentionally private per module. PR #11043 introduces a shared normalize_gfx_arch() in codegen_common.py but has not merged, so it is not importable from develop; these helpers should collapse onto it once it lands.
ozturkosu
marked this pull request as draft
September 17, 2026 00:34
ozturkosu
changed the title
ozturkosu
requested review from
ThruptiRajLakshmanaGowda
and
a balanced review from Copilot
September 18, 2026 06:15
ozturkosu
marked this pull request as ready for review
September 18, 2026 06:15
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Several enabled gfx1250 variants remain incorrect or non-buildable, and architecture validation is inconsistent.
Get a fresh assessment by requesting another Copilot review.
Review details
- Files reviewed: 7/7 changed files
- Comments generated: 6
- Review effort level: Balanced
ozturkosu
marked this pull request as draft
September 18, 2026 09:12
All gfx1250 behaviour below is GPU-verified on MI400 (gfx1250, ROCm 10.1) as
part of this change; the numbers quoted are from those runs.
- abquant fp4 on gfx1250 used warp_tile_k=32, which is a silent-wrong-answer
bug. abquant dispatches on AComputeDataType and auto_compute_type collapses
packed A==B==pk_fp4_t to fp8_t, so at K=32 the fp4 kernel selects the
unguarded MFMA Dispatcher<fp8_t,fp8_t,float,16,16,32>, whose intrinsics are
#if __gfx94__/__gfx95__ only and otherwise return CVecType{0.f}. Measured on
MI400: it builds, launches, and yields a dead accumulator -- ~50% exact zeros
and ~50% uncorrelated noise (corr vs reference -0.004). It is NOT uniformly
zero, so the all-zero guards in the GPU harnesses do not catch it. K=64 and
K=128 both verify at 4.74e-4 max relative error (fp8 control: 4.70e-4);
gfx1250 now uses 128 for every variant, so the arch has one uniform tile.
- fp8i4 / bf8i4 are now rejected on gfx1250. Across {fp8i4,bf8i4} x warp_tile_k
{16,32,64,128} x N {128,256,512}, all 24 combinations built and launched and
all 24 returned NaN, while C4/fp8 and C4/bf8 passed in the same harness with
the same host codec. The rejection lives in the config dataclass __post_init__
so it covers the public constructors, the internal builders and sweep
expansion alike. Note this contradicts the prose in the grouped sibling, which
says i4 "does NOT compile on gfx1250": it compiles cleanly and is wrong at
runtime, which is the more dangerous direction.
- grouped bquant regressed to warp_tile_k=32 on gfx1250. The merge from develop
generalized the previously hardcoded warp_tile_k=128 into _fp8_warp_tile_k()/
_preshuffleb_warp_tile_k() but gated them on gfx950 only, so
default_fp8_config_gfx1250() -- which delegates straight to default_fp8_config
-- began emitting a 16x16x32 kernel. Restored via an exact, suffix-tolerant
arch set. This is what TestGfx1250Configs was already asserting and had
started failing on.
- _validate_arch now exact-matches after stripping the feature suffix, so
gfx12500 is rejected instead of passing startswith and silently falling
through to the legacy tile. gfx1250:xnack- and gfx942:sramecc+:xnack- still
validate, and the full target id is returned unchanged for --offload-arch.
- codegen_common.fp8_warp_tile_k_for_arch aligned on gfx1250: --list-names
--gfx-arch gfx1250 reported 16x16x32 while the runtime built 16x16x128. That
helper's own docstring requires the rule to exist exactly once. It now routes
through the shared normalize_gfx_arch() that landed with #11043.
- OCP fp8 compile defines extracted into a per-module _ocp_fp8_arch_defines()
so the family-wide-gfx12 invariant is unit-tested across all five bridges;
this also fixes abquant, the one module that gated them on gfx950 only. The
device pass was already OCP via config.hpp's __gfx12__ fallback, so this
removes a host/device constexpr skew and a latent trap, not wrong numerics.
- _is_gfx1250 is evaluated first at every call site: `in` raised TypeError on
None before the helper ran, making its documented None tolerance unreachable
from the bridges. Selectors now raise on a missing arch rather than
defaulting to the gfx942 tile.
- Docstrings no longer claim the selectors mirror get_k_warp_tile on gfx1250.
That helper returns 64 there (WMMA branch at M_Warp_Tile==16, reachable
because CMakeLists sets CK_TILE_USE_WMMA=1 for gfx12); 128 is a deliberate,
now GPU-verified divergence and is documented as such.
Arch-gate suite 168 -> 317 tests. Full dispatcher suite, same node and PATH:
1418 -> 1488 passed with the two TestGfx1250Configs failures fixed and none
introduced (28 -> 26 on the GPU runner; the remainder are pre-existing and are
not in files this PR touches).
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Build-target overrides and missing standalone architecture defines can bypass the new gfx1250 correctness safeguards.
Get a fresh assessment by requesting another Copilot review.
Review effort: Balanced
Findings: 1
Open (3)
Resolved since last review (11)
fp8i4/bf8i4BQuant constructors, so passinggfx1250now…fp8i4/bf8i4AQuant constructors, so passinggfx1250…gfx1250to the bridge allowlist exposes FP8/BF8 ABQuant builds, but…arch.startswith(...), so all…gfx1250, this branch still returns32for thefp4variant. Under__gfx125__,…
TestShuffledBCacheGuardParity searched gemm_ctypes_lib.cpp for a literal
`g_shuffled_b_cache = ShuffledBCache{}` assignment and walked back to the
nearest `#if ... GEMM_KEY_PRESHUFFLE` guard. Neither exists any more:
- the cache is released with `.clear()` rather than reassigned, so the
`next()` over the source raised StopIteration and the test errored out
instead of asserting anything;
- GEMM_KEY_PRESHUFFLE was retired entirely. Preshuffle is now a capability
carried in the kernel's own metadata (SelectedKernel::Preshuffle, emitted
by unified_gemm_codegen.py) and branched on with `if constexpr`, giving the
same dead-code elimination the macro used to. The one remaining guard,
GEMM_KEY_DTYPE_A, only asks "is this a modern codegen header?".
So the test was pinning a macro spelling that the code deliberately moved away
from, while no longer checking anything.
Rewritten to assert the invariant that does still matter and is spelling
agnostic: every reference to g_shuffled_b_cache must be nested inside the
preprocessor guard its definition sits under, so no use-site can reference an
undeclared cache. A small #if/#elif/#else/#endif scanner builds the active
guard stack per line and compares each use against the definition's.
Verified both ways: it passes on the current source, and moving the `.clear()`
call one line past its `#endif` makes it fail and name the offending line,
guard and missing condition -- so it is not vacuous.
Pre-existing failure on develop; unrelated to any gfx1250 work.
Two dispatcher bridges loaded libamdhip64 with their own partial name list, and
both fail on a stock ROCm 7 node:
- grouped_conv_utils.py hardcoded ctypes.CDLL("libamdhip64.so"). The bare name
is the *development* symlink; it ships in $ROCM_PATH/lib but is generally not
in the ldconfig cache, so the load fails unless the caller happens to have
LD_LIBRARY_PATH set. The failure is then swallowed into _init_error,
is_available() returns False, every runner is dropped, and the grouped-conv
examples report "JIT build failed" -- which is doubly misleading, since
codegen and hipcc both succeeded and nothing about the build was wrong.
- fmha_utils.py tried ["libamdhip64.so", "libamdhip64.so.6"], which covers
ROCm 6 but not ROCm 7, where only .so.7 is registered.
Replaced both with a shared dispatcher_common.load_hip_runtime(), so the soname
list exists once instead of drifting per bridge. It tries the unversioned name
first (correct when the dev package is on the path), then the known versioned
sonames, then the same names under $ROCM_PATH/lib (default /opt/rocm) so the
load still succeeds when the library is present but unregistered. On failure it
raises naming every candidate tried, rather than surfacing later as a bare
"no GPU available".
Verified on an MI400 node with ROCm 7.15, with LD_LIBRARY_PATH unset:
ctypes.CDLL("libamdhip64.so") raises "cannot open shared object file", while
load_hip_runtime() resolves /opt/rocm/lib/libamdhip64.so.
Pre-existing on develop; unrelated to any gfx1250 enablement work.
…ing silently GpuGroupedConvRunner._ensure_initialized() returned early when no dispatcher library could be loaded, leaving _init_error as None. is_available() then reported False with no recorded reason, so a caller could not distinguish "no dispatcher .so was found" from "the GPU context failed to initialise" -- and the grouped-conv examples surfaced both as a bare "JIT build failed", even in runs where codegen and hipcc had demonstrably succeeded and the .so was on disk. Record the reason on that path, naming whether an explicit lib_path or auto-detection was used, and noting that a successful build does not imply the library landed where find() looks. No behaviour change: the early return and is_available() semantics are unchanged. This only populates the diagnostic that get_init_error() already promises to return. Pre-existing on develop; unrelated to any gfx1250 enablement work.
dispatcher_gemm_lib could not be built for gfx1250 at all.
The single fallback kernel generated for the Python library hardcoded warp tile
32x32x16, both in the tile-config JSON and in the expected output filename. That
is an MFMA fragment: valid on gfx9, and absent on every wave32/WMMA part --
gfx1250 needs 16x16x32, gfx1200/gfx1201 need 16x16x16. On those targets the
codegen's arch filter dropped the only requested config, reporting
INFO: standard: 0 configurations
INFO: Kernels: 0 ... Failed: 0
and exiting 0, so add_custom_target reported success. The breakage then surfaced
two steps later as dispatcher_gemm_lib failing with a bare "file not found" on
the -include of a header nothing had generated -- a confusing symptom a long way
from its cause.
Two changes:
- Select the fp16 warp tile at configure time from
arch_specs_generated.get_warp_tile_combos(), the same table the codegen filters
against, and build both the tile JSON and the expected filename from those
values. They are now derived from one source and cannot drift apart again. The
chosen tile is printed as a STATUS line, and an architecture the table does not
know is a configure-time FATAL_ERROR rather than a late compile failure.
- Assert the kernel actually appeared before the library tries to include it, so
a generated-nothing run fails at the generating step with an explanatory
message instead of at the consuming one.
Verified on MI400: configure now reports
-- GEMM Python fallback kernel: fp16 warp tile 16x16x32 for gfx1250
and emits gemm_fp16_..._2x2x1_16x16x32.hpp, where 32x32x16 previously produced
nothing.
NOTE: this unblocks the build step but does not make dispatcher_gemm_lib
buildable on gfx1250 end to end. With a valid tile the compile now reaches, and
fails in, the wave-wise WMMA path:
include/ck_tile/core/arch/mma/mma_wavewise.hpp:266: error: 'exec' following the
'template' keyword does not refer to a template
instantiated from WaveWiseMmaPipeline<fp16, fp16, float, 16, 16, 32, ...>. That
is a separate, pre-existing gfx1250 defect that the hardcoded tile was hiding;
it needs its own fix.
Pre-existing on develop; unrelated to any gfx1250 enablement work.
…tten
Four independent problems, none of them a real product defect:
1. test_batched_gemm_gpu_correctness / test_contraction_multi_abd_gpu_correctness
errored at collection with "fixture 'gfx_arch' not found". Both modules are
script-style -- a parameterized function plus a main() guard, run directly by
ctest (tests/CMakeLists.txt says so explicitly) -- but their entry points were
named test_*, so pytest collected them and tried to inject a fixture that does
not exist (conftest provides gpu_arch). They also return a (status, detail)
tuple, which pytest flags as well. Renamed the entry points to check_* so
pytest collects nothing from them; main() is unchanged and remains the
supported entry point. Verified both ways: pytest now reports "no tests ran",
and running each script still exits 77 (SKIP) without a GPU.
2. test_multi_abd_gpu_correctness hardcoded warp tile 32x32x16 while passing the
*detected* arch to codegen. That is an MFMA fragment, so on a wave32/WMMA part
the arch filter dropped the config, nothing was generated, and
setup_multiple_gemm_dispatchers returned None -- reported as a build failure
rather than "unsupported here". The tile is now taken from
arch_specs_generated.get_warp_tile_combos(), preferring 32x32x16 while it is
still legal so gfx90a/gfx942/gfx950 are byte-identical to before; only arches
that could never run pick something new (gfx1250 -> 16x16x32,
gfx1100/1200/1201 -> 16x16x16). An arch with no usable fp16 tile now skips
with a reason instead of failing.
3. test_multi_abd_bridge::test_configs_are_valid_sweeps globbed *.json in the
config directory and asserted every file contained "tile_config". The
directory also ships problem-shape files (example_problems.json is a
{"problems": [...]} list), so the test failed on a file it never meant to
check. Now filters to files that actually carry a tile_config, and still
asserts at least one sweep config is shipped.
Net effect on a CPU-only runner: two collection ERRORs and two FAILures removed,
no behaviour change on any architecture that previously passed.
Pre-existing on develop; unrelated to any gfx1250 enablement work.
github-actions
Bot
added
ck: attention-moe
Contributor
Author
Note on this PR's historyEarlier timeline comments on this PR (on-device validation runs, the Summarising the part reviewers need, so it is not only in the archive: The last 5 commits fix pre-existing dispatcher failures, not gfx1250 enablement. They reproduce identically without this PR and touch none of the files the gfx1250 work touches. They are in separate commits and can be split into their own PR against
|
1. Configs could be compiled for an architecture they were not built for. Every arch safeguard in these bridges -- the warp_tile_k selectors, the fp4 rule, the i4 rejection -- runs when the CONFIG is constructed, keyed on that config's own gfx_arch. The compile entry points take their own gfx_arch, so a config built for one arch and handed to a build for another slipped past all of them and had its literal tile emitted verbatim. Reproduced: default_fp4_config(gfx_arch="gfx950") records warp_tile_k=32, and setup_multiple_abquant_dispatchers([cfg], gfx_arch="gfx1250") emitted that 16x16x32 tile for gfx1250 -- the GPU-confirmed dead-accumulator case the fp4 rule exists to prevent. The same hole bypassed the AQuant/BQuant i4 rejection. Added dispatcher_common.validate_configs_match_arch(), called from all five setup_multiple_*_dispatchers entry points. Only a genuine mismatch raises; configs with no recorded arch are left alone. 2. The five standalone JIT compile paths emitted only the OCP encoding defines. They omitted the per-arch feature set that dispatcher_common.arch_feature_defines() supplies and that the top-level CMakeLists provides for a normal build -- notably CK_TILE_USE_WMMA, which must be passed even when 0: leaving it undefined only works because the preprocessor reads an undefined identifier as 0, which is the right answer on gfx942/gfx950 and the wrong one on every WMMA part. A gfx1250 kernel could therefore compile down the non-WMMA feature path. All five now use the shared helper, as the five grouped bridges already did; the local _ocp_fp8_arch_defines() introduced earlier in this PR was a partial reinvention of it and is removed. The shared helper is a strict superset: it also covers the gfx950 NATIVE_MX/GFX950_SUPPORT flags that were inline here. It deliberately omits CK_USE_GFX950, which ck_tile never reads. Note this does not change any emitted warp_tile_k: codegen writes a literal and the kernels never call get_k_warp_tile(). 3. Reverted the grouped-bquant preshuffle-B warp tile on gfx1250. _preshuffleb_warp_tile_k feeds six public constructors (preshuffle-B, preshuffle-quant and the combined form, fp8 and bf8), none of which has gfx1250 test coverage or on-device validation. Only _fp8_warp_tile_k was needed for the TestGfx1250Configs failures this PR fixes, so gfx1250 keeps the 64 it already receives -- instantiable there (gfx1250 has a 16x16x64 fp8 WMMA fragment), and unchanged, so no unvalidated numeric behaviour is introduced. Extending it wants GPU correctness coverage for those six paths first. Arch-gate suite 277 -> 293 tests, now covering the arch-mismatch rejection and the gfx1250 WMMA feature defines. Full dispatcher suite unchanged at 15 failed / 1465 passed; no regressions.
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Architecture safeguards remain bypassable or inconsistent in several public build paths.
Get a fresh assessment by requesting another Copilot review.
Review effort: Balanced
Findings: 3
Open (6)
Apply the shared architecture features and gfx1250 framework gate to the standalone GEMM fallback target. Build the static dispatcher prerequisite for GPU integration tests and create missing JIT output directories. Bootstrap requested GEMM libraries without requiring a prebuilt default. Include architecture and all configuration fields in cache identity, isolate generated headers, reject rebuild failures, and compile identical auto-corrected outputs once. Reject unsupported permute-N configurations before device discovery. Restore the newly enabled gfx1250 GEMM parity coverage with native WMMA tiles, the native int8 architecture entry, vector-aligned padding cases, and corrected B-preshuffle access ordering. Add a standalone CMake build-and-compute regression test and use isolated architecture-aware caching tests. Validation on MI400/gfx1250 with ROCm 10.0: - All 15 originally reported failures pass without skips. - Full dispatcher suite: 1638 passed, 14 skipped, 1192 subtests passed. - Standalone dispatcher_gemm_lib builds and verifies two fp16 GEMMs. - Independent fp16/bf16/fp8/bf8 preshuffle matrix: 12/12 passed. - CPU suite: 1467 passed, 185 skipped, 1030 subtests passed.
ozturkosu
marked this pull request as ready for review
September 19, 2026 04:53
ozturkosu
requested review from
yraparti
and
a balanced review from Copilot
September 19, 2026 07:21
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Mutable or unsupported configurations can still bypass gfx1250 safety checks, and some new regressions are not reliably exercised.
Get a fresh assessment by requesting another Copilot review.
Review effort: Balanced
Findings: 5
Open (9)
Resolved since last review (6)
Comment on lines
+2608
to
+2612
| elif lib is None: | ||
| lib = DispatcherLib.auto() | ||
| if lib is None or lib.get_kernel_name() != kernel_header.stem: | ||
| result.error = "No matching dispatcher library; enable auto_rebuild to build it" | ||
| return result |
Comment on lines
+721
to
+722
| configs = resolve_default_configs(configs, arch) | ||
| validate_configs_match_arch(configs, arch, "ABQuant") |
Comment on lines
1012
to
+1014
| arch = gfx_arch or _detect_gpu_arch() | ||
| configs = resolve_default_configs(configs, arch) | ||
| validate_configs_match_arch(configs, arch, "BQuant") |
Comment on lines
623
to
+625
| arch = gfx_arch or _detect_gpu_arch() | ||
| configs = resolve_default_configs(configs, arch) | ||
| validate_configs_match_arch(configs, arch, "RowColQuant") |
Comment on lines
+535
to
+536
| configs = resolve_default_configs(configs, arch) | ||
| validate_configs_match_arch(configs, arch, "TensorQuant") |
Comment on lines
+100
to
+117
| candidates.extend(_HIP_SONAMES[1:]) | ||
|
|
||
| rocm = Path(os.environ.get("ROCM_PATH", "/opt/rocm")).expanduser() | ||
| for libdir in (rocm / "lib", rocm / "lib64"): | ||
| candidates.append(str(libdir / _HIP_SONAMES[0])) | ||
| # Numeric ordering tries .so.10 before .so.9 and accepts full filenames | ||
| # such as .so.10.0.26306 when even the major-version symlink is absent. | ||
| versioned = [] | ||
| try: | ||
| for path in libdir.glob("libamdhip64.so.*"): | ||
| match = re.fullmatch(r"libamdhip64\.so\.(\d+(?:\.\d+)*)", path.name) | ||
| if match and path.is_file(): | ||
| versioned.append((tuple(map(int, match[1].split("."))), str(path))) | ||
| except OSError: | ||
| pass | ||
| candidates.extend(path for _, path in sorted(versioned, reverse=True)) | ||
| candidates.extend(str(libdir / soname) for soname in _HIP_SONAMES[1:]) | ||
| return list(dict.fromkeys(candidates)) |
Comment on lines
+153
to
+167
| archive = build / "libck_tile_dispatcher.a" | ||
| if not archive.exists(): | ||
| hipcc = shutil.which("hipcc") or "/opt/rocm/bin/hipcc" | ||
| commands = [] | ||
| if not (build / "CMakeCache.txt").exists(): | ||
| commands.append([ | ||
| "cmake", "-S", str(root), "-B", str(build), | ||
| f"-DCMAKE_CXX_COMPILER={hipcc}", "-DCMAKE_BUILD_TYPE=Release", | ||
| ]) | ||
| commands.append([ | ||
| "cmake", "--build", str(build), "--target", "ck_tile_dispatcher", "-j4", | ||
| ]) | ||
| for command in commands: | ||
| result = subprocess.run(command, capture_output=True, text=True, timeout=300) | ||
| assert result.returncode == 0, result.stdout + result.stderr |
Comment on lines
105
to
+106
| def _run_dtype(self, dtype: str): | ||
| from arch_specs_generated import get_warp_tile_combos |
Comment on lines
+34
to
+36
| from arch_specs_generated import get_warp_tile_combos | ||
|
|
||
| tiles = [tuple(t) for t in get_warp_tile_combos(gpu_arch.split(":", 1)[0], "fp16_fp16_fp32")] |
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.
JIRA ID: AICK-2105
JIRA ID: AICK-2138
Summary
This PR enables the five non-grouped block-scale quant GEMM bridges (
aquant,abquant,bquant,rowcolquant, andtensor_quant) on gfx1250 (MI400). Legacy MFMA tiles can compile and launch on gfx1250 but produce incorrect output.The bridges now select gfx1250 WMMA defaults and reject known unsafe requests. The PR also fixes standalone dispatcher builds, library caching, and B-preshuffle correctness needed for MI400 validation.
gfx1250 support
The table describes the non-grouped bridge configurations enabled by this PR. Availability of a dtype does not imply that every layout, tile, or preshuffle combination has been GPU-validated.
aquantrcrdefault. Codegen also acceptsrrr,crr, andccr;ccris unavailable with A-scale preshuffling.abquantFP8/BF8rcrdefaults.abquantFP4compv3,rcr. FP4 with B preshuffling is rejected at runtime.bquantrcrdefaults.rowcolquantcompv3,rcronly.tensor_quantcompv3,rcronly.All rows use FP32 scales and accumulation, with
16x16x128as the gfx1250 default warp tile. Here FP8 means OCP E4M3 and BF8 means OCP E5M2. Layout tags describe A/B/C storage:ris row-major andcis column-major, sorcrmeans row-major A, column-major B, row-major C.fp8i4/bf8i4fp8i4/bf8i4mx_bf16bf16,mx_bf16bf8,mx_bf16fp4The separate regular-GEMM INT8
16x16x64architecture entry is not an INT8 quant-bridge variant.Changes
warp_tile_k=128for gfx1250 quant defaults, including ABQuant FP4, and align the shared codegen helper and grouped BQuant's non-preshuffle FP8/BF8 defaults. Match gfx1250 exactly after stripping feature suffixes; preserve full targets in compiler flags and existing gfx9 tile rules.fp8i4andbf8i4on gfx1250, where tested configurations compile and launch but return NaNs. Validate AQuant's explicit or modified16x16FP8/BF8 tiles, and reject config/build architecture mismatches across all five bridges.16x16x64tile to architecture constraints, restore GPU test prerequisites and pytest entry points, and extend regression coverage.Design Note
Tile selection and FP8 encoding use different architecture rules: the
16x16x128quant default is specific to gfx1250, while OCP FP8 encoding applies across gfx12. Feature suffixes such asgfx1250:xnack-must not change tile selection.ABQuant FP4 dispatches through an FP8 compute type. Its old
warp_tile_k=32selects an MFMA fragment that produces incorrect results on gfx1250. The default is now the GPU-validated128; this is a deliberate choice rather than a direct copy ofget_k_warp_tile, which returns64on that WMMA path.Untargeted factory calls retain a gfx950 preview for offline naming. At build time, unchanged previews are regenerated for the actual target so architecture-dependent tiles and pipelines are selected together. Explicitly targeted configs retain their architecture checks.
Validation
Recorded MI400 validation with ROCm 10.0 at the standalone-build fix,
ff4646158be7:dispatcher_gemm_libbuilt and verified two FP16 GEMMs.CPU regression coverage also checks architecture gates, feature-suffixed targets, compile defines, config/build mismatches, deferred defaults, unsafe AQuant overrides, HIP library discovery, and CMake fallback generation.
Limitations
warp_tile_k=64, pending separate validation.