Skip to content

Commit 6f6a4cf

Browse files
[rocm-libraries] ROCm/rocm-libraries#10897 (commit 7080fbe)
feat(ck-tile): add dispatcher correctness and performance CI lanes for gfx942/gfx950 (#10897) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation ISSUE ID: 10897 A change to the CK-Tile dispatcher bridge or codegen can break numeric correctness, kernel build, or kernel launch for any of the 13 bridged GEMM operators, and today nothing catches it until someone runs the tests by hand. Three operators (stream-K, aquant, abquant) had no GPU correctness test at all, and four more were registered in ctest but never invoked from CI. This adds two opt-in Jenkins lanes that fan out across gfx942 and gfx950 and cover every operator with a bridge on `develop`, plus the missing tests. Both parameters default to **OFF**, so no existing CI run gets slower. ## Dependencies **Depends on #11591** — `fix(ck-tile): guard the gfx1250 unscaled WMMA builtin behind __gfx1250__`. `ck_tile/ops/gemm.hpp` does not compile on any non-gfx1250 target on current `develop` (19 errors, in both the host and the device pass), which breaks four of the dispatcher GPU correctness tests this PR adds. That fix is unrelated to this change and is split out into its own PR. Default CI here is unaffected — both lanes are `defaultValue: false` — but Test Plan steps 1, 2 and 4 will not pass until #11591 merges. ## Technical Details ### 1. `RUN_DISPATCHER_CORRECTNESS_TESTS` — JIT-compiled kernels vs. host reference Covers all 13 operators rather than a representative sample: a developer touching the bridge needs signal on all of them, and a lane covering four cannot give that. Cost is held down by the *search budget* instead — `--budget 64` for the standard sweep, `--budget 16` per variant sweep. | | gfx942 | gfx950 | |---|---|---| | shared | `gemm_universal` sweep + parity, `grouped`, `multi_d` (4 layouts), `multi_abd`, `batched_gemm`, `batched_contraction` (fp16+bf16+fp32), `stream_k`, `rowcolquant`, `tensorquant`, `aquant`, `abquant` | same | | arch-only | — | `mx_gemm`, `bquant` C4 (fp8/bf8), C (fp8i4/bf8i4), H3 (`mx_*`) | Only two things stay gfx950-only, both for ISA reasons rather than policy. `mx_gemm` needs the scale-MFMA builtins; it also cannot join the `--variant` sweep at all, because `mx_gemm_utils` exposes only `default_fp8_config()` / `default_fp4_config()` and never plugs into `expand_sweep`. The `bquant` H3 (`mx_*`) configs need the same builtins. `bquant`'s C4/C paths are held to gfx950 for a softer reason: its config builders are arch-aware now, so the gate is pending one validating gfx942 run rather than a known defect. Everything else — including stream-K and all three quant operators — was widened to run on gfx942 as well over the course of this branch. Both rationales are recorded inline in `ck.groovy`. There is deliberately **no gfx1201 stage**: `gemm_utils._SUPPORTED_ARCHES` excludes it, and widening that guard does not help, because every warp tile in `default_ci_config.json` is 64 wide while RDNA4 is warp_size 32, so all 16 strata enumerate 0 configs and the runner exits 1 on "no configs to run". One narrowed surface is documented rather than closed: `grouped_gemm_{a,ab,b}quant` run only on `rcr`, because all three config builders hardcode `layout="rcr"` and the ctypes libs assert packed rcr strides. ### 2. `RUN_DISPATCHER_PERF_TESTS` — ahead-of-time tile_engine instances Smoke-tier benchmark covering the **same 13 operators** as the correctness lane, so the AOT instance path gets the same operator coverage the JIT path gets. The two are gated separately and neither substitutes for the other: correctness JIT-builds each config with hipcc and checks it against a host reference, while this lane exercises tile_engine's codegen, instance builder and benchmark harness. One dtype, one layout, one problem size, and a fixed **64 AOT instances per operator**. | | gfx942 | gfx950 | |---|---|---| | fp16 / `rcr` | `gemm_universal`, `batched_gemm`, `batched_contraction`, `gemm_streamk`, `grouped_gemm` | `gemm_streamk` | | fp16 / `rcrr` | `gemm_multi_d`, `gemm_multi_abd` | — | | fp8 / `rcr` | `grouped_gemm_rowcolquant`, `grouped_gemm_tensorquant` | `gemm_aquant`, `gemm_bquant`, `gemm_abquant`, `mx_gemm` | Three details that are load-bearing rather than incidental: - **The per-op cap is explicit, not derived.** `TILE_ENGINE_SAMPLING_TIER` is a *total* budget split evenly across every op with a non-empty `<OP>_DATATYPE` (`ops/gemm/CMakeLists.txt:16-23`), so adding operators silently shrinks the coverage of the ones already there — gfx942 alone would have gone from 125 to 55 instances per op. Pinning `<OP>_MAX_INSTANCES` uses the first-class override path (same file, `:50-59`) and keeps lane cost fixed and independent of how many ops the lane builds. The tier is still passed, because each op also forwards it to the sampler as `--tier` to select the sampling strategy. - **One dtype and one layout per op is not just smoke sizing.** The per-op budget is divided again across `(dtype × layout)` combos with integer truncation, so a multi-combo op can floor a combo to zero instances and build nothing. - **`gemm_multi_d` and `gemm_multi_abd` take 4-character layouts** (A, B, D, E). Passing the `rcr` the other ops use is a parse error in the instance builder, not a narrower selection. `mx_gemm` cannot leave the gfx950 branch: its CMakeLists filters targets with a literal regex, `target MATCHES "^gfx950"`, so on any other arch `benchmark_mx_gemm_all` is never created and ninja fails on an unknown target rather than skipping. It also has no fp16 path. `gemm_streamk` builds on both arches. `DESIRED_TARGETS` was widened to include gfx950 earlier on this branch, but that CMake change is the one part that could not be exercised locally; this lane is its first real confirmation, and it matches the correctness lane, which already runs stream-K on both. If the gfx950 AOT build turns out to be broken, that `DESIRED_TARGETS` line is revertible on its own without touching the lane. ### 3. New search-space runner — `dispatcher/tests/test_gemm_search_space.py` One runner covering every variant `expand_sweep` understands (standard, grouped, multi_d, multi_abd, stream_k), selected with `--variant` — they share the same enumerate/sample/build/report machinery and differ only in which `Gpu*Runner` to drive and how the numpy reference is computed. It enumerates the `GemmKernelConfig` space against tile_engine's `default_ci_config.json`, samples a budget-limited subset with a daily rotating seed **per (dtype, layout) stratum** — so even budget 64 spreads ~4 configs across all 16 combinations rather than testing one deeply — compiles via the bridge, runs on GPU, and reports correctness + TFLOPS as JSON. This is the dispatcher equivalent of tile_engine's per-op benchmark scripts, with one deliberate difference: those never verify numerics and always exit 0. This one validates every kernel against a numpy reference and **exits 1 on any mismatch or build failure**. ### 4. New GPU correctness tests - **`test_streamk_gpu_correctness.py`** — all three reduction strategies. The shared sweep only ever builds `atomic` (`default_ci_config.json` carries no `streamk_config`, so `expand_sweep` falls back to `["atomic"]`), leaving `linear`/`tree` unverified through the ctypes bridge. Two shapes: an occupancy-bound 4096³ case, and a Stream-K-heavy `M=N=128, K=8192` case that only linear/tree can assert tightly, since `estimate_num_wgs_per_tile` is gated on Atomic and returns 1 otherwise. Tolerance is CK's own split-K-aware element-wise bound, ported from `streamk_driver_common.hpp` + `check_err.hpp`. - **`test_aquant_gpu_correctness.py`** — A-side per-(m,k)-group scales; fp8, bf8, and a tiled shape spanning a 4×4 output-tile grid that the 1×1-tile cases cannot catch. - **`test_abquant_gpu_correctness.py`** — scales on *both* operands, completing the quant trio; fp8 + bf8 on `compv3`, plus the gfx950-native `eightwaves` pipeline (192×256 tiles, `TransposeC=True`, `bquant_group_n=128`), a different MFMA path and B-scale granularity than compv3. All four quant/stream-K bridge tests are now registered in ctest with `SKIP_RETURN_CODE 77`. They intentionally take **no** `--gfx` from CMake: each hipcc-compiles with `--offload-arch` and then runs, so the arch must match the *physical* GPU, and the configured `GPU_TARGETS` may legitimately differ from the card in the box. ### 5. Skip semantics — exit 77 instead of exit 0 The script-style tests previously returned 0 when no GPU was present, so a CPU-only runner reported a green PASS for a test that never touched the GPU. They now return 77, which ctest maps to *skipped* via `SKIP_RETURN_CODE` and the Jenkins lane maps via a `run_ok` shell helper. Only 77 is swallowed — any other non-zero still fails the lane. This matters because `execute_cmd` is one `&&`-joined string: without it, a single unsupported arch would abort the lane and every later operator would silently never run. `--gfx` also loses its hardcoded `gfx950` default in the quant tests, so "user asked for gfx950" stays distinguishable from "we are on an unrelated box" and the skip can fire. ### 6. Fixes and coverage extensions found while wiring this up - `gemm_utils.expand_sweep`: skip non-`cshuffle` epilogues for `stream_k` — the codegen emits only cshuffle for that variant, so those configs surfaced as spurious build failures. - `grouped_gemm_aquant_utils._compile_aquant_kernel`: add the arch defines (`CK_USE_OCP_FP8`, `CK_TILE_USE_OCP_FP8`, `CK_USE_NATIVE_MX_SUPPORT`, `CK_GFX950_SUPPORT`) that CMake normally injects but the standalone hipcc path omitted, matching the bquant path. gfx942 is FNUZ, gfx950 is OCP; without this the encoded inputs and the compiled kernel disagree and the device returns NaN. - `tile_engine/ops/gemm_streamk/CMakeLists.txt`: widen AOT `DESIRED_TARGETS` to include gfx950, and correct the skip warning, which still listed only `(gfx90a, gfx942)` afterwards. - `streamk_gemm_ctypes_lib.cpp`: correct a stale comment claiming the lib is Atomic-only. The strategy is a codegen-time template constant and `GetWorkSpaceSize` returns 0 for Atomic / non-zero for Linear/Tree, so the one allocate-and-launch path serves all three. - Widened two existing GPU tests that were narrower than the operator they cover: batched-contraction now exercises bf16 and fp32 alongside fp16, and the Multi-D GEMM test exercises all four layouts rather than one. - `Dockerfile`: add `ml_dtypes`, required for host-side fp8/bf8 encoding. Both lanes archive their per-kernel JSON from a `finally` block — a red lane is exactly when that artifact is worth having, and `buildAndTest` throws on failure. ## Test Plan 1. **gfx942 correctness lane** — `RUN_DISPATCHER_CORRECTNESS_TESTS=true`: the `gemm_universal` sweep at budget 64 across all 16 dtype × layout strata, the parity test, the grouped/multi_d/multi_abd/batched_gemm/batched_contraction/rowcolquant/ tensorquant/aquant/abquant/bquant bridge tests, the `grouped`/`multi_d`/`multi_abd`/ `stream_k` variant sweeps at budget 16, and the stream-K registry + bridge tests. 2. **gfx950 correctness lane** — the same set, plus `mx_gemm` and the `bquant` H3 (`mx_*`) configs. 3. **Both perf lanes** — `RUN_DISPATCHER_PERF_TESTS=true` on gfx942 and gfx950; confirm all 9 gfx942 and all 5 gfx950 `benchmark_*_all` targets build, every script runs, and each `*_results.json` lands as an archived artifact. The gfx950 `gemm_streamk` result is the specific new signal — it is the first exercise of the `DESIRED_TARGETS` widening. 4. **Skip path** — run the ctest suite on a non-gfx942/gfx950 box and confirm the arch-gated tests report *Skipped*, not Passed and not Failed. 5. **Artifact archiving** — confirm `dispatcher_*_results.json` and `*_results.json` appear as build artifacts, including on a deliberately failed run. ## Test Result - Will be run on CI ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
1 parent 6f9eefd commit 6f6a4cf

49 files changed

Lines changed: 4736 additions & 1308 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

Dockerfile

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -100,7 +100,15 @@ RUN set -x && \
100100
wget https://github.com/Yelp/dumb-init/releases/download/v1.2.0/dumb-init_1.2.0_amd64.deb && \
101101
dpkg -i dumb-init_*.deb && rm dumb-init_*.deb && \
102102
# Install packages for processing the performance results
103-
pip3 install --break-system-packages --upgrade pytest pymysql pandas==2.2.3 sqlalchemy==2.0.3 setuptools-rust setuptools sshtunnel==0.4.0 && \
103+
# ml_dtypes is pinned like its neighbours, and for a sharper reason than most:
104+
# it supplies both float8_e4m3fn/float8_e5m2 (OCP) and float8_e4m3fnuz/
105+
# float8_e5m2fnuz (FNUZ), and the dispatcher tests pick between those pairs per
106+
# arch. An unpinned upgrade would silently move fp8 encoding semantics -- and
107+
# the reference values every parity/correctness test compares against -- from
108+
# under CI on an image rebuild. 0.6.0 is what this image already resolved on
109+
# Ubuntu 24.04 / Python 3.12, so this records current behaviour rather than
110+
# bumping it.
111+
pip3 install --break-system-packages --upgrade pytest pymysql pandas==2.2.3 sqlalchemy==2.0.3 setuptools-rust setuptools sshtunnel==0.4.0 ml_dtypes==0.6.0 && \
104112
# Add render group
105113
groupadd -f render && \
106114
# Install the new rocm-cmake version

Jenkinsfile

Lines changed: 164 additions & 416 deletions
Large diffs are not rendered by default.

dispatcher/python/batched_contraction_utils.py

Lines changed: 32 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -693,11 +693,40 @@ def expand_sweep(config: dict, dtype: str = "fp16", layout: str = "rcr") -> List
693693
return out
694694

695695

696-
def default_fp16_config(gfx_arch: Optional[str] = None) -> BatchedContractionKernelConfig:
696+
# Warp tile per dtype, chosen from the XDL allow-list in is_valid(). fp16 and
697+
# bf16 share the same MFMA shape; fp32 has no 32x32x16 instruction and needs
698+
# 32x32x8 instead. Picking a shape outside the allow-list is not an error you
699+
# would notice -- is_valid() just returns False and the config is silently
700+
# dropped from the sweep (test_batched_contraction_bridge.py asserts exactly
701+
# that), so a "default" config with the wrong tile would look like a dtype that
702+
# simply has no instances.
703+
_DEFAULT_WARP_TILE = {
704+
"fp16": (32, 32, 16),
705+
"bf16": (32, 32, 16),
706+
"fp32": (32, 32, 8),
707+
}
708+
709+
710+
def default_config(
711+
dtype: str = "fp16", gfx_arch: Optional[str] = None
712+
) -> BatchedContractionKernelConfig:
713+
"""A known-valid single config for ``dtype``, for smoke tests and defaults."""
714+
try:
715+
wt_m, wt_n, wt_k = _DEFAULT_WARP_TILE[dtype]
716+
except KeyError:
717+
raise ValueError(
718+
f"no default warp tile for dtype {dtype!r}; "
719+
f"known: {', '.join(sorted(_DEFAULT_WARP_TILE))}"
720+
) from None
697721
gfx_arch = _validate_arch(gfx_arch) if gfx_arch else _get_arch()
698722
return BatchedContractionKernelConfig(
699-
dtype="fp16", layout="rcr", pipeline="compv3", epilogue="cshuffle", scheduler="intrawave",
723+
dtype=dtype, layout="rcr", pipeline="compv3", epilogue="cshuffle", scheduler="intrawave",
700724
tile_m=128, tile_n=128, tile_k=64, warp_m=2, warp_n=2, warp_k=1,
701-
warp_tile_m=32, warp_tile_n=32, warp_tile_k=16,
725+
warp_tile_m=wt_m, warp_tile_n=wt_n, warp_tile_k=wt_k,
702726
num_dim_g=1, num_dim_m=1, num_dim_n=1, num_dim_k=1, gfx_arch=gfx_arch,
703727
)
728+
729+
730+
def default_fp16_config(gfx_arch: Optional[str] = None) -> BatchedContractionKernelConfig:
731+
"""Thin alias kept for existing callers; prefer :func:`default_config`."""
732+
return default_config("fp16", gfx_arch=gfx_arch)

dispatcher/python/dispatcher_common.py

Lines changed: 114 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,120 @@ def detect_gpu_arch(fallback: str = "gfx942") -> str:
101101
return fallback
102102

103103

104+
# ============================================================================
105+
# fp8 / bf8 encoding format per architecture
106+
# ============================================================================
107+
#
108+
# CK has two incompatible 8-bit float encodings and picks between them per arch:
109+
#
110+
# OCP (gfx950, gfx12): e4m3fn / e5m2 -- exponent bias 7 / 15
111+
# FNUZ (everything else, notably gfx942): e4m3fnuz / e5m2fnuz -- bias 8 / 16
112+
#
113+
# include/ck_tile/core/config.hpp:361-371 resolves this at compile time, but only
114+
# the *device* pass sees __gfx950__; the host pass of the very same header falls
115+
# back to FNUZ. So a host-side reference encoder that guesses from the header ends
116+
# up disagreeing with the kernel it is validating. Every caller must therefore
117+
# decide from the target arch string, and pass -DCK_TILE_USE_OCP_FP8 explicitly so
118+
# both compiler passes agree. These helpers are the single source of truth for
119+
# both halves of that contract.
120+
121+
122+
def normalize_arch(arch: Optional[str]) -> str:
123+
"""Lowercase `arch` and strip the target-feature suffix, if any.
124+
125+
rocminfo, hipcc and the HSA runtime all hand back full target triples --
126+
``gfx950:sramecc+:xnack-`` -- while configs, CI parameters and these helpers
127+
are written in terms of the bare name. Every arch predicate below matches on
128+
a prefix, so a triple happens to survive unnormalized, but anything carrying
129+
a leading component (an ``amdgcn-amd-amdhsa--gfx950`` offload target, say)
130+
silently falls through to the FNUZ default. That miss drops
131+
``-DCK_USE_OCP_FP8``, and host and device then disagree on the fp8 encoding
132+
with no diagnostic at all -- the kernel builds and returns wrong numbers.
133+
134+
Normalizing in one place, up front, is what keeps that from depending on
135+
which spelling of the arch a given caller happened to be handed.
136+
"""
137+
a = (arch or "").lower().strip()
138+
# Feature suffix: gfx950:sramecc+:xnack- -> gfx950
139+
a = a.split(":", 1)[0]
140+
# Offload-target prefix: amdgcn-amd-amdhsa--gfx950 -> gfx950
141+
idx = a.rfind("gfx")
142+
return a[idx:] if idx > 0 else a
143+
144+
145+
def fp8_uses_ocp(arch: Optional[str]) -> bool:
146+
"""True iff `arch` uses the OCP fp8/bf8 encoding rather than FNUZ.
147+
148+
Mirrors the __gfx950__ / __gfx12__ test in include/ck_tile/core/config.hpp.
149+
Use this to select the host-side codec (ml_dtypes.float8_e4m3fn vs
150+
float8_e4m3fnuz) so it matches the bytes the kernel actually produces.
151+
152+
Accepts bare names and full target triples alike; see `normalize_arch`.
153+
"""
154+
a = normalize_arch(arch)
155+
return a.startswith("gfx950") or a.startswith("gfx12")
156+
157+
158+
def ocp_arch_defines(arch: Optional[str]) -> List[str]:
159+
"""hipcc defines that pin the fp8/bf8 encoding for `arch`.
160+
161+
Returned for OCP archs only; FNUZ is the default for both compiler passes and
162+
needs no define. Passing these makes the host pass agree with the device pass
163+
instead of silently falling back to FNUZ.
164+
"""
165+
if not fp8_uses_ocp(arch):
166+
return []
167+
return ["-DCK_USE_OCP_FP8", "-DCK_TILE_USE_OCP_FP8"]
168+
169+
170+
def arch_feature_defines(arch: Optional[str]) -> List[str]:
171+
"""`ocp_arch_defines` plus the per-arch feature-enablement defines.
172+
173+
The top-level CMakeLists.txt (:456-512) sets these for a normal build; they
174+
are absent in the standalone hipcc JIT path, so any bridge that compiles a
175+
kernel out-of-tree has to re-supply them. This mirrors that block for the
176+
arches the dispatcher targets:
177+
178+
gfx950 -> CK_USE_NATIVE_MX_SUPPORT, CK_GFX950_SUPPORT
179+
gfx1250 -> CK_USE_GFX1250, CK_USE_NATIVE_MX_SUPPORT, CK_GFX1250_SUPPORT
180+
gfx11/gfx12 -> CK_TILE_USE_WMMA=1 (gfx12 also CK_GFX12_SUPPORT)
181+
182+
CK_TILE_USE_WMMA is the one that must be passed even when it is 0: CMake
183+
always defines it (:480), and the JIT path leaving it undefined only happens
184+
to work because the preprocessor reads an undefined identifier as 0, which is
185+
the right answer on gfx942/gfx950 and the wrong one on every WMMA part.
186+
187+
CK_USE_GFX950 is deliberately *not* emitted -- CMake sets it, but it is read
188+
only by the legacy ck/library conv instances, never by ck_tile, so it is
189+
dead weight on this path.
190+
191+
On the warp tile: these defines do select the branch of
192+
tile_gemm_shape.hpp get_k_warp_tile() (CK_TILE_USE_WMMA outermost, then
193+
CK_USE_GFX1250 / CK_GFX950_SUPPORT), but the emitted kernels never *call*
194+
that function -- codegen writes a literal warp_tile_k, mirroring it in Python
195+
via codegen_common.fp8_warp_tile_k_for_arch. The hazard is therefore a
196+
mismatch between that literal and the warp-gemm the target arch actually has:
197+
there is no 16x16x128 fp8 warp-gemm on gfx942, and asking for one compiles
198+
cleanly and returns all zeros. Keep warp_tile_k arch-aware alongside these
199+
defines, and derive it from the one helper rather than a local copy.
200+
"""
201+
a = normalize_arch(arch)
202+
defines = ocp_arch_defines(arch)
203+
if a.startswith("gfx950"):
204+
defines = defines + ["-DCK_USE_NATIVE_MX_SUPPORT", "-DCK_GFX950_SUPPORT"]
205+
elif a.startswith("gfx11") or a.startswith("gfx12"):
206+
defines = defines + ["-DCK_TILE_USE_WMMA=1"]
207+
if a.startswith("gfx12"):
208+
defines = defines + ["-DCK_GFX12_SUPPORT"]
209+
if a.startswith("gfx1250"):
210+
defines = defines + [
211+
"-DCK_USE_GFX1250",
212+
"-DCK_USE_NATIVE_MX_SUPPORT",
213+
"-DCK_GFX1250_SUPPORT",
214+
]
215+
return defines
216+
217+
104218
# ============================================================================
105219
# Architecture Filter Data
106220
# ============================================================================

0 commit comments

Comments
 (0)