[rocm-libraries] ROCm/rocm-libraries#12216 (commit 72144cb)

fix(ck-tile): honor the gfx1250 unified-framework gate in
 bridge kernel builds (#12216)

## Summary

`projects/composablekernel/CMakeLists.txt:230-232` force-defines
`USE_NEW_UNIFIED_FRAMEWORK=0` for gfx1250 targets, because the unified
ck_tile framework does not support gfx1250 yet. That gate is an
`add_compile_definitions` call, so it only reaches targets of that CMake
project.

The dispatcher bridges never saw it: they assemble their own per-kernel
`hipcc` command lines, outside CMake entirely.

This was harmless while the framework default was `0`. PR #11646 flipped
that default to `1`, after which every bridge kernel compiled for
gfx1250 failed on the `getCMakeCompilerTarget` static assertion in
`arch.hpp:509`, leaving `tile_engine` unable to build or benchmark any
GEMM on gfx1250.

This adds a shared `unified_framework_flags()` helper in
`dispatcher_common.py` and applies it at each of the 21 sites that build
a kernel compile line. The flag is emitted **only** for gfx1250.

Fixes the blocker reported in AICK-2268.

## Motivation

Without this, the documented gfx1250 tuning workflow does not function
at all: the shipped `configs/default_ci_config_gfx1250.json` sweep
reports `0/4 configs -> .so` and produces no measurements.

Only `gemm_utils.py` is exercised by that sweep, but all 21 compile-line
construction sites have the identical gap, so the remaining operators
would fail the same way as soon as anyone benchmarks them on gfx1250.
Fixing them together via one helper avoids leaving 17 latent copies of
the same bug.

## Design note

The helper mirrors the CMake gate rather than replacing it. A
longer-term cleanup would move the gate out of `CMakeLists.txt` into a
header, so that every compile path observes one consistent value instead
of two places having to agree. That is deliberately out of scope here.

Worth noting for reviewers: the dispatcher static library is built by a
separate CMake project (`project(ck_tile_dispatcher ...)` at
`dispatcher/CMakeLists.txt:6`), so it also does not inherit the gfx1250
gate and is compiled with the framework default of 1, while the kernels
this PR fixes are compiled with 0. That mix was verified to work
correctly on hardware -- the dispatcher translation units do not
instantiate the affected templates -- so no CMake change is required
here.

This restores the legacy WarpGemm path on gfx1250, which is what CMake
already intends today. Making the unified framework genuinely work on
gfx1250 is separate follow-up work.

## Changes

- `dispatcher_common.py`: new `unified_framework_flags(arch)` helper.
- 18 bridge modules: emit the helper's flags alongside `-DGFX_ARCH` at
each compile-line site.

19 files changed, 54 insertions, 2 deletions.

## Test plan

Verified on gfx1250 hardware against a pristine tree at the commit named
in AICK-2268, with this PR as the only local change, and with the
dispatcher built exactly as the ticket documents (no extra CMake flags).
Reproduced on a second gfx1250 node on a different ROCm build to rule
out version specificity.

- [x] Baseline reproduces the reported failure: `0/4 configs -> .so`
- [x] With this PR applied: `4/4 configs -> .so`, **0 compile failures**
- [x] Full bf16 sweep across all four layouts (rcr, rrr, crr, ccr): **64
configs, 0 compile failures, 320/320 measurements, 0 failures**
- [x] Numerical correctness via `--verify` against the fp32 reference:
**320/320 `verified=True`**, worst `max_rel` 0.003789, zero
all-zero-output rows
- [x] Kernels execute on device and emit native matrix instructions: 40
x `v_wmma_f32_16x16x32_bf16`, no fallback dummy-exec path
- [x] Non-gfx1250 unaffected: generated compile command for gfx942 and
gfx950 is byte-identical to before (27 args, no added flag); gfx1250
gets 28
- [x] All 21 dispatcher python modules import cleanly
- [ ] CI sweep on gfx942 / gfx950 to confirm no regression
- [ ] CI smoke check that builds at least one gfx1250 tile_engine
kernel, so this cannot silently regress again

Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>

Author: ozturkosu

dispatcher/python/batched_contraction_utils.py

@@ -22,6 +22,7 @@
 3-char a/b/e layout, num_dim_g/m/n/k, num_d_tensors == 0 (PassThrough).
 """
 
+from dispatcher_common import unified_framework_flags
 import concurrent.futures
 import ctypes
 import functools
@@ -557,7 +558,7 @@ def _compile_kernel(hpp: Path, so: Path, arch: str) -> bool:
         _HIPCC, "-c", "-fPIC", "-O3", "-std=c++17",
         f"-I{ck_root}/include", f"-I{ck_root}",
         "-DCK_TILE_SINGLE_KERNEL_INCLUDE", f"-include{hpp}",
-        "-D__HIP_PLATFORM_AMD__", f"--offload-arch={arch}", f'-DGFX_ARCH="{arch}"',
+        "-D__HIP_PLATFORM_AMD__", f"--offload-arch={arch}", f'-DGFX_ARCH="{arch}"', *unified_framework_flags(arch),
         # Match Tile Engine's AMDGPU codegen flags exactly so the bridge .so
         # produces the same machine code as Old-TE (inlining, register
         # allocation, occupancy).  Without these, persistent kernels size their

dispatcher/python/batched_gemm_utils.py

@@ -32,6 +32,7 @@
 """
 
 from __future__ import annotations
+from dispatcher_common import unified_framework_flags
 
 import ctypes
 import multiprocessing
@@ -604,6 +605,7 @@ def _build_batched_compile_jobs(
         "-D__HIP_PLATFORM_AMD__",
         f"--offload-arch={gfx_arch}",
         f'-DGFX_ARCH="{gfx_arch}"',
+        *unified_framework_flags(gfx_arch),
         # Byte-identical AMDGPU backend flags to the single-problem bridge and
         # Old-TE (see gemm_utils._tile_engine_codegen_flags) -- required for a
         # fair A/B parity comparison.

dispatcher/python/contraction_multi_abd_utils.py

@@ -26,6 +26,7 @@
   result = runner.run(As, Bs, Ds, problem)
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import json
 import logging
@@ -734,6 +735,7 @@ def _compile_kernel(
         "-DCK_TILE_SINGLE_KERNEL_INCLUDE", "-w",
         f"--offload-arch={gfx_arch}",
         f"-DGFX_ARCH=\"{gfx_arch}\"",
+        *unified_framework_flags(gfx_arch),
         *arch_defines,
         "-include", str(hpp_path),
         str(_CTYPES_LIB_SRC),

dispatcher/python/ctypes_utils.py

@@ -26,6 +26,7 @@
     check = validator.check(result.C, C_reference)
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import subprocess
 import numpy as np
@@ -2039,6 +2040,7 @@ def _rebuild_library_for_config(
             "-D__HIP_PLATFORM_AMD__",
             f"--offload-arch={config.gfx_arch}",
             f'-DGFX_ARCH="{config.gfx_arch}"',  # Pass arch as string for gemm_ctypes_lib.cpp
+            *unified_framework_flags(config.gfx_arch),
             "-mllvm",
             "-enable-noalias-to-md-conversion=0",
             "-Wno-undefined-func-template",
@@ -2131,6 +2133,7 @@ def build_libraries_parallel(
                 "-D__HIP_PLATFORM_AMD__",
                 f"--offload-arch={config.gfx_arch}",
                 f'-DGFX_ARCH="{config.gfx_arch}"',
+                *unified_framework_flags(config.gfx_arch),
                 "-mllvm",
                 "-enable-noalias-to-md-conversion=0",
                 "-Wno-undefined-func-template",
@@ -2891,6 +2894,7 @@ def setup_multiple_gemm_dispatchers(
             "-D__HIP_PLATFORM_AMD__",
             f"--offload-arch={c.gfx_arch}",
             f'-DGFX_ARCH="{c.gfx_arch}"',
+            *unified_framework_flags(c.gfx_arch),
             "-mllvm",
             "-enable-noalias-to-md-conversion=0",
             "-Wno-undefined-func-template",

dispatcher/python/dispatcher_common.py

@@ -101,6 +101,21 @@ def detect_gpu_arch(fallback: str = "gfx942") -> str:
     return fallback
 
 
+def unified_framework_flags(arch: Optional[str]) -> List[str]:
+    """Extra defines a per-kernel hipcc line needs to match CK's CMake gate.
+
+    ``projects/composablekernel/CMakeLists.txt`` force-defines
+    ``USE_NEW_UNIFIED_FRAMEWORK=0`` for gfx1250 targets, because the unified
+    ck_tile framework does not support gfx1250 yet. That gate is an
+    ``add_compile_definitions`` call, so it only reaches targets of that CMake
+    project. The bridges build their own hipcc command lines outside it and
+    would otherwise pick up the header default of 1, which does not compile.
+    """
+    if normalize_arch(arch) == "gfx1250":
+        return ["-DUSE_NEW_UNIFIED_FRAMEWORK=0"]
+    return []
+
+
 # ============================================================================
 # fp8 / bf8 encoding format per architecture
 # ============================================================================

dispatcher/python/fmha_utils.py

@@ -16,6 +16,7 @@
     result = runner.run(Q, K, V, problem)
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import json
 import os
@@ -1206,6 +1207,7 @@ def fmha_compile_flags(arch: str, hipcc: str = "", family: str = "") -> List[str
     - CK_USE_XDL: enables MFMA (matrix fused multiply-add) instructions
     - CK_TILE_USE_WMMA: 0 for CDNA (uses MFMA instead)
     - CK_TILE_FLOAT_TO_BFLOAT16_DEFAULT=3: BWD bf16 conversion mode
+    - USE_NEW_UNIFIED_FRAMEWORK=0: preserves the gfx1250 CMake gate for every TU
     """
     if not hipcc:
         hipcc = _find_hipcc()
@@ -1217,6 +1219,7 @@ def fmha_compile_flags(arch: str, hipcc: str = "", family: str = "") -> List[str
         "-O3",
         "-DNDEBUG",
         f"--offload-arch={arch}",
+        *unified_framework_flags(arch),
         "-std=c++17",
         f"-I{root.parent / 'include'}",
         f"-I{root / 'include'}",

dispatcher/python/gemm_abquant_utils.py

@@ -34,6 +34,7 @@
   result = runner.run(A, B, AQ, BQ, ABQuantGemmProblem(M=128, N=128, K=256))
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import json
 import logging
@@ -626,6 +627,7 @@ def _compile_abquant_kernel(
                    "-DCK_TILE_SINGLE_KERNEL_INCLUDE", "-w",
                    f"--offload-arch={gfx_arch}",
                    f"-DGFX_ARCH=\"{gfx_arch}\"",
+                   *unified_framework_flags(gfx_arch),
                    *arch_defines,
                    *perf_flags,
                    "-include", str(hpp_path),

dispatcher/python/gemm_aquant_utils.py

@@ -23,6 +23,7 @@
   result = runner.run(A, AQ, B, AQuantGemmProblem(M=16, N=64, K=256))
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import json
 import logging
@@ -510,6 +511,7 @@ def _compile_aquant_kernel(
                    "-DCK_TILE_SINGLE_KERNEL_INCLUDE", "-w",
                    f"--offload-arch={gfx_arch}",
                    f"-DGFX_ARCH=\"{gfx_arch}\"",
+                   *unified_framework_flags(gfx_arch),
                    *arch_defines,
                    *_te_perf_flags(hipcc),
                    "-include", str(hpp_path),

dispatcher/python/gemm_bquant_utils.py

@@ -25,6 +25,7 @@
   result = runner.run(A, B, BQ, BQuantGemmProblem(M=16, N=64, K=256))
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import json
 import functools
@@ -876,6 +877,7 @@ def _compile_bquant_kernel(
                    "-DCK_TILE_SINGLE_KERNEL_INCLUDE", "-w",
                    f"--offload-arch={gfx_arch}",
                    f"-DGFX_ARCH=\"{gfx_arch}\"",
+                   *unified_framework_flags(gfx_arch),
                    *arch_defines,
                    *codegen_flags,
                    "-include", str(hpp_path),

dispatcher/python/gemm_rowcolquant_utils.py

@@ -28,6 +28,7 @@
   result = runner.run(A, B, AQ, BQ, RowColQuantGemmProblem(M=16, N=64, K=256))
 """
 
+from dispatcher_common import unified_framework_flags
 import ctypes
 import json
 import logging
@@ -528,6 +529,7 @@ def _compile_rowcolquant_kernel(
                    "-DCK_TILE_SINGLE_KERNEL_INCLUDE", "-w",
                    f"--offload-arch={gfx_arch}",
                    f"-DGFX_ARCH=\"{gfx_arch}\"",
+                   *unified_framework_flags(gfx_arch),
                    *arch_defines,
                    *_te_perf_flags(hipcc),
                    "-include", str(hpp_path),