Repository navigation
feat(ck-tile): Enabled the unification framework - #11646
Conversation
Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>
ThruptiRajLakshmanaGowda
left a comment
There was a problem hiding this comment.
Hey Folks, is this an urgent change? Can we wait till freeze to merge this PR?
|
@ThruptiRajLakshmanaGowda Definitely. We either way need to validate no performance regression is introduced by this. |
JP-Fernando
left a comment
There was a problem hiding this comment.
This PR enables the MMA-unified framework. The additional correctness fix for gfx12 WMMA is justified and documented. Approved without further comments from my side.
|
Verified that the performance of ck_tile fmha and ck_tile gemms is not affected by these changes. |
The merge freeze is over.
feat(ck-tile): Enabled the unification framework This PR should be used to do performance measurements and, compared to its baseline `develop`, validate that no performance regression is introduced by the Unification Framework. Closes: #11646
TheRock Submodule Bump ActivityNewest first
|
…el 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>
This PR should be used to do performance measurements and, compared to its baseline
develop, validate that no performance regression is introduced by the Unification Framework.Closes: #11646