[P3-1] Tile-config table design (consteval get_tile_config) by chris-tsiaousis-hpc · Pull Request #7538 · ROCm/rocm-libraries

chris-tsiaousis-hpc · 2026-05-18T15:20:08Z

Workplan task ID: P3-1

Files: dqdkdv_spec.hpp; dqdkdv_dev.hpp

Description: Design a consteval tile config lookup: get_tile_config(hdim_q, hdim_v, dtype, arch) returning {block_size, block_n0, BlockTile, GemmBlockWarps, GemmWarpTile, occupancy}. Source data from fmha_bwd.py's get_dq_dk_dv_tiles(). Two options: (a) embed the 5 gfx9 configs as a consteval table in the spec, or (b) add tile config fields to FmhaBwdDQDKDVAlgorithm (27 fields — heavy). Recommend (a).

Example to follow: GEMM tile-config dispatch in include/ck_tile/ops/gemm/

CK Tile reference: example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py (source of truth)

CK Tile feature check (required): Before implementing this task, verify the feature is present in the CK Tile source tree in BOTH locations:

fmha_bwd example — especially codegen: projects/composablekernel/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py (and supporting host files: fmha_bwd.hpp, fmha_bwd_runner.hpp, example_fmha_bwd.cpp).
kernels: projects/composablekernel/include/ck_tile/ops/fmha/{kernel,pipeline,block}/.

If the feature is missing or only partially wired in either location, raise it as a blocker before writing the rocm_ck wrapper — the wrapper cannot expose what CK Tile does not implement.

Tests: experimental/rocm_ck/tests/test_fmha_bwd_consteval.cpp (config dispatch)

Blocked by: (none)

Closes: #7506

Copilot

Pull request overview

Note

Copilot was unable to run its full agentic suite in this review.

Replaces the hardcoded d128 tile geometry in the dQ/dK/dV backward bridge with a consteval lookup table that returns tile configs for all GFX9 fp16/bf16 head dimensions (32/64/96/128/256), and wires that table through makeSpec() and the device-side FmhaBwdDQDKDVTypes.

Changes:

Add FmhaBwdDQDKDVTileConfig struct + GFX9_FP16_DQDKDV_TILES table + consteval getTileConfig() lookup in dqdkdv_spec.hpp, and consume it from makeSpec() (block_size, bn0, occupancy).
Replace hardcoded sequences/asserts in dqdkdv_dev.hpp with kTile = getTileConfig(...), deriving BlockTile, all GemmNBlockWarps/GemmNWarpTile, GEMM4 wk = min(wk0, bk4), kMaxSeqLenQ, and IsWG32 from the lookup.
Add test_fmha_bwd_consteval.cpp covering each gfx9 tile entry, BF16 ≡ FP16, block_per_cu auto-resolution, and update existing dqdkdv tests to use hdim_q == hdim_v.

Reviewed changes

Copilot reviewed 5 out of 5 changed files in this pull request and generated 7 comments.

Show a summary per file

File	Description
projects/composablekernel/experimental/rocm_ck/include/rocm_ck/ops/fmha_bwd/dqdkdv_spec.hpp	Adds tile config struct, gfx9 table, consteval lookup; switches `makeSpec` to use it.
projects/composablekernel/experimental/rocm_ck/include/rocm_ck/ops/fmha_bwd/dqdkdv_dev.hpp	Derives all tile sequences and dropout/maxSeqQ from `getTileConfig`; relaxes d128-only static_assert.
projects/composablekernel/experimental/rocm_ck/tests/test_fmha_bwd_consteval.cpp	New tests for `getTileConfig` table, BF16 parity, and `makeSpec` integration.
projects/composablekernel/experimental/rocm_ck/tests/test_fmha_bwd_dqdkdv.cpp	Updates existing tests to use matching `hdim_q == hdim_v`.
projects/composablekernel/experimental/rocm_ck/tests/CMakeLists.txt	Registers the new `test_fmha_bwd_consteval` target.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

+    // --- tile geometry (from consteval lookup table) ---
+    // getTileConfig() returns the architecture-specific tile geometry for
+    // the given (hdim_q, hdim_v, dtype, target). Currently only GFX9 fp16/bf16
+    // configs are populated.
+    constexpr GpuTarget target = GpuTarget::gfx942;
+    auto tile                  = getTileConfig(sig.hdim_q, sig.hdim_v, sig.dtype, target);


+    constexpr GpuTarget target = GpuTarget::gfx942;
+    auto tile                  = getTileConfig(sig.hdim_q, sig.hdim_v, sig.dtype, target);


+consteval FmhaBwdDQDKDVTileConfig
+getTileConfig(int hdim_q, int hdim_v, DataType dtype, GpuTarget target)
+{


+TEST(TileConfig, GFX9_FP16_D64)
+{
+    constexpr auto t = getTileConfig(64, 64, DataType::FP16, GpuTarget::gfx942);
+
+    EXPECT_EQ(t.hdim_q, 64);
+    EXPECT_EQ(t.hdim_v, 64);
+    EXPECT_EQ(t.bm0, 32);
+    EXPECT_EQ(t.bn0, 128);
+    EXPECT_EQ(t.bk0, 64);
+    EXPECT_EQ(t.bk1, 32);
+    EXPECT_EQ(t.bk2, 64);
+    EXPECT_EQ(t.bk3, 32);
+    EXPECT_EQ(t.bk4, 32);
+    EXPECT_EQ(t.rm2, 1);
+    EXPECT_EQ(t.rn2, 4);
+    EXPECT_EQ(t.rk2, 1);


+{
+    EXPECT_EQ(wavefrontSize(GpuTarget::gfx1100), 32);
+    EXPECT_EQ(wavefrontSize(GpuTarget::gfx1101), 32);
+}


Add FmhaArch to DqDkDv specs, resolve arch-dependent tile geometry for gfx9/gfx950/gfx11/gfx12. Include arch in registry matching. Extend existing FMHA BWD tests for arch propagation and dispatch behavior.

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

…hpc/ck/rocm-ck-p3-1-tile-config

…ariant Three follow-ups from PR review: 1. Add an early throw in makeSpec for the four asymmetric tuples (32,128), (96,128), (32,256), (96,256) so the failure names the asymmetric case rather than bubbling up as "no valid GEMM4 warp distribution" from a consteval helper several layers down. 2. Split the throw strings in computeGemm4Warps and getBaseTile so each compile-time error points at the precondition that fired (dtype vs target vs missing-entry vs bm0 multiple-of-16 vs no warp split) instead of a single generic message. 3. Add a static_assert over GFX9_FP16_DQDKDV_BASE_TILES enforcing bn0 % bk4 == 0 (the bk4 doc-stated invariant), so adding a new row that breaks it fails the build with a pointed message instead of surfacing later as an unrelated tile-derivation error. Signed-off-by: Adam Osewski <Adam.Osewski@amd.com> Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

… if constexpr (#7638) Closes #7560 Note: Part of it was done in #7506 --------- Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com> Signed-off-by: Adam Osewski <Adam.Osewski@amd.com> Co-authored-by: Adam Osewski <Adam.Osewski@amd.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

CK Tile's fmha_bwd.py get_dq_dk_dv_tiles() defines tuned dQ/dK/dV tile configs only for symmetric head dimensions (hdim_q == hdim_v). The prior table keyed on max(hdim_q, hdim_v) and synthesized configs for asymmetric pairs by extrapolation, which has no validated tile and could select an illegal GEMM4 warp distribution. Reject asymmetric configs instead. Addresses PR #7538 review findings: - Collapse generateTileConfig(hdim_q, hdim_v) to generateTileConfig(hdim); getTileConfig now throws on hdim_q != hdim_v as the single gate, so makeSpec's duplicate asymmetric reject-list is removed (#1, #4). - Keep the consteval derivation engine, symmetric-only (#3). - Add validateWaveTiles() so a future base-tile edit that yields an illegal MFMA/WMMA warp shape fails the build; uses the previously unused isValidWaveTile() as the single source of truth (#6). - Fix the "occupancy -1 = auto" doc contradiction and assert the table stores occupancy >= 1 (#2). - Remove the dead GEMM4 (4,1) warp branch; it is unreachable for the shipped gfx9 base tiles (#8). - Reword computeGemm4Warps as a divisibility rule, not a head-dim-size heuristic (#11). - Drop the now self-referential asymmetric tests from the consteval and compat suites; add two compile_fail/ negative TUs covering both asymmetric directions (#7, #9). Verified: positive smoke compiles; both negatives fail to compile with "requires hdim_q == hdim_v"; gtest green (consteval 21/21, dqdkdv 34/34); clang-format-18 applied. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Resolve conflicts from base advancing past the PR fork point: - tests/test_fmha_bwd_consteval.cpp (add/add): base [P1-7] (#7504) added a same-named file with SWA/sliding-window mask-slot consteval tests, while this PR's file holds tile-config dispatch tests. Combine both into one file -- the TileConfig suite (tile-config dispatch, symmetric-only) plus the FmhaBwdConsteval suite (SWA/CMaskBR mask-slot pins) -- with a unified preamble (union of includes/usings + the mask slot-layout static_asserts). - tests/CMakeLists.txt: keep a single test_fmha_bwd_consteval target with the examples/06_rocm_ck_fmha_bwd include dir, which the merged file needs for rocm_fmha_bwd_registry.hpp. - tests/test_fmha_bwd_compat.cpp: git text-merged 5 byte-identical frozen-baseline tests that both branches added independently (OGradDotO D32/D96/D256, ConvertDQ D64/D256) into duplicate TEST definitions that broke the build; drop the redundant copies. Verified host-only: consteval 32/32 (21 TileConfig + 11 FmhaBwdConsteval), compat 93/93, dqdkdv 34/34; both asymmetric compile_fail negatives still fail with "requires hdim_q == hdim_v"; clang-format-18 applied. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Resolve conflicts between PR #7807 (hardcoded d32/d96/d256 DqDkDv tile configs) and the base's PR #7538 (consteval getTileConfig tile-config table), which generalized the same tile geometry PR #7807 hardcoded. Resolution: - dqdkdv_dev.hpp / dqdkdv_spec.hpp: take the base's table-driven design. Its GFX9_FP16_DQDKDV_BASE_TILES table already covers d32/d64/d96/d128/d256 and generateTileConfig() reproduces the exact tile values PR #7807 wired by hand (verified d32/d96/d256), so the device bridge supersedes both the old hardcoded std::conditional_t and the review-time alias refactor. block_per_cu now comes from tile.occupancy (table-driven). - test_fmha_bwd_compat.cpp: keep PR #7807's BF16 d32/d96/d256 spec tests and the d32/d96/d256 registry-lookup tests; drop PR #7807's FP16 spec-test copies (the base added equivalents) and the arch assertions (the base removed the spec 'arch' field, #7633/#7856). Set ALL_DQDKDV_VARIANTS_COUNT to 28 (merged registry: base's variants + PR #7807's six, all unique). - main.cpp: the runner OOB guard for non-d128 hdims is preserved. PR #7807's durable contribution after the merge is the six d32/d96/d256 .hip entrypoints + registry/CMake/pack.py wiring + tests; the device tile geometry is now provided by the base table. Verified: host compat suite compiles and passes 102/102 (incl. variant count 28). The kpack device archive remains broken at the base by pre-existing CK Tile / __AMDGCN_WAVEFRONT_SIZE drift, unaffected by this PR. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

github-actions Bot added the project: composablekernel label May 18, 2026

chris-tsiaousis-hpc added the organization: streamhpc contributors from streamhpc label May 18, 2026

chris-tsiaousis-hpc requested a review from bidlekm May 18, 2026 15:21

chris-tsiaousis-hpc linked an issue May 18, 2026 that may be closed by this pull request

[P3-1] Tile-config table design (consteval get_tile_config) #7506

Closed

chris-tsiaousis-hpc requested a review from Copilot May 19, 2026 09:50

Copilot started reviewing on behalf of chris-tsiaousis-hpc May 19, 2026 09:56 View session

Copilot AI reviewed May 19, 2026

View reviewed changes

chris-tsiaousis-hpc marked this pull request as ready for review May 20, 2026 14:51

chris-tsiaousis-hpc requested a review from a team as a code owner May 20, 2026 14:51

chris-tsiaousis-hpc requested a review from aledudek May 20, 2026 15:01

chris-tsiaousis-hpc added 3 commits May 28, 2026 09:54

Added tile config with consteval lookup

14ecf26

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

Added tests for consteval tile config lookup

0cfd402

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

Address copilot review comments

7bc25ce

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

chris-tsiaousis-hpc force-pushed the users/chris-tsiaousis-hpc/ck/rocm-ck-p3-1-tile-config branch from 1bcb6e3 to 7bc25ce Compare May 28, 2026 07:54

Merge branch 'users/shumway/ck/exp-kpack' into users/chris-tsiaousis-…

c6bd381

…hpc/ck/rocm-ck-p3-1-tile-config

aosewski mentioned this pull request May 29, 2026

[P3-8] Runtime smoke tests for new dQ/dK/dV tile configs (d32/d64/d96/d256 + asymmetric) #7886

Open

aosewski mentioned this pull request May 29, 2026

[CK rocm_ck] GpuTarget for FMHA BWD DQDKDV tile selection (P4-1)(#7629) #7823

Closed

1 task

chris-tsiaousis-hpc and others added 3 commits June 1, 2026 14:01

aosewski merged commit 1a217a6 into users/shumway/ck/exp-kpack Jun 1, 2026
1 of 2 checks passed

aosewski deleted the users/chris-tsiaousis-hpc/ck/rocm-ck-p3-1-tile-config branch June 1, 2026 14:38

bidlekm mentioned this pull request Jun 4, 2026

[P3-7] IsWG32 derivation from tile config #7808

Open

chris-tsiaousis-hpc added a commit that referenced this pull request Jun 4, 2026

Import the changes from PRs #7538, #7274 and #7928

7e21ee1

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

Navigation Menu

[P3-1] Tile-config table design (consteval get_tile_config) - #7538

[P3-1] Tile-config table design (consteval get_tile_config)#7538
aosewski merged 8 commits into
users/shumway/ck/exp-kpackfrom
users/chris-tsiaousis-hpc/ck/rocm-ck-p3-1-tile-config

chris-tsiaousis-hpc commented May 18, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

		constexpr GpuTarget target = GpuTarget::gfx942;
		auto tile = getTileConfig(sig.hdim_q, sig.hdim_v, sig.dtype, target);

Conversation

chris-tsiaousis-hpc commented May 18, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants