From 0679817e6e4012234dcf553ececd58ca6dd6965a Mon Sep 17 00:00:00 2001 From: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Date: Mon, 17 Aug 2026 21:16:34 -0400 Subject: [PATCH 1/6] feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) The regular-GEMM TE->Dispatcher bridge (#8997 fp16/bf16, #8998 fp8/bf8/int8) was CDNA-only (gfx90a/gfx942/gfx950, MFMA). This enables it on gfx1250 (MI400, RDNA4/WMMA): - gemm_utils: add gfx1250 to _SUPPORTED_ARCHES + wave-combo table; add an OCP fp8/bf8 host codec (float8_e4m3fn / float8_e5m2) selected by arch (gfx950/gfx12xx -> OCP, gfx942 -> FNUZ), since gfx1250's device fp8_t is OCP. - gemm_validation_utils: gfx1250 warp-configuration combos. - arch_specs_generated: gfx1250 fp8/bf8 WMMA warp tile (16x16x64). - gemm_full_benchmark: allow --dtype fp8/bf8. - configs: gfx1250 CI configs (fp16/bf16 warp_tile 16x16x32; fp8 16x16x64). - tests: CPU-only gfx1250/OCP-fp8 codec + arch tests. Validated on MI400/gfx1250 (ROCm 10.1): fp16/bf16 and fp8 kernels build and pass --verify (fp8 max_rel ~3.6e-2 vs fp32 ref, i.e. E4M3 quantization only). --- .../codegen/arch_specs_generated.py | 2 + .../dispatcher/python/gemm_utils.py | 40 +++++++- .../dispatcher/tests/test_gemm_utils.py | 48 ++++++++++ .../configs/default_ci_config_gfx1250.json | 92 +++++++++++++++++++ .../default_ci_config_gfx1250_fp8.json | 92 +++++++++++++++++++ .../ops/gemm/gemm_full_benchmark.py | 2 +- .../ops/gemm/gemm_validation_utils.py | 1 + 7 files changed, 271 insertions(+), 6 deletions(-) create mode 100644 projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250.json create mode 100644 projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250_fp8.json diff --git a/projects/composablekernel/dispatcher/codegen/arch_specs_generated.py b/projects/composablekernel/dispatcher/codegen/arch_specs_generated.py index c3137cf32b2f..5bdca2cccd73 100644 --- a/projects/composablekernel/dispatcher/codegen/arch_specs_generated.py +++ b/projects/composablekernel/dispatcher/codegen/arch_specs_generated.py @@ -109,6 +109,8 @@ "gfx1250": { "fp16_fp16_fp32": [[16, 16, 32]], "bf16_bf16_fp32": [[16, 16, 32]], + "fp8_fp8_fp32": [[16, 16, 64]], + "bf8_bf8_fp32": [[16, 16, 64]], }, } diff --git a/projects/composablekernel/dispatcher/python/gemm_utils.py b/projects/composablekernel/dispatcher/python/gemm_utils.py index 42d69a297fac..aed408d336d3 100644 --- a/projects/composablekernel/dispatcher/python/gemm_utils.py +++ b/projects/composablekernel/dispatcher/python/gemm_utils.py @@ -50,7 +50,7 @@ # --- Bridge shared helpers (canonical superset; byte-identical across bridges) --- # Supported GPU architectures for the bridge (single source of truth). -_SUPPORTED_ARCHES = ("gfx90a", "gfx942", "gfx950") +_SUPPORTED_ARCHES = ("gfx90a", "gfx942", "gfx950", "gfx1250") # Single source of truth for the preshuffle B-shuffle permutation used by the # bridge. The bridge codegen only emits the NON-permuteN preshuffle pipeline @@ -968,6 +968,33 @@ def _bf8_u8_to_fp32(u8: np.ndarray) -> np.ndarray: return _fnuz_decode_table(5, 2)[u8.astype(np.intp)] +def _fp32_to_fp8_ocp_u8(x): + """Encode fp32 -> fp8 E4M3 **OCP** (gfx950/gfx12xx) bit pattern (uint8).""" + import ml_dtypes + return np.ascontiguousarray(x, dtype=np.float32).astype(ml_dtypes.float8_e4m3fn).view(np.uint8) + + +def _fp32_to_bf8_ocp_u8(x): + """Encode fp32 -> bf8 E5M2 **OCP** (gfx950/gfx12xx) bit pattern (uint8).""" + import ml_dtypes + return np.ascontiguousarray(x, dtype=np.float32).astype(ml_dtypes.float8_e5m2).view(np.uint8) + + +_OCP_FP8_CACHE = {} + + +def _use_ocp_fp8(): + """True on archs whose device fp8_t is OCP (gfx950/MI350, gfx12xx/RDNA) rather + than the FNUZ encoding used by gfx942/MI300. Detected once via rocminfo.""" + if "v" not in _OCP_FP8_CACHE: + try: + a = _get_arch() + except Exception: + a = "" + _OCP_FP8_CACHE["v"] = a.startswith("gfx12") or a == "gfx950" + return _OCP_FP8_CACHE["v"] + + # Output (C) element dtype for an A/B element dtype, mirroring the codegen's # CommonTypeMappings.get_output_dtype: fp8/bf8 accumulate into fp16, int8 into # int32, everything else stores in its own dtype. @@ -1047,11 +1074,13 @@ def run( A_h = _fp32_to_bf16_u16(A_lay) B_h = _fp32_to_bf16_u16(B_lay) elif dtype == "fp8": - A_h = _fp32_to_fp8_u8(A_lay) - B_h = _fp32_to_fp8_u8(B_lay) + _enc = _fp32_to_fp8_ocp_u8 if _use_ocp_fp8() else _fp32_to_fp8_u8 + A_h = _enc(A_lay) + B_h = _enc(B_lay) elif dtype == "bf8": - A_h = _fp32_to_bf8_u8(A_lay) - B_h = _fp32_to_bf8_u8(B_lay) + _enc = _fp32_to_bf8_ocp_u8 if _use_ocp_fp8() else _fp32_to_bf8_u8 + A_h = _enc(A_lay) + B_h = _enc(B_lay) elif dtype == "int8": A_h = np.ascontiguousarray(A_lay, dtype=np.int8) B_h = np.ascontiguousarray(B_lay, dtype=np.int8) @@ -2130,6 +2159,7 @@ def _dim_bad(repeat: int, warp_tile: int) -> bool: "gfx90a": [[1, 4, 1], [2, 2, 1], [4, 1, 1]], "gfx942": [[1, 4, 1], [2, 2, 1], [4, 1, 1]], "gfx950": [[1, 4, 1], [2, 2, 1], [4, 1, 1]], + "gfx1250": [[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]], "gfx1201": [[2, 4, 1], [1, 8, 1], [8, 1, 1], [4, 2, 1]], } diff --git a/projects/composablekernel/dispatcher/tests/test_gemm_utils.py b/projects/composablekernel/dispatcher/tests/test_gemm_utils.py index 14afc24a2b19..e4e6793c1b72 100644 --- a/projects/composablekernel/dispatcher/tests/test_gemm_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_gemm_utils.py @@ -285,3 +285,51 @@ def test_multi_abd_runner_has_no_stray_multi_d_code(self): if __name__ == "__main__": unittest.main() + + +# --- gfx1250 (MI400 / RDNA4-WMMA) enablement ------------------------------- +# The regular-GEMM bridge historically allow-listed only CDNA (gfx90a/942/950) +# and carried FNUZ-only fp8 codecs. gfx1250 uses WMMA + OCP fp8, so it needs an +# arch entry and an OCP codec path. These CPU-only tests lock that surface in. +from gemm_utils import ( # noqa: E402 + _SUPPORTED_ARCHES, + _fp32_to_fp8_ocp_u8, + _fp32_to_bf8_ocp_u8, + _use_ocp_fp8, +) + + +class TestGfx1250Fp8Ocp(unittest.TestCase): + def test_gfx1250_in_supported_arches(self): + self.assertIn("gfx1250", _SUPPORTED_ARCHES) + + def test_ocp_fp8_shape_and_dtype(self): + x = (np.random.RandomState(0).randn(8, 16) * 0.1).astype(np.float32) + u8 = _fp32_to_fp8_ocp_u8(x) + self.assertEqual(u8.dtype, np.uint8) + self.assertEqual(u8.shape, x.shape) + + def test_ocp_fp8_roundtrip_exact_values(self): + # 1.0, 0.5, 2.0, -1.0 are exactly representable in fp8 E4M3 (OCP). + import ml_dtypes + + x = np.array([[1.0, 0.5, 2.0, -1.0]], dtype=np.float32) + u8 = _fp32_to_fp8_ocp_u8(x) + back = u8.view(ml_dtypes.float8_e4m3fn).astype(np.float32) + np.testing.assert_array_equal(back, x) + + def test_ocp_differs_from_fnuz(self): + # OCP and FNUZ are distinct encodings; the byte patterns must not match + # for a generic value (guards against silently reusing the FNUZ codec). + x = np.array([[0.1, 0.3, 1.5]], dtype=np.float32) + self.assertFalse( + np.array_equal(_fp32_to_fp8_ocp_u8(x), _fp32_to_fp8_u8(x)) + ) + + def test_ocp_bf8_shape(self): + x = (np.random.RandomState(1).randn(4, 4) * 0.1).astype(np.float32) + self.assertEqual(_fp32_to_bf8_ocp_u8(x).shape, x.shape) + + def test_use_ocp_fp8_is_callable(self): + # Returns a bool; on a CPU-only box arch detection may fail -> False. + self.assertIsInstance(_use_ocp_fp8(), bool) diff --git a/projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250.json b/projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250.json new file mode 100644 index 000000000000..d82536909948 --- /dev/null +++ b/projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250.json @@ -0,0 +1,92 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 32 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv3", + "compv4", + "mem" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle", + "default" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "k_block_per_cu": 1 +} \ No newline at end of file diff --git a/projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250_fp8.json b/projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250_fp8.json new file mode 100644 index 000000000000..630912d5ae47 --- /dev/null +++ b/projects/composablekernel/tile_engine/ops/gemm/configs/default_ci_config_gfx1250_fp8.json @@ -0,0 +1,92 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 64 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv3", + "compv4", + "mem" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle", + "default" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "k_block_per_cu": 1 +} \ No newline at end of file diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_full_benchmark.py b/projects/composablekernel/tile_engine/ops/gemm/gemm_full_benchmark.py index db171f8d5e9c..27c3f4bfb038 100644 --- a/projects/composablekernel/tile_engine/ops/gemm/gemm_full_benchmark.py +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_full_benchmark.py @@ -86,7 +86,7 @@ {"M": 257, "N": 257, "K": 257}, ] -SUPPORTED_DTYPES = ("fp16", "bf16") +SUPPORTED_DTYPES = ("fp16", "bf16", "fp8", "bf8") # Row-major C only: ck_tile's universal GEMM rejects column-major C at build. # The 4-char codes (rcrr, ...) are the multi_abd A,B,E,D layouts; TE gemm_multi_abd # only supports rcrr today. diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_validation_utils.py b/projects/composablekernel/tile_engine/ops/gemm/gemm_validation_utils.py index 183a5824588a..ad262a249532 100644 --- a/projects/composablekernel/tile_engine/ops/gemm/gemm_validation_utils.py +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_validation_utils.py @@ -45,6 +45,7 @@ def get_warp_size_for_gpu(gpu_target: str) -> int: WARP_SUPPORTED_COMBINATIONS = { + "gfx1250": [[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]], "gfx90a": [ [1, 4, 1], [2, 2, 1], From ff22c8852ecdc134315421c3fc31e23314805380 Mon Sep 17 00:00:00 2001 From: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Date: Tue, 18 Aug 2026 01:33:51 -0400 Subject: [PATCH 2/6] feat(ck-tile): add grouped/multi-D/multi-ABD gfx1250 configs + tests Fold the grouped (#9000), multi-D (#9308), and multi-ABD (#9305) variant enablement onto the shared gfx1250 foundation so gemm_universal + these three "full" variants ship in one PR: - grouped: gfx1250 CI config (WMMA 16x16x32) + CPU tests - multi-D: gfx1250 CI config (WMMA 16x16x32) + CPU tests + --arch help - multi-ABD: gfx1250 CI config (WMMA 16x16x32, fp16-only) + CPU tests Each merged variant pinned the CDNA MFMA 32x32x16 tile in its CI config, which does not run on gfx1250 WMMA; the new configs use 16x16x32. No variant codegen change was needed beyond the shared arch enablement. Validated on MI400/gfx1250 (ROCm 10.1): grouped 64/64 fp16 OK; multi-D 80 measurements VERIFY (max_rel ~5e-4); multi-ABD 32/32 fp16 VERIFY. 63 CPU tests pass. --- .../dispatcher/tests/test_gemm_utils.py | 34 ++++++ .../dispatcher/tests/test_multi_abd_bridge.py | 51 +++++++++ .../dispatcher/tests/test_multi_d_bridge.py | 35 ++++++ .../configs/default_ci_config_gfx1250.json | 92 +++++++++++++++ .../configs/default_ci_config_gfx1250.json | 106 ++++++++++++++++++ .../ops/gemm/gemm_multi_d_full_benchmark.py | 5 +- .../configs/default_ci_config_gfx1250.json | 92 +++++++++++++++ 7 files changed, 413 insertions(+), 2 deletions(-) create mode 100644 projects/composablekernel/tile_engine/ops/gemm/gemm_multi_abd/configs/default_ci_config_gfx1250.json create mode 100644 projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config_gfx1250.json create mode 100644 projects/composablekernel/tile_engine/ops/gemm/grouped_gemm/configs/default_ci_config_gfx1250.json diff --git a/projects/composablekernel/dispatcher/tests/test_gemm_utils.py b/projects/composablekernel/dispatcher/tests/test_gemm_utils.py index e4e6793c1b72..bcf7fafe0cd9 100644 --- a/projects/composablekernel/dispatcher/tests/test_gemm_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_gemm_utils.py @@ -333,3 +333,37 @@ def test_ocp_bf8_shape(self): def test_use_ocp_fp8_is_callable(self): # Returns a bool; on a CPU-only box arch detection may fail -> False. self.assertIsInstance(_use_ocp_fp8(), bool) + + +# --- grouped GEMM on gfx1250 (MI400) --------------------------------------- +# The grouped bridge (#9000) shares the arch gate + WMMA warp tiles with the +# regular bridge, so it runs on gfx1250 once gfx1250 is enabled. Its CI config +# already uses the gfx1250-valid WMMA warp tile 16x16x32. These CPU-only checks +# lock that in (multi-problem plumbing + gfx1250 arch acceptance). +from gemm_utils import GroupedGemmProblem, _resolve_arch # noqa: E402 + + +class TestGroupedGfx1250(unittest.TestCase): + def test_grouped_problem_roundtrip_and_flops(self): + groups = [(1024, 1024, 1024), (512, 2048, 256)] + p = GroupedGemmProblem(groups=groups) + self.assertEqual(p.group_count, 2) + self.assertEqual(GroupedGemmProblem.from_dict(p.to_dict()).groups, groups) + self.assertEqual(p.flops, sum(2.0 * m * n * k for (m, n, k) in groups)) + + def test_gfx1250_arch_accepted(self): + self.assertEqual(_resolve_arch("gfx1250"), "gfx1250") + + def test_grouped_gfx1250_ci_config_uses_wmma_warp_tile(self): + import json as _json + from pathlib import Path as _Path + + cfg = ( + _Path(__file__).parent.parent.parent + / "tile_engine/ops/gemm/grouped_gemm/configs/default_ci_config_gfx1250.json" + ) + tc = _json.load(open(cfg))["tile_config"] + # gfx1250 fp16/bf16 WMMA is 16x16x32 (not the CDNA MFMA 32x32x16). + self.assertEqual(tc["warp_tile_m"]["values"], [16]) + self.assertEqual(tc["warp_tile_n"]["values"], [16]) + self.assertEqual(tc["warp_tile_k"]["values"], [32]) diff --git a/projects/composablekernel/dispatcher/tests/test_multi_abd_bridge.py b/projects/composablekernel/dispatcher/tests/test_multi_abd_bridge.py index 9f503a1fb5f0..a15ecb8502e5 100644 --- a/projects/composablekernel/dispatcher/tests/test_multi_abd_bridge.py +++ b/projects/composablekernel/dispatcher/tests/test_multi_abd_bridge.py @@ -258,5 +258,56 @@ def test_run_forwards_nonzero_status(self): self.assertEqual(status, -3) +class TestMultiAbdGfx1250Config(unittest.TestCase): + """The gfx1250 (MI400) multi_abd sweep config must be WMMA-shaped. + + gfx1250 runs WMMA, not MFMA: the fp16/bf16 warp tile is 16x16x32, and the + 32x32x16 MFMA tile shipped in ``default_ci_config.json`` never instantiates + a runnable kernel there. ``default_ci_config_gfx1250.json`` is the arch + variant that swaps in the WMMA warp tile while keeping every other sweep + axis identical, so the bridge codegens kernels that actually launch on MI400. + """ + + _GFX1250_CONFIG = _CONFIG_DIR / "default_ci_config_gfx1250.json" + + def test_gfx1250_config_exists(self): + self.assertTrue( + self._GFX1250_CONFIG.is_file(), + f"missing gfx1250 multi_abd config: {self._GFX1250_CONFIG}", + ) + + def test_gfx1250_config_uses_wmma_warp_tile(self): + with open(self._GFX1250_CONFIG) as f: + tc = json.load(f)["tile_config"] + # gfx1250 fp16/bf16 WMMA warp tile is 16x16x32 (NOT the 32x32x16 MFMA + # tile the CDNA CI config ships); a stale MFMA tile here would silently + # produce zero runnable kernels on MI400. + self.assertEqual(tc["warp_tile_m"]["values"], [16]) + self.assertEqual(tc["warp_tile_n"]["values"], [16]) + self.assertEqual(tc["warp_tile_k"]["values"], [32]) + + def test_gfx1250_config_wave_triple_is_supported(self): + # The swept warps-per-block triple must be one gfx1250 actually supports, + # or expand_sweep gates the whole sweep away to an empty kernel set. + sys.path.insert( + 0, str(REPO_ROOT / "tile_engine" / "ops" / "gemm") + ) + from gemm_validation_utils import ( # noqa: E402 + WARP_SUPPORTED_COMBINATIONS, + ) + + self.assertIn("gfx1250", WARP_SUPPORTED_COMBINATIONS) + with open(self._GFX1250_CONFIG) as f: + tc = json.load(f)["tile_config"] + for wm in tc["warp_m"]["values"]: + for wn in tc["warp_n"]["values"]: + for wk in tc["warp_k"]["values"]: + self.assertIn( + [wm, wn, wk], + WARP_SUPPORTED_COMBINATIONS["gfx1250"], + f"wave triple [{wm},{wn},{wk}] not supported on gfx1250", + ) + + if __name__ == "__main__": unittest.main() diff --git a/projects/composablekernel/dispatcher/tests/test_multi_d_bridge.py b/projects/composablekernel/dispatcher/tests/test_multi_d_bridge.py index 61ebf4e06633..5742154c9e83 100644 --- a/projects/composablekernel/dispatcher/tests/test_multi_d_bridge.py +++ b/projects/composablekernel/dispatcher/tests/test_multi_d_bridge.py @@ -171,5 +171,40 @@ def test_configs_are_valid_sweeps(self): self.assertIn(key, tr, f"{path.name} missing {key}") +class TestMultiDGfx1250Config(unittest.TestCase): + """The gfx1250 (MI400) CI config must use WMMA warp tiles. + + gfx1250 has no MFMA units: it runs the RDNA4 WMMA path whose fp16/bf16 warp + tile is 16x16x32 (see arch_specs_generated.py). The merged multi_d bridge + (#9308) shipped only an MFMA CI config (32x32x16, valid on gfx942/gfx950), + which the kernel reports as unsupported (status -2/-1) on gfx1250. The + gfx1250 CI config therefore pins the WMMA warp tile so the sweep produces + kernels that actually run on MI400. + """ + + _GFX1250_CONFIG = _CONFIG_DIR / "default_ci_config_gfx1250.json" + + def test_gfx1250_config_exists(self): + self.assertTrue(self._GFX1250_CONFIG.is_file(), self._GFX1250_CONFIG) + + def test_gfx1250_config_uses_wmma_warp_tile(self): + with open(self._GFX1250_CONFIG) as f: + tc = json.load(f)["tile_config"] + # WMMA (RDNA4) fp16/bf16 warp tile, not MFMA (CDNA) 32x32x16. + self.assertEqual(tc["warp_tile_m"]["values"], [16]) + self.assertEqual(tc["warp_tile_n"]["values"], [16]) + self.assertEqual(tc["warp_tile_k"]["values"], [32]) + + def test_gfx1250_config_keeps_multi_d_block(self): + # gfx1250 enablement must not drop the D-fusion knobs that make this a + # multi_d config rather than a plain GEMM config. + with open(self._GFX1250_CONFIG) as f: + data = json.load(f) + self.assertIn("multi_d_config", data) + md = data["multi_d_config"] + self.assertIn("elementwise_ops", md) + self.assertIn("num_d_tensors", md) + + if __name__ == "__main__": unittest.main() diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_abd/configs/default_ci_config_gfx1250.json b/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_abd/configs/default_ci_config_gfx1250.json new file mode 100644 index 000000000000..38376a410b01 --- /dev/null +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_abd/configs/default_ci_config_gfx1250.json @@ -0,0 +1,92 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 32 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv3", + "compv4", + "mem" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle", + "default" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "k_block_per_cu": 1 +} diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config_gfx1250.json b/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config_gfx1250.json new file mode 100644 index 000000000000..9dfe7eb323cc --- /dev/null +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config_gfx1250.json @@ -0,0 +1,106 @@ +{ + "tile_config": { + "tile_m": { + "values": [ + 128 + ] + }, + "tile_n": { + "values": [ + 128 + ] + }, + "tile_k": { + "values": [ + 64 + ] + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 32 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv3", + "compv4", + "mem" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle", + "default" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "multi_d_config": { + "elementwise_ops": { + "values": [ + "MultiDAdd", + "MultiDMultiply" + ] + }, + "num_d_tensors": { + "values": [ + 1, + 2 + ] + } + }, + "k_block_per_cu": 1 +} diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py b/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py index 74037d1b1e61..9cb12c331500 100644 --- a/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py @@ -245,8 +245,9 @@ def main(): parser.add_argument( "--arch", default=None, - help="GPU arch (gfx90a/gfx942/gfx950); default: auto-detect via rocminfo " - "(no silent gfx942 default). Flows into expand_sweep, which resolves None.", + help="GPU arch (gfx90a/gfx942/gfx950/gfx1250); default: auto-detect via " + "rocminfo (no silent gfx942 default). Flows into expand_sweep, which " + "resolves None.", ) parser.add_argument( "--dtype", diff --git a/projects/composablekernel/tile_engine/ops/gemm/grouped_gemm/configs/default_ci_config_gfx1250.json b/projects/composablekernel/tile_engine/ops/gemm/grouped_gemm/configs/default_ci_config_gfx1250.json new file mode 100644 index 000000000000..d82536909948 --- /dev/null +++ b/projects/composablekernel/tile_engine/ops/gemm/grouped_gemm/configs/default_ci_config_gfx1250.json @@ -0,0 +1,92 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 32 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv3", + "compv4", + "mem" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle", + "default" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "k_block_per_cu": 1 +} \ No newline at end of file From 82b53687a9e1c35400400e429530a9f35df70a21 Mon Sep 17 00:00:00 2001 From: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Date: Tue, 18 Aug 2026 03:16:03 -0400 Subject: [PATCH 3/6] feat(ck-tile): enable grouped AQuant/ABQuant/BQuant GEMM bridges on gfx1250 (MI400) Enables the three merged grouped quant GEMM TE->dispatcher bridges on gfx1250, stacked on the shared gfx1250 foundation (PR #10921). C++ (small, required): the three ctypes-lib runtime arch allow-lists rejected gfx1250; add gfx1250 to each gate and to the "(supported: ...)" message. clang-format-18 applied. Python: add gfx1250 default-config helpers that pin warp_tile_m/n/k=16 so the codegen selects the gfx12 WMMA instruction (warp_tile_k=32/128 select the gfx9-only MFMA/FlatMM paths, which silently return zeros on gfx12). ABQuant uses the standard CompV3 path (not the gfx950-native eightwaves/FlatMM path). Tests: CPU unit tests assert the gfx1250 configs use WMMA warp_tile_k=16, propagate gfx_arch, and (abquant) use compv3 rather than eightwaves. CPU + codegen validated (176/176 quant utils tests pass); on-GPU validation pending gfx1250 node recovery. --- .../grouped_gemm_abquant_ctypes_lib.cpp | 4 +- .../ctypes/grouped_gemm_aquant_ctypes_lib.cpp | 4 +- .../ctypes/grouped_gemm_bquant_ctypes_lib.cpp | 4 +- .../python/grouped_gemm_abquant_utils.py | 61 ++++++++++ .../python/grouped_gemm_aquant_utils.py | 105 ++++++++++++++++++ .../python/grouped_gemm_bquant_utils.py | 101 +++++++++++++++++ .../tests/test_grouped_gemm_abquant_utils.py | 47 ++++++++ .../tests/test_grouped_gemm_aquant_utils.py | 56 ++++++++++ .../tests/test_grouped_gemm_bquant_utils.py | 52 +++++++++ 9 files changed, 428 insertions(+), 6 deletions(-) diff --git a/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_abquant_ctypes_lib.cpp b/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_abquant_ctypes_lib.cpp index 844e54796eae..90d6ad570e60 100644 --- a/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_abquant_ctypes_lib.cpp +++ b/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_abquant_ctypes_lib.cpp @@ -79,10 +79,10 @@ int dispatcher_initialize() } g_gfx_arch = props.gcnArchName; if(g_gfx_arch.rfind("gfx950", 0) != 0 && g_gfx_arch.rfind("gfx942", 0) != 0 && - g_gfx_arch.rfind("gfx90a", 0) != 0) + g_gfx_arch.rfind("gfx90a", 0) != 0 && g_gfx_arch.rfind("gfx1250", 0) != 0) { std::cerr << "dispatcher_initialize: unsupported GPU architecture '" << g_gfx_arch - << "' (supported: gfx90a, gfx942, gfx950)\n"; + << "' (supported: gfx90a, gfx942, gfx950, gfx1250)\n"; return -1; } g_initialized = true; diff --git a/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_aquant_ctypes_lib.cpp b/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_aquant_ctypes_lib.cpp index 0597217ff1ec..1e6c84ff9ea5 100644 --- a/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_aquant_ctypes_lib.cpp +++ b/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_aquant_ctypes_lib.cpp @@ -77,10 +77,10 @@ int dispatcher_initialize() } g_gfx_arch = props.gcnArchName; if(g_gfx_arch.rfind("gfx950", 0) != 0 && g_gfx_arch.rfind("gfx942", 0) != 0 && - g_gfx_arch.rfind("gfx90a", 0) != 0) + g_gfx_arch.rfind("gfx90a", 0) != 0 && g_gfx_arch.rfind("gfx1250", 0) != 0) { std::cerr << "dispatcher_initialize: unsupported GPU architecture '" << g_gfx_arch - << "' (supported: gfx90a, gfx942, gfx950)\n"; + << "' (supported: gfx90a, gfx942, gfx950, gfx1250)\n"; return -1; } g_initialized = true; diff --git a/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_bquant_ctypes_lib.cpp b/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_bquant_ctypes_lib.cpp index 22e0b7db9b89..756c1760dd07 100644 --- a/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_bquant_ctypes_lib.cpp +++ b/projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_bquant_ctypes_lib.cpp @@ -145,10 +145,10 @@ int dispatcher_run_bquant_gemm(const void* A, } const std::string arch(props.gcnArchName); if(arch.rfind("gfx950", 0) != 0 && arch.rfind("gfx942", 0) != 0 && - arch.rfind("gfx90a", 0) != 0) + arch.rfind("gfx90a", 0) != 0 && arch.rfind("gfx1250", 0) != 0) { std::cerr << "dispatcher_run_bquant_gemm: unsupported GPU architecture '" << arch - << "' (supported: gfx90a, gfx942, gfx950)\n"; + << "' (supported: gfx90a, gfx942, gfx950, gfx1250)\n"; return -1; } } diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py index a55458104fac..bd7738c74e7a 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py @@ -857,3 +857,64 @@ def default_bf8_preshuffleb_config( k_block_per_cu=2, gfx_arch=gfx_arch, ) + + +# ============================================================================= +# gfx1250 (MI400 / RDNA WMMA) default configs +# ============================================================================= +# gfx1250 uses the standard CompV3 pipeline (NOT the gfx950-native eightwaves / +# FlatMM / preshuffleb paths, which require CK_GFX950_SUPPORT and warp_tile_k=128 +# and TransposeC=True). On gfx1250 the codegen selects the gfx12 WMMA instruction +# when warp_tile_k=16 (warp_tile_k=32 selects the gfx9-only MFMA path, which +# silently returns zeros on gfx12). These helpers mirror default_*_compv3_config +# but pin warp_tile_m=warp_tile_n=warp_tile_k=16 (WMMA) and keep transpose_c=False. +# The compile path additionally injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 +# for gfx12 archs, so fp8/bf8 use the OCP encoding on gfx1250. + +_GFX1250_ARCH = "gfx1250" + + +def default_fp8_compv3_config_gfx1250( + quant_group_k: int = 128, + bquant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> ABQuantKernelConfig: + """fp8 ABQuant CompV3 config for gfx1250 (WMMA warp_tile_k=16, TransposeC=False).""" + return ABQuantKernelConfig( + variant_key="fp8", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=128, tile_n=128, tile_k=128, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + aquant_group_m=1, aquant_group_n=1, aquant_group_k=quant_group_k, + bquant_group_m=1, bquant_group_n=bquant_group_n, bquant_group_k=quant_group_k, + preshuffle_b=False, preshuffle_aq=False, preshuffle_bq=False, + transpose_c=False, + gfx_arch=gfx_arch, + ) + + +def default_bf8_compv3_config_gfx1250( + quant_group_k: int = 128, + bquant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> ABQuantKernelConfig: + """bf8 ABQuant CompV3 config for gfx1250 (WMMA warp_tile_k=16, TransposeC=False).""" + return ABQuantKernelConfig( + variant_key="bf8", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=128, tile_n=128, tile_k=128, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + aquant_group_m=1, aquant_group_n=1, aquant_group_k=quant_group_k, + bquant_group_m=1, bquant_group_n=bquant_group_n, bquant_group_k=quant_group_k, + preshuffle_b=False, preshuffle_aq=False, preshuffle_bq=False, + transpose_c=False, + gfx_arch=gfx_arch, + ) diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py index 60e529446e2c..5eda55da6e08 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py @@ -785,3 +785,108 @@ def default_bf8_preshuffleaq_config( preshuffle_aq=True, gfx_arch=gfx_arch, ) + + +# ============================================================================= +# gfx1250 (MI400 / RDNA WMMA) default configs +# ============================================================================= +# On gfx1250 the quant MFMA path is unavailable; the codegen selects the gfx12 +# WMMA instruction when warp_tile_k=16 (warp_tile_k=32 selects the gfx9-only +# MFMA path, which silently returns zeros on gfx12). These helpers mirror the +# gfx9 defaults above but pin warp_tile_m=warp_tile_n=warp_tile_k=16 (WMMA). +# The compile path additionally injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 +# for gfx12 archs, so fp8/bf8 use the OCP encoding on gfx1250. + +_GFX1250_ARCH = "gfx1250" + + +def default_fp8_config_gfx1250( + quant_group_k: int = 128, + quant_group_m: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> AQuantKernelConfig: + """fp8 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" + return AQuantKernelConfig( + variant_key="fp8", + layout="rcr", + pipeline="mem", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=quant_group_m, + quant_group_n=1, + quant_group_k=quant_group_k, + preshuffle_aq=False, + gfx_arch=gfx_arch, + ) + + +def default_bf8_config_gfx1250( + quant_group_k: int = 128, + quant_group_m: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> AQuantKernelConfig: + """bf8 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" + return AQuantKernelConfig( + variant_key="bf8", + layout="rcr", + pipeline="mem", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=quant_group_m, + quant_group_n=1, + quant_group_k=quant_group_k, + preshuffle_aq=False, + gfx_arch=gfx_arch, + ) + + +def default_fp8i4_config_gfx1250( + quant_group_k: int = 128, + quant_group_m: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> AQuantKernelConfig: + """fp8i4 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" + return AQuantKernelConfig( + variant_key="fp8i4", + layout="rcr", + pipeline="mem", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=quant_group_m, + quant_group_n=1, + quant_group_k=quant_group_k, + preshuffle_aq=False, + gfx_arch=gfx_arch, + ) + + +def default_bf8i4_config_gfx1250( + quant_group_k: int = 128, + quant_group_m: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> AQuantKernelConfig: + """bf8i4 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" + return AQuantKernelConfig( + variant_key="bf8i4", + layout="rcr", + pipeline="mem", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=quant_group_m, + quant_group_n=1, + quant_group_k=quant_group_k, + preshuffle_aq=False, + gfx_arch=gfx_arch, + ) diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py index 92fb0038c8f7..a8582338f193 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py @@ -1214,3 +1214,104 @@ def default_mx_bf16fp4_config( quant_group_k=quant_group_k, gfx_arch=gfx_arch, ) + + +# ============================================================================= +# gfx1250 (MI400 / RDNA WMMA) default configs +# ============================================================================= +# On gfx1250 the quant path uses the gfx12 WMMA instruction, selected by the +# codegen when warp_tile_k=16 (warp_tile_k=32/128 select the gfx9-only MFMA/FlatMM +# paths, which silently return zeros on gfx12). These helpers mirror the standard +# gfx9 compv3 defaults but pin warp_tile_m=warp_tile_n=warp_tile_k=16 (WMMA). The +# compile path additionally injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for +# gfx12 archs, so fp8/bf8 use the OCP encoding on gfx1250. + +_GFX1250_ARCH = "gfx1250" + + +def default_fp8_config_gfx1250( + quant_group_k: int = 128, + quant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> BQuantKernelConfig: + """fp8 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" + return BQuantKernelConfig( + variant_key="fp8", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=1, + quant_group_n=quant_group_n, + quant_group_k=quant_group_k, + gfx_arch=gfx_arch, + ) + + +def default_bf8_config_gfx1250( + quant_group_k: int = 128, + quant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> BQuantKernelConfig: + """bf8 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" + return BQuantKernelConfig( + variant_key="bf8", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=1, + quant_group_n=quant_group_n, + quant_group_k=quant_group_k, + gfx_arch=gfx_arch, + ) + + +def default_fp8i4_config_gfx1250( + quant_group_k: int = 128, + quant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> BQuantKernelConfig: + """fp8i4 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" + return BQuantKernelConfig( + variant_key="fp8i4", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=1, + quant_group_n=quant_group_n, + quant_group_k=quant_group_k, + gfx_arch=gfx_arch, + ) + + +def default_bf8i4_config_gfx1250( + quant_group_k: int = 128, + quant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> BQuantKernelConfig: + """bf8i4 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" + return BQuantKernelConfig( + variant_key="bf8i4", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + quant_group_m=1, + quant_group_n=quant_group_n, + quant_group_k=quant_group_k, + gfx_arch=gfx_arch, + ) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py index 1d102e28f00c..2ca1f074c739 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py @@ -361,3 +361,50 @@ def test_transpose_c_in_name_only_when_true(self): without_tc = default_fp8_compv3_config() assert "transposec" in with_tc.name assert "transposec" not in without_tc.name + + +# ============================================================================= +# gfx1250 (MI400 / WMMA) default configs +# ============================================================================= + + +from grouped_gemm_abquant_utils import ( # noqa: E402 + default_fp8_compv3_config_gfx1250, + default_bf8_compv3_config_gfx1250, +) + + +class TestGfx1250Configs: + + def _all(self): + return [ + default_fp8_compv3_config_gfx1250(), + default_bf8_compv3_config_gfx1250(), + ] + + def test_gfx1250_uses_wmma_warp_tile_k_16(self): + # gfx1250 must use the standard CompV3 pipeline with the WMMA instruction + # (warp_tile_k=16), NOT the gfx950-native eightwaves/FlatMM path. + for cfg in self._all(): + assert cfg.warp_tile_m == 16 + assert cfg.warp_tile_n == 16 + assert cfg.warp_tile_k == 16, f"{cfg.name} must use WMMA warp_tile_k=16" + + def test_gfx1250_uses_compv3_not_eightwaves(self): + # eightwaves / preshuffleb are gfx950-only; gfx1250 uses standard compv3. + for cfg in self._all(): + assert cfg.pipeline == "compv3" + assert cfg.transpose_c is False + assert cfg.preshuffle_b is False + + def test_gfx1250_arch_propagated(self): + for cfg in self._all(): + assert cfg.gfx_arch == "gfx1250" + + def test_gfx1250_names_16x16x16(self): + for cfg in self._all(): + assert "16x16x16" in cfg.name + + def test_gfx1250_unique_names(self): + names = [cfg.name for cfg in self._all()] + assert len(names) == len(set(names)) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py index be95e88c0822..89b80fec9dce 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py @@ -408,3 +408,59 @@ def test_mem_and_compv3_pipelines_produce_different_names(self): assert mem_cfg.name != compv3_cfg.name assert "mem" in mem_cfg.name assert "compv3" in compv3_cfg.name + + +# ============================================================================= +# gfx1250 (MI400 / WMMA) default configs +# ============================================================================= + + +from grouped_gemm_aquant_utils import ( # noqa: E402 + default_fp8_config_gfx1250, + default_bf8_config_gfx1250, + default_fp8i4_config_gfx1250, + default_bf8i4_config_gfx1250, +) + + +class TestGfx1250Configs: + + def _all(self): + return [ + default_fp8_config_gfx1250(), + default_bf8_config_gfx1250(), + default_fp8i4_config_gfx1250(), + default_bf8i4_config_gfx1250(), + ] + + def test_gfx1250_uses_wmma_warp_tile_k_16(self): + # gfx1250 must use the WMMA instruction (warp_tile_k=16); warp_tile_k=32 + # selects the gfx9 MFMA path which silently returns zeros on gfx12. + for cfg in self._all(): + assert cfg.warp_tile_m == 16 + assert cfg.warp_tile_n == 16 + assert cfg.warp_tile_k == 16, f"{cfg.name} must use WMMA warp_tile_k=16" + + def test_gfx1250_arch_propagated(self): + for cfg in self._all(): + assert cfg.gfx_arch == "gfx1250" + + def test_gfx1250_names_16x16x16(self): + for cfg in self._all(): + assert "16x16x16" in cfg.name + + def test_gfx1250_pipeline_is_mem(self): + for cfg in self._all(): + assert cfg.pipeline == "mem" + + def test_gfx1250_layout_is_rcr(self): + for cfg in self._all(): + assert cfg.layout == "rcr" + + def test_gfx1250_no_preshuffle(self): + for cfg in self._all(): + assert cfg.preshuffle_aq is False + + def test_gfx1250_unique_names(self): + names = [cfg.name for cfg in self._all()] + assert len(names) == len(set(names)) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py index 1f96dd70f9b4..dc63b1e3c4cc 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py @@ -724,3 +724,55 @@ def test_name_matches_kernel_config_name(self, tmp_path): quant_group_m=1, quant_group_n=1, quant_group_k=128, ) assert configs[0].name == manual.name + + +# ============================================================================= +# gfx1250 (MI400 / WMMA) default configs +# ============================================================================= + + +from grouped_gemm_bquant_utils import ( # noqa: E402 + default_fp8_config_gfx1250, + default_bf8_config_gfx1250, + default_fp8i4_config_gfx1250, + default_bf8i4_config_gfx1250, +) + + +class TestGfx1250Configs: + + def _all(self): + return [ + default_fp8_config_gfx1250(), + default_bf8_config_gfx1250(), + default_fp8i4_config_gfx1250(), + default_bf8i4_config_gfx1250(), + ] + + def test_gfx1250_uses_wmma_warp_tile_k_16(self): + # gfx1250 must use the WMMA instruction (warp_tile_k=16); the gfx9 fp8/bf8 + # defaults use warp_tile_k=128 (FlatMM) which is invalid on gfx12. + for cfg in self._all(): + assert cfg.warp_tile_m == 16 + assert cfg.warp_tile_n == 16 + assert cfg.warp_tile_k == 16, f"{cfg.name} must use WMMA warp_tile_k=16" + + def test_gfx1250_arch_propagated(self): + for cfg in self._all(): + assert cfg.gfx_arch == "gfx1250" + + def test_gfx1250_names_16x16x16(self): + for cfg in self._all(): + assert "16x16x16" in cfg.name + + def test_gfx1250_pipeline_is_compv3(self): + for cfg in self._all(): + assert cfg.pipeline == "compv3" + + def test_gfx1250_layout_is_rcr(self): + for cfg in self._all(): + assert cfg.layout == "rcr" + + def test_gfx1250_unique_names(self): + names = [cfg.name for cfg in self._all()] + assert len(names) == len(set(names)) From ad8a03201fbc1b7c5448fb942c21cda2976025a4 Mon Sep 17 00:00:00 2001 From: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Date: Tue, 18 Aug 2026 03:56:30 -0400 Subject: [PATCH 4/6] fix(ck-tile): correct gfx1250 grouped-quant configs to FlatMM warp_tile_k=128 On-GPU validation on MI400/gfx1250 (fp32-dequant reference) shows the previous warp_tile_k=16 (WMMA) gfx1250 overrides return all-zeros for every grouped-quant variant. The correct gfx1250 config for fp8/bf8 is the FlatMM tile (warp_tile_k=128): - aquant fp8/bf8: 16x16x16 -> all-zero; stock 16x16x32 -> WRONG (max_rel 1.0); 16x16x128 -> PASS (max_rel ~3e-4 / ~1e-4) - abquant compv3 fp8/bf8: 16x16x16 and 16x16x32 -> all-zero; 16x16x128 -> PASS (max_rel ~3e-4) - bquant fp8/bf8: stock is already 16x16x128 -> PASS (max_rel ~5e-4); 16x16x16 -> all-zero. gfx1250 helper now returns the stock config. Changes: - aquant/abquant gfx1250 helpers now pin warp_tile_k=128 (FlatMM). - bquant gfx1250 helpers delegate to the stock fp8/bf8 config. - Remove fp8i4/bf8i4 gfx1250 helpers: the packed-int4 quant path does not compile on gfx1250 at any warp_tile_k (no gfx12 instruction). The C++ ctypes arch-gate stays open for i4, but on-GPU correctness is deferred. - CPU tests updated to assert warp_tile_k=128 / 16x16x128 names. C++ ctypes arch-gates (add gfx1250) are unchanged and remain correct. --- .../python/grouped_gemm_abquant_utils.py | 38 +++++-- .../python/grouped_gemm_aquant_utils.py | 87 +++++---------- .../python/grouped_gemm_bquant_utils.py | 104 ++++++------------ .../tests/test_grouped_gemm_abquant_utils.py | 17 ++- .../tests/test_grouped_gemm_aquant_utils.py | 19 ++-- .../tests/test_grouped_gemm_bquant_utils.py | 20 ++-- 6 files changed, 117 insertions(+), 168 deletions(-) diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py index bd7738c74e7a..5248d34b5ab8 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py @@ -860,16 +860,23 @@ def default_bf8_preshuffleb_config( # ============================================================================= -# gfx1250 (MI400 / RDNA WMMA) default configs +# gfx1250 (MI400) default configs # ============================================================================= # gfx1250 uses the standard CompV3 pipeline (NOT the gfx950-native eightwaves / -# FlatMM / preshuffleb paths, which require CK_GFX950_SUPPORT and warp_tile_k=128 -# and TransposeC=True). On gfx1250 the codegen selects the gfx12 WMMA instruction -# when warp_tile_k=16 (warp_tile_k=32 selects the gfx9-only MFMA path, which -# silently returns zeros on gfx12). These helpers mirror default_*_compv3_config -# but pin warp_tile_m=warp_tile_n=warp_tile_k=16 (WMMA) and keep transpose_c=False. -# The compile path additionally injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 -# for gfx12 archs, so fp8/bf8 use the OCP encoding on gfx1250. +# preshuffleb paths, which require CK_GFX950_SUPPORT and TransposeC=True). +# +# Empirically (on-GPU, MI400/gfx1250, HIP_VISIBLE_DEVICES=0, fp32-dequant +# reference), the ABQuant CompV3 fp8/bf8 kernels are correct on gfx1250 only +# with warp_tile_k=128 (the FlatMM tile). The gfx9 stock CompV3 default uses +# warp_tile_k=32 (plain MFMA) which returns all-zeros on gfx1250; warp_tile_k=16 +# (gfx12 WMMA) also returns all-zeros. So the only correct gfx1250 override for +# the CompV3 pipeline is 16x16x128 with transpose_c=False. +# +# fp8: warp_tile 16x16x128 -> nonzero, max_rel_err ~3e-4 (PASS) +# bf8: warp_tile 16x16x128 -> nonzero, max_rel_err ~3e-4 (PASS) +# +# The compile path injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for gfx12 +# archs, so fp8/bf8 use the OCP encoding on gfx1250. _GFX1250_ARCH = "gfx1250" @@ -879,7 +886,11 @@ def default_fp8_compv3_config_gfx1250( bquant_group_n: int = 1, gfx_arch: str = _GFX1250_ARCH, ) -> ABQuantKernelConfig: - """fp8 ABQuant CompV3 config for gfx1250 (WMMA warp_tile_k=16, TransposeC=False).""" + """fp8 ABQuant CompV3 config for gfx1250 (FlatMM warp_tile_k=128, TransposeC=False). + + GPU-verified on MI400/gfx1250: max_rel_err ~3e-4 vs. fp32 dequant reference. + warp_tile_k=32 (gfx9 stock) and warp_tile_k=16 (gfx12 WMMA) both zero out. + """ return ABQuantKernelConfig( variant_key="fp8", layout="rcr", @@ -888,7 +899,7 @@ def default_fp8_compv3_config_gfx1250( scheduler="intrawave", tile_m=128, tile_n=128, tile_k=128, warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, aquant_group_m=1, aquant_group_n=1, aquant_group_k=quant_group_k, bquant_group_m=1, bquant_group_n=bquant_group_n, bquant_group_k=quant_group_k, preshuffle_b=False, preshuffle_aq=False, preshuffle_bq=False, @@ -902,7 +913,10 @@ def default_bf8_compv3_config_gfx1250( bquant_group_n: int = 1, gfx_arch: str = _GFX1250_ARCH, ) -> ABQuantKernelConfig: - """bf8 ABQuant CompV3 config for gfx1250 (WMMA warp_tile_k=16, TransposeC=False).""" + """bf8 ABQuant CompV3 config for gfx1250 (FlatMM warp_tile_k=128, TransposeC=False). + + GPU-verified on MI400/gfx1250: max_rel_err ~3e-4 vs. fp32 dequant reference. + """ return ABQuantKernelConfig( variant_key="bf8", layout="rcr", @@ -911,7 +925,7 @@ def default_bf8_compv3_config_gfx1250( scheduler="intrawave", tile_m=128, tile_n=128, tile_k=128, warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, aquant_group_m=1, aquant_group_n=1, aquant_group_k=quant_group_k, bquant_group_m=1, bquant_group_n=bquant_group_n, bquant_group_k=quant_group_k, preshuffle_b=False, preshuffle_aq=False, preshuffle_bq=False, diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py index 5eda55da6e08..8ae843cc0c43 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py @@ -788,14 +788,26 @@ def default_bf8_preshuffleaq_config( # ============================================================================= -# gfx1250 (MI400 / RDNA WMMA) default configs +# gfx1250 (MI400) default configs # ============================================================================= -# On gfx1250 the quant MFMA path is unavailable; the codegen selects the gfx12 -# WMMA instruction when warp_tile_k=16 (warp_tile_k=32 selects the gfx9-only -# MFMA path, which silently returns zeros on gfx12). These helpers mirror the -# gfx9 defaults above but pin warp_tile_m=warp_tile_n=warp_tile_k=16 (WMMA). -# The compile path additionally injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 -# for gfx12 archs, so fp8/bf8 use the OCP encoding on gfx1250. +# Empirically (on-GPU, MI400/gfx1250, HIP_VISIBLE_DEVICES=0, fp32-dequant +# reference), the AQuant fp8/bf8 kernels are correct on gfx1250 only with +# warp_tile_k=128 (the FlatMM path, same tile the gfx9 fp8/bf8 FlatMM decode +# uses). The gfx9 stock AQuant default uses warp_tile_k=32 (plain MFMA), which +# produces a WRONG result on gfx1250 (max_rel_err=1.0); warp_tile_k=16 (gfx12 +# WMMA) returns all-zeros. So the only correct gfx1250 override is 16x16x128. +# +# fp8: warp_tile 16x16x128 -> nonzero, max_rel_err ~3e-4 (PASS) +# bf8: warp_tile 16x16x128 -> nonzero, max_rel_err ~1e-4 (PASS) +# +# The compile path injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for gfx12 +# archs, so fp8/bf8 use the OCP encoding on gfx1250. +# +# fp8i4 / bf8i4 (pk_int4 weights) do NOT compile on gfx1250 at any warp_tile_k +# (no gfx12 WMMA/FlatMM instruction for the packed-int4 quant path). Those +# variants are therefore not enabled on gfx1250: the C++ ctypes arch-gate is +# opened, but on-GPU correctness is deferred pending kernel work. No gfx1250 +# config helper is provided for i4 because it would emit a non-buildable kernel. _GFX1250_ARCH = "gfx1250" @@ -805,7 +817,11 @@ def default_fp8_config_gfx1250( quant_group_m: int = 1, gfx_arch: str = _GFX1250_ARCH, ) -> AQuantKernelConfig: - """fp8 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" + """fp8 AQuant decode config for gfx1250 (FlatMM, warp_tile_k=128). + + GPU-verified on MI400/gfx1250: max_rel_err ~3e-4 vs. fp32 dequant reference. + warp_tile_k=32 (gfx9 stock) is WRONG on gfx1250; warp_tile_k=16 zeros out. + """ return AQuantKernelConfig( variant_key="fp8", layout="rcr", @@ -814,7 +830,7 @@ def default_fp8_config_gfx1250( scheduler="intrawave", tile_m=16, tile_n=64, tile_k=256, warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, quant_group_m=quant_group_m, quant_group_n=1, quant_group_k=quant_group_k, @@ -828,62 +844,19 @@ def default_bf8_config_gfx1250( quant_group_m: int = 1, gfx_arch: str = _GFX1250_ARCH, ) -> AQuantKernelConfig: - """bf8 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" - return AQuantKernelConfig( - variant_key="bf8", - layout="rcr", - pipeline="mem", - epilogue="cshuffle", - scheduler="intrawave", - tile_m=16, tile_n=64, tile_k=256, - warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, - quant_group_m=quant_group_m, - quant_group_n=1, - quant_group_k=quant_group_k, - preshuffle_aq=False, - gfx_arch=gfx_arch, - ) - - -def default_fp8i4_config_gfx1250( - quant_group_k: int = 128, - quant_group_m: int = 1, - gfx_arch: str = _GFX1250_ARCH, -) -> AQuantKernelConfig: - """fp8i4 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" - return AQuantKernelConfig( - variant_key="fp8i4", - layout="rcr", - pipeline="mem", - epilogue="cshuffle", - scheduler="intrawave", - tile_m=16, tile_n=64, tile_k=256, - warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, - quant_group_m=quant_group_m, - quant_group_n=1, - quant_group_k=quant_group_k, - preshuffle_aq=False, - gfx_arch=gfx_arch, - ) - + """bf8 AQuant decode config for gfx1250 (FlatMM, warp_tile_k=128). -def default_bf8i4_config_gfx1250( - quant_group_k: int = 128, - quant_group_m: int = 1, - gfx_arch: str = _GFX1250_ARCH, -) -> AQuantKernelConfig: - """bf8i4 AQuant decode config for gfx1250 (WMMA, warp_tile_k=16).""" + GPU-verified on MI400/gfx1250: max_rel_err ~1e-4 vs. fp32 dequant reference. + """ return AQuantKernelConfig( - variant_key="bf8i4", + variant_key="bf8", layout="rcr", pipeline="mem", epilogue="cshuffle", scheduler="intrawave", tile_m=16, tile_n=64, tile_k=256, warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, quant_group_m=quant_group_m, quant_group_n=1, quant_group_k=quant_group_k, diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py index a8582338f193..6f68a5b37573 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py @@ -1217,14 +1217,26 @@ def default_mx_bf16fp4_config( # ============================================================================= -# gfx1250 (MI400 / RDNA WMMA) default configs +# gfx1250 (MI400) default configs # ============================================================================= -# On gfx1250 the quant path uses the gfx12 WMMA instruction, selected by the -# codegen when warp_tile_k=16 (warp_tile_k=32/128 select the gfx9-only MFMA/FlatMM -# paths, which silently return zeros on gfx12). These helpers mirror the standard -# gfx9 compv3 defaults but pin warp_tile_m=warp_tile_n=warp_tile_k=16 (WMMA). The -# compile path additionally injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for -# gfx12 archs, so fp8/bf8 use the OCP encoding on gfx1250. +# Empirically (on-GPU, MI400/gfx1250, HIP_VISIBLE_DEVICES=0, via the +# test_bquant_gpu_correctness fp32-dequant reference), the BQuant fp8/bf8 kernels +# are correct on gfx1250 with the SAME config that gfx9 uses: the CompV3 FlatMM +# tile with warp_tile_k=128. warp_tile_k=16 (gfx12 WMMA) returns all-zeros on +# gfx1250. So the gfx1250 helpers simply return the stock fp8/bf8 config +# (warp_tile_k=128); no distinct gfx1250 tile is required for BQuant. +# +# fp8: warp_tile 16x16x128 -> nonzero, max_rel_err ~5e-4 (PASS) +# bf8: warp_tile 16x16x128 -> nonzero, max_rel_err ~5e-4 (PASS) +# +# The compile path injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for gfx12 +# archs, so fp8/bf8 use the OCP encoding on gfx1250. +# +# fp8i4 / bf8i4 (pk_int4 weights) do NOT compile on gfx1250 at any warp_tile_k +# (no gfx12 instruction for the packed-int4 quant path). Those variants are not +# enabled on gfx1250: the C++ ctypes arch-gate is opened, but on-GPU correctness +# is deferred pending kernel work. No gfx1250 config helper is provided for i4 +# because it would emit a non-buildable kernel. _GFX1250_ARCH = "gfx1250" @@ -1234,19 +1246,14 @@ def default_fp8_config_gfx1250( quant_group_n: int = 1, gfx_arch: str = _GFX1250_ARCH, ) -> BQuantKernelConfig: - """fp8 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" - return BQuantKernelConfig( - variant_key="fp8", - layout="rcr", - pipeline="compv3", - epilogue="cshuffle", - scheduler="intrawave", - tile_m=16, tile_n=64, tile_k=256, - warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, - quant_group_m=1, - quant_group_n=quant_group_n, + """fp8 BQuant config for gfx1250 (compv3, FlatMM warp_tile_k=128). + + Identical to the stock fp8 config; GPU-verified on MI400/gfx1250 with + max_rel_err ~5e-4 vs. fp32 dequant reference. warp_tile_k=16 zeros out. + """ + return default_fp8_config( quant_group_k=quant_group_k, + quant_group_n=quant_group_n, gfx_arch=gfx_arch, ) @@ -1256,62 +1263,13 @@ def default_bf8_config_gfx1250( quant_group_n: int = 1, gfx_arch: str = _GFX1250_ARCH, ) -> BQuantKernelConfig: - """bf8 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" - return BQuantKernelConfig( - variant_key="bf8", - layout="rcr", - pipeline="compv3", - epilogue="cshuffle", - scheduler="intrawave", - tile_m=16, tile_n=64, tile_k=256, - warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, - quant_group_m=1, - quant_group_n=quant_group_n, - quant_group_k=quant_group_k, - gfx_arch=gfx_arch, - ) - + """bf8 BQuant config for gfx1250 (compv3, FlatMM warp_tile_k=128). -def default_fp8i4_config_gfx1250( - quant_group_k: int = 128, - quant_group_n: int = 1, - gfx_arch: str = _GFX1250_ARCH, -) -> BQuantKernelConfig: - """fp8i4 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" - return BQuantKernelConfig( - variant_key="fp8i4", - layout="rcr", - pipeline="compv3", - epilogue="cshuffle", - scheduler="intrawave", - tile_m=16, tile_n=64, tile_k=256, - warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, - quant_group_m=1, - quant_group_n=quant_group_n, + Identical to the stock bf8 config; GPU-verified on MI400/gfx1250 with + max_rel_err ~5e-4 vs. fp32 dequant reference. + """ + return default_bf8_config( quant_group_k=quant_group_k, - gfx_arch=gfx_arch, - ) - - -def default_bf8i4_config_gfx1250( - quant_group_k: int = 128, - quant_group_n: int = 1, - gfx_arch: str = _GFX1250_ARCH, -) -> BQuantKernelConfig: - """bf8i4 BQuant config for gfx1250 (compv3, WMMA warp_tile_k=16).""" - return BQuantKernelConfig( - variant_key="bf8i4", - layout="rcr", - pipeline="compv3", - epilogue="cshuffle", - scheduler="intrawave", - tile_m=16, tile_n=64, tile_k=256, - warp_m=1, warp_n=4, warp_k=1, - warp_tile_m=16, warp_tile_n=16, warp_tile_k=16, - quant_group_m=1, quant_group_n=quant_group_n, - quant_group_k=quant_group_k, gfx_arch=gfx_arch, ) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py index 2ca1f074c739..86f67b087b04 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py @@ -375,6 +375,11 @@ def test_transpose_c_in_name_only_when_true(self): class TestGfx1250Configs: + # gfx1250 correct config was determined empirically on MI400/gfx1250: + # ABQuant CompV3 fp8/bf8 verify only with warp_tile_k=128 (FlatMM tile); + # warp_tile_k=32 (gfx9 MFMA) and warp_tile_k=16 (WMMA) both produce all-zeros + # on gfx12. Still uses the standard CompV3 pipeline with transpose_c=False + # (NOT the gfx950-native eightwaves/preshuffleb path). def _all(self): return [ @@ -382,13 +387,13 @@ def _all(self): default_bf8_compv3_config_gfx1250(), ] - def test_gfx1250_uses_wmma_warp_tile_k_16(self): - # gfx1250 must use the standard CompV3 pipeline with the WMMA instruction - # (warp_tile_k=16), NOT the gfx950-native eightwaves/FlatMM path. + def test_gfx1250_uses_flatmm_warp_tile_k_128(self): + # GPU-verified: CompV3 fp8/bf8 are correct on gfx1250 only with + # warp_tile_k=128; warp_tile_k=32 and warp_tile_k=16 both zero out. for cfg in self._all(): assert cfg.warp_tile_m == 16 assert cfg.warp_tile_n == 16 - assert cfg.warp_tile_k == 16, f"{cfg.name} must use WMMA warp_tile_k=16" + assert cfg.warp_tile_k == 128, f"{cfg.name} must use FlatMM warp_tile_k=128" def test_gfx1250_uses_compv3_not_eightwaves(self): # eightwaves / preshuffleb are gfx950-only; gfx1250 uses standard compv3. @@ -401,9 +406,9 @@ def test_gfx1250_arch_propagated(self): for cfg in self._all(): assert cfg.gfx_arch == "gfx1250" - def test_gfx1250_names_16x16x16(self): + def test_gfx1250_names_16x16x128(self): for cfg in self._all(): - assert "16x16x16" in cfg.name + assert "16x16x128" in cfg.name def test_gfx1250_unique_names(self): names = [cfg.name for cfg in self._all()] diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py index 89b80fec9dce..9f9da689a226 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py @@ -418,36 +418,35 @@ def test_mem_and_compv3_pipelines_produce_different_names(self): from grouped_gemm_aquant_utils import ( # noqa: E402 default_fp8_config_gfx1250, default_bf8_config_gfx1250, - default_fp8i4_config_gfx1250, - default_bf8i4_config_gfx1250, ) class TestGfx1250Configs: + # gfx1250 correct config was determined empirically on MI400/gfx1250: + # fp8/bf8 verify only with warp_tile_k=128 (FlatMM); warp_tile_k=16 (WMMA) + # and warp_tile_k=32 (gfx9 MFMA) both produce wrong/zero output on gfx12. + # i4 variants do not compile on gfx1250, so no gfx1250 helper exists for them. def _all(self): return [ default_fp8_config_gfx1250(), default_bf8_config_gfx1250(), - default_fp8i4_config_gfx1250(), - default_bf8i4_config_gfx1250(), ] - def test_gfx1250_uses_wmma_warp_tile_k_16(self): - # gfx1250 must use the WMMA instruction (warp_tile_k=16); warp_tile_k=32 - # selects the gfx9 MFMA path which silently returns zeros on gfx12. + def test_gfx1250_uses_flatmm_warp_tile_k_128(self): + # GPU-verified: fp8/bf8 are correct on gfx1250 only with warp_tile_k=128. for cfg in self._all(): assert cfg.warp_tile_m == 16 assert cfg.warp_tile_n == 16 - assert cfg.warp_tile_k == 16, f"{cfg.name} must use WMMA warp_tile_k=16" + assert cfg.warp_tile_k == 128, f"{cfg.name} must use FlatMM warp_tile_k=128" def test_gfx1250_arch_propagated(self): for cfg in self._all(): assert cfg.gfx_arch == "gfx1250" - def test_gfx1250_names_16x16x16(self): + def test_gfx1250_names_16x16x128(self): for cfg in self._all(): - assert "16x16x16" in cfg.name + assert "16x16x128" in cfg.name def test_gfx1250_pipeline_is_mem(self): for cfg in self._all(): diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py index dc63b1e3c4cc..1f04ea74dcb1 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py @@ -734,36 +734,36 @@ def test_name_matches_kernel_config_name(self, tmp_path): from grouped_gemm_bquant_utils import ( # noqa: E402 default_fp8_config_gfx1250, default_bf8_config_gfx1250, - default_fp8i4_config_gfx1250, - default_bf8i4_config_gfx1250, ) class TestGfx1250Configs: + # gfx1250 correct config was determined empirically on MI400/gfx1250: + # BQuant fp8/bf8 verify with the SAME stock config (warp_tile_k=128, FlatMM); + # warp_tile_k=16 (WMMA) produces all-zeros on gfx12. i4 variants do not + # compile on gfx1250, so no gfx1250 helper exists for them. def _all(self): return [ default_fp8_config_gfx1250(), default_bf8_config_gfx1250(), - default_fp8i4_config_gfx1250(), - default_bf8i4_config_gfx1250(), ] - def test_gfx1250_uses_wmma_warp_tile_k_16(self): - # gfx1250 must use the WMMA instruction (warp_tile_k=16); the gfx9 fp8/bf8 - # defaults use warp_tile_k=128 (FlatMM) which is invalid on gfx12. + def test_gfx1250_uses_flatmm_warp_tile_k_128(self): + # GPU-verified: fp8/bf8 are correct on gfx1250 with warp_tile_k=128 + # (== stock). warp_tile_k=16 silently returns zeros on gfx12. for cfg in self._all(): assert cfg.warp_tile_m == 16 assert cfg.warp_tile_n == 16 - assert cfg.warp_tile_k == 16, f"{cfg.name} must use WMMA warp_tile_k=16" + assert cfg.warp_tile_k == 128, f"{cfg.name} must use FlatMM warp_tile_k=128" def test_gfx1250_arch_propagated(self): for cfg in self._all(): assert cfg.gfx_arch == "gfx1250" - def test_gfx1250_names_16x16x16(self): + def test_gfx1250_names_16x16x128(self): for cfg in self._all(): - assert "16x16x16" in cfg.name + assert "16x16x128" in cfg.name def test_gfx1250_pipeline_is_compv3(self): for cfg in self._all(): From 08466e476e33188195b32db825bb7eee2c7d90b2 Mon Sep 17 00:00:00 2001 From: Muhammed Emin Ozturk Date: Mon, 31 Aug 2026 17:46:06 -0400 Subject: [PATCH 5/6] fix(ck-tile): remove bad-merge duplicate _populate_launch_bquant A develop merge reintroduced a second _populate_launch_bquant definition that passes TransposeC (bool) into the GemmBQuantPipelineProblem ComputeDataType_ (type) slot, so grouped bquant failed to build on every arch (0 kernels). Remove the broken duplicate; the correct definition (compute slot = ADataType) is retained. gfx1250: 383 kernels build, 1747 verify pass. --- .../ops/gemm/gemm_instance_builder.py | 90 ------------------- 1 file changed, 90 deletions(-) diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py b/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py index 1026990fdf2a..8e1414ae7ae8 100644 --- a/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py @@ -1799,96 +1799,6 @@ def populate_cshuffle_gemm_aquant(self): using GemmEpilogue = ck_tile::CShuffleEpilogue;""" return instance_code - def _populate_launch_bquant( - self, - scheduler_type_map, - scheduler, - pipeline_impl_map, - pipeline, - epilogue, - k_block_per_cu, - ): - instance_code = """ - - // Launch function - static float launch(const ck_tile::QuantGemmHostArgs& args, const ck_tile::stream_config& stream) { - - // Hot loop detection - const ck_tile::index_t K_split = ck_tile::integer_least_multiple(args.K, TileShape::kK); - const ck_tile::index_t num_loop = TilePartitioner::GetLoopNum(K_split); - const bool has_hot_loop = BaseGemmPipeline::BlockHasHotloop(num_loop); - const ck_tile::TailNumber tail_num = BaseGemmPipeline::GetBlockLoopTailNum(num_loop);""" - - instance_code += f""" - - const auto Run = [&](const auto has_hot_loop_, const auto tail_number_) {{ - constexpr bool has_hot_loop_v = has_hot_loop_.value; - constexpr auto tail_number_v = tail_number_.value; - - // Full pipeline problem with hot loop and tail number - using PipelineProblem = ck_tile::GemmBQuantPipelineProblem< - ADataType, - BDataType, - BQDataType, - AccDataType, - TileShape, - Traits, - QuantGroupSize, - TransposeC, - BDataType, - {scheduler_type_map.get(scheduler)}, - has_hot_loop_v, - tail_number_v>; - - using GemmPipeline = {pipeline_impl_map.get(pipeline)};""" - - instance_code += """ - - // Epilogue""" - if epilogue == "cshuffle": - instance_code += self.populate_cshuffle_gemm_bquant() - else: - instance_code += self.populate_default_gemm_bquant() - - instance_code += f""" - - // Kernel type - using Kernel = ck_tile::QuantGemmKernel< - TilePartitioner, GemmPipeline, GemmEpilogue, - ck_tile::QuantType::BQuantGrouped>; - - // Kernel arguments - auto kargs = Kernel::MakeKernelArgs(args); - - if (!Kernel::IsSupportedArgument(kargs)) {{ - throw std::runtime_error("Unsupported kernel arguments; skipping GEMM launch."); - }} - - // Get grid and block sizes - const dim3 grids = Kernel::GridSize(args.M, args.N, args.k_batch); - const dim3 blocks = Kernel::BlockSize(); - - if(stream.log_level_ > 0) {{ - std::cout << "Launching kernel: " << Kernel::GetName() << '\\n' - << "grid: {{" << grids.x << ", " << grids.y << ", " << grids.z << "}}" - << ", blocks: {{" << blocks.x << ", " << blocks.y << ", " << blocks.z << "}}" - << std::endl; - }} - - // Launch kernel - constexpr int kBlockPerCu = {k_block_per_cu}; - float ave_time = ck_tile::launch_kernel( - stream, - ck_tile::make_kernel(Kernel{{}}, grids, blocks, 0, kargs)); - - return ave_time; - }}; - return BaseGemmPipeline::TailHandler(Run, has_hot_loop, tail_num); - }} -}}; -""" - return instance_code - def populate_default_gemm_bquant(self): instance_code = """ using EpilogueProblem = ck_tile::DefaultGemm2DEpilogueProblem< From a5aa436d2abd6d7e1d250ed2101986b0a98ebaed Mon Sep 17 00:00:00 2001 From: Muhammed Emin Ozturk Date: Tue, 1 Sep 2026 05:58:25 +0000 Subject: [PATCH 6/6] fix(ck-tile): remove duplicate aquant/bquant trait block in populate_trait_config (merge artifact using undefined persistent_or_preshuffle_quant broke abquant/bquant codegen on gfx1250) --- .../tile_engine/ops/gemm/gemm_instance_builder.py | 14 -------------- 1 file changed, 14 deletions(-) diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py b/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py index 8e1414ae7ae8..90a119f90e64 100644 --- a/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py @@ -780,20 +780,6 @@ def populate_trait_config(self, trait_combo): instance_code += """ static constexpr bool Preshuffle = false;""" - if self.kernel_name_prefix == "gemm_aquant": - instance_code += f""" - static constexpr bool APreshuffleQuant = {"true" if persistent_or_preshuffle_quant in [True, "true"] else "false"}; - static constexpr bool BPreshuffleQuant = false; - static constexpr bool PreshuffleB = false; - static constexpr ck_tile::index_t GroupSizeK = {self.group_size_k};""" - - elif self.kernel_name_prefix == "gemm_bquant": - instance_code += f""" - static constexpr bool APreshuffleQuant = false; - static constexpr bool BPreshuffleQuant = {"true" if persistent_or_preshuffle_quant in [True, "true"] else "false"}; - static constexpr bool PreshuffleB = false; - static constexpr ck_tile::index_t GroupSizeK = {self.group_size_k};""" - return instance_code def populate_initialization(self, base_pipeline_map, pipeline):