Skip to content

feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 - #11043

Merged
ozturkosu merged 28 commits into
developfrom
users/muozturk/ck/rowcolquant-tensorquant-gfx1250-enable
Sep 18, 2026
Merged

ozturkosu merged 28 commits into
developfrom
users/muozturk/ck/rowcolquant-tensorquant-gfx1250-enable

Conversation

@ozturkosu

@ozturkosu ozturkosu commented Aug 19, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID: AICK-2108
JIRA ID: AICK-2141

What this PR does

This PR enables the grouped RowColQuant and TensorQuant GEMM dispatcher bridges on gfx1250 for the existing FP8/BF8 rcr path.

It adds the correct gfx1250 WMMA tile, wires architecture handling through the runtime and codegen paths, and rejects configurations that can compile or launch but return incorrect output.

Motivation

The existing default uses the gfx9 MFMA 32x32x16 warp fragment. gfx1250 uses WMMA and does not provide that fragment. The kernel still compiles and launches, but returns incorrect or all-zero output.

Because this fails silently, adding gfx1250 to the architecture allowlist is not enough. Every entry point must select a valid tile and fail closed when the architecture, tile, or shape is unsupported.

Summary

  • Adds the gfx1250 FlatMM default:
    • tile: 16x64x256
    • warp layout: 1x4x1
    • warp tile: 16x16x128
  • Uses one shared tile selector for the runtime and standalone codegen paths.
  • Adds --gfx-arch support to both code generators and rejects unknown targets.
  • Normalizes feature-suffixed architecture names such as gfx1250:xnack-.
  • Adds gfx1250 validation in both Python and C++, including compile-time guards for invalid WMMA tiles and warp layouts.
  • Improves runtime diagnostics by reading TileN and PadN from the compiled kernel instead of hardcoding them.
  • Versions the .so cache ABI so older libraries cannot be reused after exported symbols change.
  • Extends the GPU harnesses and registers the new host-side regression tests with CTest.

Design Note

The tile selection intentionally matches gfx1250 exactly, not the full gfx12 family. gfx1200 and gfx1201 also use WMMA, but their 8-bit fragment is 16x16x16; giving them the gfx1250 tile can compile successfully and return incorrect output. OCP-FP8 encoding remains gfx12-family-wide because that behavior is shared across gfx12 targets.

The Python validator provides an early, readable error. The C++ static_assert remains as a second line of defense for generated-header and CMake builds. This duplication is intentional because an invalid gfx1250 tile does not reliably fail during compilation.

For RowColQuant shapes where M % 4 != 0, this PR rejects the call instead of modifying the shared quant kernel without gfx942/gfx950 hardware coverage.

Limitations

Area Current limitation
Target This change enables gfx1250 only. It does not enable gfx1200 or gfx1201.
Supported surface FP8/BF8 inputs, FP16 output, and rcr layout only.
Shape constraints K % 16 == 0; N % TileN == 0 when PadN=false (TileN=64 for the default gfx1250 tile). RowColQuant on gfx12 also requires M % 4 == 0; TensorQuant has no M restriction.
Generated-header build The header must be generated with the matching --gfx-arch gfx1250 option before the CMake build. A mismatched tile now fails at compile time.
Validation scope On-device coverage used k_batch=1 on gfx1250. gfx942/gfx950 were covered by host-side non-regression checks but were not rerun on hardware in this round.
Performance Correctness only. The tile is validated, not performance-tuned, and Old-TE has no gfx1250 backend for a parity baseline.

Validation

Hardware: gfx1250 , ROCm 10.1, single GPU.

  • Compiled and ran all four default paths: RowColQuant and TensorQuant × FP8 and BF8.
  • Tested five shapes per path, including aligned, non-aligned, and non-square cases.
  • Compared GPU output against an FP32 dequantized reference; worst global relative error was 4.22e-4.
  • Verified that RowColQuant shapes with M % 4 != 0 are cleanly rejected, while TensorQuant remains correct for the same M values.
  • Ran a 14,208-row configuration sweep. Every silent wrong-result row used the invalid 32x32x32 warp tile; 8-warp layouts were also launch-unstable and are now rejected.
  • Verified bare and feature-suffixed gfx1250 names end to end.
  • Verified codegen defaults, runtime/codegen name alignment, architecture validation, cache ABI handling, and compile-time tile guards with host-side tests.
  • Dispatcher suite result: 212 passed, 0 failed.

Negative controls were also run by reverting each required enablement independently:

  • Reverting the gfx1250 tile produced incorrect output (max_rel=1.0).
  • Removing gfx1250 from the C++ architecture table caused initialization to fail.
  • Restoring the hardcoded gfx950 test target triggered the compile/runtime architecture mismatch guard.

@ozturkosu
ozturkosu requested a review from a team as a code owner August 19, 2026 22:08
@ozturkosu ozturkosu self-assigned this Aug 19, 2026
@ozturkosu
ozturkosu changed the base branch from users/tlakshma/ck/te_bs_bridge to develop August 19, 2026 22:15
@therock-pr-bot

therock-pr-bot Bot commented Aug 19, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@sa-faizal sa-faizal added the ci:asan Opt-in to building ASAN label Aug 19, 2026
@sa-faizal sa-faizal added ci:host-asan Opt-in to running multi-arch host-asan CI on a pull request and removed ci:asan Opt-in to building ASAN labels Aug 20, 2026
@therock-pr-bot

therock-pr-bot Bot commented Aug 21, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

… gfx1250 (MI400)

Re-applies the gfx1250 enablement of PR #11043 on top of current develop. The
original branch predates #10439, which merged the bridges themselves into
develop, so only the gfx1250 delta is carried over here.

- codegen_common.py: add ROWCOL_TENSOR_QUANT_DEFAULT_TILE_GFX12 and the
  rowcol_tensor_quant_default_tile(gfx_arch) selector. The stock default tile is
  the gfx9 MFMA 32x32x16 fragment, which does not exist on gfx12 WMMA hardware:
  the kernel compiles but produces all-zero output. gfx12 needs the FlatMM 8-bit
  tile 16x64x256 / warp 1x4x1 / warp_tile 16x16x128.
- grouped_gemm_{rowcolquant,tensorquant}_utils.py: _default_config() selects the
  tile through that helper. gfx942/gfx950 kernel names are unchanged.
- grouped_gemm_{rowcolquant,tensorquant}_ctypes_lib.cpp: add gfx1250 to
  kSupportedArchs, which otherwise rejects the .so in dispatcher_initialize with
  "compile-time GFX_ARCH 'gfx1250' is not a supported architecture".
- test_{rowcolquant,tensorquant}_gpu_correctness.py: add gfx1250 to
  _SUPPORTED_ARCHES, and default --gfx to the detected device instead of a
  hardcoded gfx950 (which built a gfx950 image and aborted on any other GPU).

Validated on gfx1250 / MI400 (ROCm 10.1): both harnesses PASS fp8 + bf8.
@ozturkosu
ozturkosu force-pushed the users/muozturk/ck/rowcolquant-tensorquant-gfx1250-enable branch from db87a19 to c2ccdc2 Compare September 1, 2026 23:14
ozturkosu added a commit that referenced this pull request Sep 2, 2026
JIRA ID : AICK-2101
JIRA ID : AICK-2133

## Summary
This PR enables the **multi-ABD** (`a2_b2_d2`) GEMM TE→dispatcher bridge
(#9305) on **gfx1250 / MI400**. It was split out of #10921 so the other
three ops (`gemm_universal`, `grouped_gemm`, `gemm_multi_d`) weren't
held back.



Comparison of this PR's original file set vs `develop` (branch head
`56e7bc3`):



**So the substantive multi-ABD gfx1250 enablement is already on
`develop`; the remaining unique delta in this PR is a single config file
(`example_problems.json`).**

## Motivation (why the enablement was needed)
gfx1250 runs **WMMA**, not MFMA: the valid fp16/bf16 warp tile is
**16x16x32**, whereas the CDNA MFMA **32x32x16** tile in
`default_ci_config.json` instantiates zero runnable kernels there.

## Supported surface (gfx1250)
The `a2_b2_d2` variant is intentionally narrow — the codegen, ctypes
lib, and runner are all **fp16-only** and **rcrr-only** (`expand_sweep`
rejects other dtypes/layouts). So **fp16 × rcrr is the complete
supported matrix**.

## Validation (on-device, MI400 / gfx1250)
Old-TE has no gfx1250 back end, so this is **correctness-only** (GPU
output vs an fp32 CPU reference), not a performance baseline.

Full `default_config` sweep of the fp16/rcrr surface, 3 tile-divisible
shapes (768³ / 1536³ / 3072³):

| op | supported matrix | kernels | result rows | fail |
|----|------------------|--------:|------------:|-----:|
| `gemm_multi_abd` | fp16 × rcrr (only) | 1,640 | 4,920 | **0** |

Every kernel in the supported surface is correct on gfx1250 (rel err ≈
4e-4).

## Note on #11161 (compv3 + intrawave + 8-warp)
The full sweep initially surfaced 552 failures, all in the `compv3 +
intrawave + 8-warp (2x4x1 / 4x2x1)` class — the **#11161** defect (a
wave64 assumption in the `comp_v3` pipeline, wrong on gfx1250 wave32).
An interim gate excludes that class so the supported surface reports
clean; the real fix, **#11175** (`Use8WarpSchedule`), is already on
`develop`.

## Notes / dependencies
- Built on #10921's shared gfx1250 foundation (now on `develop` via the
sibling series).
- Supersedes the earlier closed attempt **#10917**.
- Sibling gfx1250 bridge PRs: #10921 (universal / grouped / multi-D),
#10927 (grouped quant), #10928 (batched GEMM + contraction), #11043
(rowcol / tensor quant), #10934 (contraction multi-ABD).

---------

Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
@ozturkosu
ozturkosu requested a lite review from Copilot September 3, 2026 03:34

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

The updated GPU correctness tests can still incorrectly skip on supported devices because the arch check uses exact matching rather than prefix/suffix-tolerant matching.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

Enables grouped RowColQuant and TensorQuant GEMM dispatcher bridges on gfx1250 (MI400) by making the default tile selection arch-aware for gfx12xx WMMA and extending the correctness harnesses/bridges to recognize the new arch.

Changes:

  • Add a gfx12xx-specific default tile (FlatMM 8-bit) and a shared rowcol_tensor_quant_default_tile(gfx_arch) selector in codegen_common.py.
  • Update RowColQuant/TensorQuant Python runtime helpers to use the shared arch-aware tile selector when building default configs.
  • Extend RowColQuant/TensorQuant GPU correctness harnesses and the ctypes bridge supported-arch lists to include gfx1250.
File summaries
File Description
projects/composablekernel/dispatcher/tests/test_tensorquant_gpu_correctness.py Adds gfx1250 support and makes the standalone runner default to the detected device arch.
projects/composablekernel/dispatcher/tests/test_rowcolquant_gpu_correctness.py Adds gfx1250 support and makes the standalone runner default to the detected device arch.
projects/composablekernel/dispatcher/python/grouped_gemm_tensorquant_utils.py Switches default tile selection from a constant dict to an arch-aware selector.
projects/composablekernel/dispatcher/python/grouped_gemm_rowcolquant_utils.py Switches default tile selection from a constant dict to an arch-aware selector.
projects/composablekernel/dispatcher/codegen/codegen_common.py Introduces gfx12xx WMMA default tile and a shared helper to select the proper tile by arch.
projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_tensorquant_ctypes_lib.cpp Adds gfx1250 to the supported-arch list for the bridge.
projects/composablekernel/dispatcher/bindings/ctypes/grouped_gemm_rowcolquant_ctypes_lib.cpp Adds gfx1250 to the supported-arch list for the bridge.
Review details
  • Files reviewed: 7/7 changed files
  • Comments generated: 3
  • Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

…silent failures

Addresses review feedback on #11043 plus two defects found while validating the
enablement on gfx1250 hardware. All three share a failure mode: the code was
wrong in a way that reported success.

1. Arch gate silently skipped on supported devices.
   The GPU correctness tests gated on `_GFX_ARCH in _SUPPORTED_ARCHES`, an exact
   string match. rocm_agent_enumerator may report feature suffixes, so a
   supported device enumerating as "gfx942:sramecc+:xnack-" caused the tests to
   SKIP. A skip is not a failure, so CI reported green while running none of the
   GPU coverage. Both test modules now normalize the arch (strip at the first
   ':') and compare the whole base token, matching the prefix tolerance the C++
   bridges already had in is_supported_arch(). Token comparison rather than
   startswith, so a hypothetical "gfx9421" is not accepted on the strength of the
   "gfx942" prefix. Pinned by a new CPU-only test module.

2. RowColQuant silently zeroed row M-1 of C when M % 4 != 0 on gfx1250.
   MakeAQBlockWindow() builds the AQ view as a broadcast
   make_naive_tensor_view({M, N}, {1, 0}) and then creates the tile window with
   no padding transform, unlike the A/B/C windows which pad from
   kPadM/kPadN/kPadK. The M axis therefore has an unguarded tail. TensorQuant is
   correct at the same M because it passes scalar scales; the BQ view has the
   same unguarded tail but is unreachable because pad_n=false already forces
   N % 64 == 0.

   The root-cause fix belongs in the shared ck_tile quant kernel header, which is
   also compiled for gfx942/gfx950 -- untested here. Rather than change tile
   distribution for architectures that were not measured, the bridge now rejects
   these shapes with an explicit message. The guard is scoped to gfx12 targets
   via the compile-time GFX_ARCH and to RowColQuant only, and its comment records
   the condition for removing it.

3. Undocumented argument constraints returned a bare -2.
   N % TileN is now pre-checked in both bridges, derived from the kernel's own
   kPadN/TileN constants rather than hardcoded. The -2 message enumerates the
   known constraints, including that K % 16 == 0 is required even with
   pad_k=true (pad_k covers the K-loop tail, not the global-load vector width).
   Constraints documented on the ABI comment and in both run() docstrings, and
   the Python RuntimeError is now self-contained instead of deferring to stderr.

   K % 16 is reported but not hard-enforced: it is checked inside ck_tile's
   IsSupportedArgument() and has not been shown to be general across dtypes and
   tiles, so enforcing it could reject shapes the kernel would accept.

Verified on gfx1250 / MI400: both GPU correctness harnesses 2/2 at
max_rel_err=0.0005; every M % 4 == 0 shape unchanged at 2.2e-4 to 3.9e-4; every
M % 4 != 0 shape now a clean rejection instead of a wrong answer; TensorQuant
still correct at those same M. Host-side suite 59 -> 95 passing, 0 failures.
clang-format-18 clean.

Known gap, unchanged: the codegen's _default_config() is still arch-agnostic, so
a CMake-built lib on gfx1250 would compile the gfx9 tile. Only the Python runtime
path is arch-aware. Tracked as a follow-up.
ozturkosu and others added 2 commits September 9, 2026 14:41
The RowColQuant/TensorQuant default-tile selector matched the whole gfx12
family. gfx1200/gfx1201 are WMMA parts too, but their 8-bit warp fragment is
16x16x16, not gfx1250's 16x16x64/16x16x128, so they would have inherited a
tile that compiles cleanly and returns garbage. Nothing upstream stops that:
the runtime helpers have no arch allowlist. Match gfx1250 exactly instead.

The OCP-FP8 hipcc define stays family-wide on purpose -- every gfx12xx part
uses OCP FP8 encoding -- and is now written as an explicit family test with a
comment, so the two sites cannot be mistaken for the same rule.

Add a shared normalize_gfx_arch() in codegen_common as the single source of
truth for stripping feature suffixes ("gfx1250:xnack-"), use it at the arch
detection boundary so suffixes never reach --offload-arch/-DGFX_ARCH (the C++
side prefix-matches, so the bare target still matches a suffixed device), and
drop the duplicated local copies in the two GPU correctness tests.

The C++ derives its N-divisibility bound from SelectedKernel::TileN, but the
Python docstring and RuntimeError hardcoded 64. Export dispatcher_get_tile_n()
and dispatcher_get_pad_n() from both ctypes libs and use the real values, so
the diagnostics stay true if the compiled tile changes.

Tests cover suffix normalization, exact-vs-family arch matching for the tile
selector, and the intentional gfx12-family OCP-FP8 behaviour.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

The CMake path can silently execute an incompatible tile, while architecture normalization, cache compatibility, and standalone test targeting remain incomplete.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Review details
  • Files reviewed: 11/11 changed files
  • Comments generated: 9
  • Review effort level: Balanced

Preserve shared architecture feature flags, gfx1250 validation, explicit-target input encoding, exit-77 skip reporting, and rectangular GPU cases. Keep all host CTest registrations and complete the dispatcher_common properties block.

Validation: 406 host tests and 1010 subtests passed; all 9 selected CTest suites passed. Added six regression cases for standalone skip and suffix-override behavior.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Grouped-only validation affects another operator, and standalone CMake builds omit required gfx1250 feature definitions.

Get a fresh assessment by requesting another Copilot review.

Review details

Suppressed comments (1)

projects/composablekernel/dispatcher/bindings/ctypes/CMakeLists.txt:428

  • The standalone CMake path still does not mirror the gfx1250 feature defines used by the working JIT path. This branch adds only USE_NEW_UNIFIED_FRAMEWORK=0, while arch_feature_defines() also supplies OCP FP8, WMMA, gfx12, and gfx1250 defines; dispatcher_common.py:113-119 notes that at least CK_TILE_USE_OCP_FP8 must be explicit so host and device compilation agree. A standalone generated-header build can therefore compile with a different FP8/architecture configuration than the validated Python path.
    if(CK_TILE_TENSORQUANT_GFX_ARCH STREQUAL "gfx1250")
        target_compile_definitions(dispatcher_grouped_gemm_tensorquant_lib PRIVATE USE_NEW_UNIFIED_FRAMEWORK=0)
    endif()
  • Files reviewed: 20/20 changed files
  • Comments generated: 2
  • Review effort level: Balanced

Apply target-scoped feature definitions matching the Python quant builds for gfx942, gfx950 and gfx1250. Preserve the gfx1250 legacy framework override and validate explicit/inferred, bare/suffixed targets against arch_feature_defines().

Regression: eight parity cases fail before the fix. Afterward 97 tests and 45 subtests plus 9/9 selected CTest suites pass. Fresh standalone CMake builds on MI400 pass all six RowColQuant/TensorQuant reference cases.
Move the gfx1250 warp-count and warp_k gates to the operator routing boundary so plain gemm_rowcolquant retains shared validation without grouped codegen restrictions. Keep common fragment checks for all three operators.

Add routing regressions for FP8/BF8 and bare/suffixed gfx1250 targets: eight plain RowColQuant cases reproduce the bug before the fix. Afterward 113 tests and 423 subtests plus all nine selected CTest suites pass.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Top-level multi-architecture builds can leak incompatible feature macros into the selected single-architecture quant bridge.

Get a fresh assessment by requesting another Copilot review.

Review details
  • Files reviewed: 20/20 changed files
  • Comments generated: 1
  • Review effort level: Balanced

@ozturkosu

Copy link
Copy Markdown
Contributor Author

For both RowColQuant and TensorQuant codegen, custom tile_configs supplied through --config or --config-json require an explicit --gfx-arch; omitting it stops generation with an error before any headers are written.

python unified_grouped_gemm_rowcolquant_codegen.py \
  --config my_tiles.json \
  --gfx-arch gfx1250 \
  --output-dir generated_kernels

@ozturkosu
ozturkosu requested a review from yraparti September 18, 2026 15:39
@ozturkosu
ozturkosu merged commit 6cea783 into develop Sep 18, 2026
116 checks passed
@ozturkosu
ozturkosu deleted the users/muozturk/ck/rowcolquant-tensorquant-gfx1250-enable branch September 18, 2026 20:38
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Sep 18, 2026
feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM
 bridges on gfx1250  (#11043)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

JIRA ID: AICK-2108
JIRA ID: AICK-2141

## What this PR does

This PR enables the grouped `RowColQuant` and `TensorQuant` GEMM
dispatcher bridges on `gfx1250` for the existing FP8/BF8 `rcr` path.

It adds the correct gfx1250 WMMA tile, wires architecture handling
through the runtime and codegen paths, and rejects configurations that
can compile or launch but return incorrect output.

## Motivation

The existing default uses the gfx9 MFMA `32x32x16` warp fragment.
gfx1250 uses WMMA and does not provide that fragment. The kernel still
compiles and launches, but returns incorrect or all-zero output.

Because this fails silently, adding gfx1250 to the architecture
allowlist is not enough. Every entry point must select a valid tile and
fail closed when the architecture, tile, or shape is unsupported.

## Summary

- Adds the gfx1250 FlatMM default:
  - tile: `16x64x256`
  - warp layout: `1x4x1`
  - warp tile: `16x16x128`
- Uses one shared tile selector for the runtime and standalone codegen
paths.
- Adds `--gfx-arch` support to both code generators and rejects unknown
targets.
- Normalizes feature-suffixed architecture names such as
`gfx1250:xnack-`.
- Adds gfx1250 validation in both Python and C++, including compile-time
guards for invalid WMMA tiles and warp layouts.
- Improves runtime diagnostics by reading `TileN` and `PadN` from the
compiled kernel instead of hardcoding them.
- Versions the `.so` cache ABI so older libraries cannot be reused after
exported symbols change.
- Extends the GPU harnesses and registers the new host-side regression
tests with CTest.

## Design Note

The tile selection intentionally matches `gfx1250` exactly, not the full
gfx12 family. `gfx1200` and `gfx1201` also use WMMA, but their 8-bit
fragment is `16x16x16`; giving them the gfx1250 tile can compile
successfully and return incorrect output. OCP-FP8 encoding remains
gfx12-family-wide because that behavior is shared across gfx12 targets.

The Python validator provides an early, readable error. The C++
`static_assert` remains as a second line of defense for generated-header
and CMake builds. This duplication is intentional because an invalid
gfx1250 tile does not reliably fail during compilation.

For RowColQuant shapes where `M % 4 != 0`, this PR rejects the call
instead of modifying the shared quant kernel without gfx942/gfx950
hardware coverage.

## Limitations

| Area | Current limitation |
|---|---|
| Target | This change enables `gfx1250` only. It does not enable
`gfx1200` or `gfx1201`. |
| Supported surface | FP8/BF8 inputs, FP16 output, and `rcr` layout
only. |
| Shape constraints | `K % 16 == 0`; `N % TileN == 0` when `PadN=false`
(`TileN=64` for the default gfx1250 tile). RowColQuant on gfx12 also
requires `M % 4 == 0`; TensorQuant has no M restriction. |
| Generated-header build | The header must be generated with the
matching `--gfx-arch gfx1250` option before the CMake build. A
mismatched tile now fails at compile time. |
| Validation scope | On-device coverage used `k_batch=1` on gfx1250.
gfx942/gfx950 were covered by host-side non-regression checks but were
not rerun on hardware in this round. |
| Performance | Correctness only. The tile is validated, not
performance-tuned, and Old-TE has no gfx1250 backend for a parity
baseline. |

## Validation

Hardware: `gfx1250` , ROCm 10.1, single GPU.

- Compiled and ran all four default paths: RowColQuant and TensorQuant ×
FP8 and BF8.
- Tested five shapes per path, including aligned, non-aligned, and
non-square cases.
- Compared GPU output against an FP32 dequantized reference; worst
global relative error was `4.22e-4`.
- Verified that RowColQuant shapes with `M % 4 != 0` are cleanly
rejected, while TensorQuant remains correct for the same M values.
- Ran a 14,208-row configuration sweep. Every silent wrong-result row
used the invalid `32x32x32` warp tile; 8-warp layouts were also
launch-unstable and are now rejected.
- Verified bare and feature-suffixed gfx1250 names end to end.
- Verified codegen defaults, runtime/codegen name alignment,
architecture validation, cache ABI handling, and compile-time tile
guards with host-side tests.
- Dispatcher suite result: `212 passed, 0 failed`.

Negative controls were also run by reverting each required enablement
independently:

- Reverting the gfx1250 tile produced incorrect output (`max_rel=1.0`).
- Removing gfx1250 from the C++ architecture table caused initialization
to fail.
- Restoring the hardcoded `gfx950` test target triggered the
compile/runtime architecture mismatch guard.
ozturkosu added a commit that referenced this pull request Sep 18, 2026
All gfx1250 behaviour below is GPU-verified on MI400 (gfx1250, ROCm 10.1) as
part of this change; the numbers quoted are from those runs.

- abquant fp4 on gfx1250 used warp_tile_k=32, which is a silent-wrong-answer
  bug. abquant dispatches on AComputeDataType and auto_compute_type collapses
  packed A==B==pk_fp4_t to fp8_t, so at K=32 the fp4 kernel selects the
  unguarded MFMA Dispatcher<fp8_t,fp8_t,float,16,16,32>, whose intrinsics are
  #if __gfx94__/__gfx95__ only and otherwise return CVecType{0.f}. Measured on
  MI400: it builds, launches, and yields a dead accumulator -- ~50% exact zeros
  and ~50% uncorrelated noise (corr vs reference -0.004). It is NOT uniformly
  zero, so the all-zero guards in the GPU harnesses do not catch it. K=64 and
  K=128 both verify at 4.74e-4 max relative error (fp8 control: 4.70e-4);
  gfx1250 now uses 128 for every variant, so the arch has one uniform tile.

- fp8i4 / bf8i4 are now rejected on gfx1250. Across {fp8i4,bf8i4} x warp_tile_k
  {16,32,64,128} x N {128,256,512}, all 24 combinations built and launched and
  all 24 returned NaN, while C4/fp8 and C4/bf8 passed in the same harness with
  the same host codec. The rejection lives in the config dataclass __post_init__
  so it covers the public constructors, the internal builders and sweep
  expansion alike. Note this contradicts the prose in the grouped sibling, which
  says i4 "does NOT compile on gfx1250": it compiles cleanly and is wrong at
  runtime, which is the more dangerous direction.

- grouped bquant regressed to warp_tile_k=32 on gfx1250. The merge from develop
  generalized the previously hardcoded warp_tile_k=128 into _fp8_warp_tile_k()/
  _preshuffleb_warp_tile_k() but gated them on gfx950 only, so
  default_fp8_config_gfx1250() -- which delegates straight to default_fp8_config
  -- began emitting a 16x16x32 kernel. Restored via an exact, suffix-tolerant
  arch set. This is what TestGfx1250Configs was already asserting and had
  started failing on.

- _validate_arch now exact-matches after stripping the feature suffix, so
  gfx12500 is rejected instead of passing startswith and silently falling
  through to the legacy tile. gfx1250:xnack- and gfx942:sramecc+:xnack- still
  validate, and the full target id is returned unchanged for --offload-arch.

- codegen_common.fp8_warp_tile_k_for_arch aligned on gfx1250: --list-names
  --gfx-arch gfx1250 reported 16x16x32 while the runtime built 16x16x128. That
  helper's own docstring requires the rule to exist exactly once. It now routes
  through the shared normalize_gfx_arch() that landed with #11043.

- OCP fp8 compile defines extracted into a per-module _ocp_fp8_arch_defines()
  so the family-wide-gfx12 invariant is unit-tested across all five bridges;
  this also fixes abquant, the one module that gated them on gfx950 only. The
  device pass was already OCP via config.hpp's __gfx12__ fallback, so this
  removes a host/device constexpr skew and a latent trap, not wrong numerics.

- _is_gfx1250 is evaluated first at every call site: `in` raised TypeError on
  None before the helper ran, making its documented None tolerance unreachable
  from the bridges. Selectors now raise on a missing arch rather than
  defaulting to the gfx942 tile.

- Docstrings no longer claim the selectors mirror get_k_warp_tile on gfx1250.
  That helper returns 64 there (WMMA branch at M_Warp_Tile==16, reachable
  because CMakeLists sets CK_TILE_USE_WMMA=1 for gfx12); 128 is a deliberate,
  now GPU-verified divergence and is documented as such.

Arch-gate suite 168 -> 317 tests. Full dispatcher suite, same node and PATH:
1418 -> 1488 passed with the two TestGfx1250Configs failures fixed and none
introduced (28 -> 26 on the GPU runner; the remainder are pre-existing and are
not in files this PR touches).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

ThruptiRajLakshmanaGowda added a commit that referenced this pull request Sep 22, 2026
…12392)

## Proposed changes

Backs the dispatcher CI lanes out of `develop` after they broke the
pipeline, by restoring the three pipeline files to their state
immediately before #10897:

- `projects/composablekernel/Jenkinsfile`
- `projects/composablekernel/Dockerfile`
- `projects/composablekernel/groovy/vars/ck.groovy`

All three are **byte-identical to `7080fbe1ca^`**, so the pipeline
returns to a previously-green configuration.

This is a **CI-files-only** revert. The dispatcher Python, codegen and
test sources are deliberately left untouched on `develop`. It is the
minimal-blast-radius alternative to #12390, which reverts those three
PRs in full (51 files).

Reverts the pipeline changes from:

| PR | Merge commit |
|---|---|
| #10897 — feat(ck-tile): add dispatcher correctness and performance CI
lanes for gfx942/gfx950 | `7080fbe1ca` |
| #12320 — fix(ck): restore declarative stages so skipped CI lanes stop
reporting success | `40995238ed` |
| #11879 — feat(ck): add 7 operators to dispatcher correctness and perf
lanes | `5873d0a0f6` |

### Why restoring only these three files is safe

- **No other author's work is affected.** Every commit that touched
these three files since #10897 is one of the three PRs being backed out.
Verified per file with `git log 7080fbe..develop -- <file>`. This
matters because three gfx1250 PRs (#11043, #10934, #12216) landed on the
*dispatcher* files after #10897 — but none of them touched the pipeline
files, so this revert cannot disturb them.
- **The leftover dispatcher sources are inert for the main build.**
`dispatcher/` is a standalone CMake project and is never reached by an
`add_subdirectory()` from the top-level `CMakeLists.txt`.
`CK_TILE_DISPATCHER` only adds a compile definition and an include
directory, so `dispatcher/tests/CMakeLists.txt` is not configured by a
normal build.
- **The one pre-existing `CK_TILE_DISPATCHER=ON` lane is unaffected.**
`runFullGroupedConvTileTests` predates #10897, uses
`DISPATCHER_RULE_SET=tests` for conv only, and builds and runs three
named conv binaries rather than invoking `ctest`. It does not reach
`dispatcher/tests`.
- **No dangling callers.** `ck.groovy` defines the same 54 top-level
functions before and after; the reverted PRs changed function bodies,
not the API.

### Verification

- The three files are byte-identical to `7080fbe1ca^`, and to the
corresponding files on `users/tlakshma/ck/revert-dispatcher-ci` (#12390)
- Diff is confined to `projects/composablekernel/`; no dispatcher or
tile_engine source is touched
- Per-file authorship of every commit since #10897 confirmed
- `ck.groovy` top-level function set compared before/after

Not verified: the `Jenkinsfile` and `ck.groovy` could not be
syntax-checked locally (no `groovy`/`groovyc` available). Mitigated by
both being byte-identical to a revision that was previously green.

### Relationship to #12390

#12390 and this PR are **alternatives, not companions** — they overlap
on these three files. Whichever lands second will need a rebase, or
should be closed. Pick based on whether the leftover dispatcher sources
on `develop` are acceptable:

- **This PR** — smallest change, restores CI immediately, but leaves
dispatcher Python and its CTest registrations on `develop` with no lane
exercising them.
- **#12390** — also removes that now-unexercised code, at the cost of a
51-file diff with hand-resolved conflicts around the gfx1250 work.

### Follow-up

#12320 was itself a *fix* for skipped lanes falsely reporting success.
Reverting it restores that false-green behaviour. If the breakage is
traced to #10897 or #11879 alone, #12320 should be re-landed on top
rather than left reverted.

Issue IDs: AICK-2170 (#12320), AICK-2171 (#11879). #10897 had no issue
ID.

## Checklist

- [ ] I have added tests relevant to the introduced functionality, and
the unit tests are passing locally
- [ ] I have added the test to REGRESSION_TESTS list defined at the top
of CMakeLists.txt in tests/CMakeLists.txt, **IF** the test takes more
than 30 seconds to run.
- [x] I have added inline documentation which enables the maintainers
with understanding the motivation
- [ ] I have removed the stale documentation which is no longer relevant
after this pull request
- [ ] (If this change is user-facing) I have added release notes which
provide the end users with a brief summary of the improvement from this
pull request
- [ ] I have run `clang-format` on all changed files
- [x] Any dependent changes have been merged

No tests are added: this PR only restores CI pipeline configuration to a
prior state and changes no library or test source.

## Discussion

The alternative is the full revert in #12390. A full revert is tidier in
that it does not leave unexercised code behind, but it requires
hand-resolving conflicts against three gfx1250 PRs that landed on the
dispatcher files after #10897, where a naive resolution silently deletes
another author's work. This PR avoids that risk entirely by touching
only files that no one else has modified, at the cost of leaving the
dispatcher sources in place until they are cleaned up or re-landed.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
evedovelli pushed a commit that referenced this pull request Sep 23, 2026
…12392)

## Proposed changes

Backs the dispatcher CI lanes out of `develop` after they broke the
pipeline, by restoring the three pipeline files to their state
immediately before #10897:

- `projects/composablekernel/Jenkinsfile`
- `projects/composablekernel/Dockerfile`
- `projects/composablekernel/groovy/vars/ck.groovy`

All three are **byte-identical to `7080fbe1ca^`**, so the pipeline
returns to a previously-green configuration.

This is a **CI-files-only** revert. The dispatcher Python, codegen and
test sources are deliberately left untouched on `develop`. It is the
minimal-blast-radius alternative to #12390, which reverts those three
PRs in full (51 files).

Reverts the pipeline changes from:

| PR | Merge commit |
|---|---|
| #10897 — feat(ck-tile): add dispatcher correctness and performance CI
lanes for gfx942/gfx950 | `7080fbe1ca` |
| #12320 — fix(ck): restore declarative stages so skipped CI lanes stop
reporting success | `40995238ed` |
| #11879 — feat(ck): add 7 operators to dispatcher correctness and perf
lanes | `5873d0a0f6` |

### Why restoring only these three files is safe

- **No other author's work is affected.** Every commit that touched
these three files since #10897 is one of the three PRs being backed out.
Verified per file with `git log 7080fbe..develop -- <file>`. This
matters because three gfx1250 PRs (#11043, #10934, #12216) landed on the
*dispatcher* files after #10897 — but none of them touched the pipeline
files, so this revert cannot disturb them.
- **The leftover dispatcher sources are inert for the main build.**
`dispatcher/` is a standalone CMake project and is never reached by an
`add_subdirectory()` from the top-level `CMakeLists.txt`.
`CK_TILE_DISPATCHER` only adds a compile definition and an include
directory, so `dispatcher/tests/CMakeLists.txt` is not configured by a
normal build.
- **The one pre-existing `CK_TILE_DISPATCHER=ON` lane is unaffected.**
`runFullGroupedConvTileTests` predates #10897, uses
`DISPATCHER_RULE_SET=tests` for conv only, and builds and runs three
named conv binaries rather than invoking `ctest`. It does not reach
`dispatcher/tests`.
- **No dangling callers.** `ck.groovy` defines the same 54 top-level
functions before and after; the reverted PRs changed function bodies,
not the API.

### Verification

- The three files are byte-identical to `7080fbe1ca^`, and to the
corresponding files on `users/tlakshma/ck/revert-dispatcher-ci` (#12390)
- Diff is confined to `projects/composablekernel/`; no dispatcher or
tile_engine source is touched
- Per-file authorship of every commit since #10897 confirmed
- `ck.groovy` top-level function set compared before/after

Not verified: the `Jenkinsfile` and `ck.groovy` could not be
syntax-checked locally (no `groovy`/`groovyc` available). Mitigated by
both being byte-identical to a revision that was previously green.

### Relationship to #12390

#12390 and this PR are **alternatives, not companions** — they overlap
on these three files. Whichever lands second will need a rebase, or
should be closed. Pick based on whether the leftover dispatcher sources
on `develop` are acceptable:

- **This PR** — smallest change, restores CI immediately, but leaves
dispatcher Python and its CTest registrations on `develop` with no lane
exercising them.
- **#12390** — also removes that now-unexercised code, at the cost of a
51-file diff with hand-resolved conflicts around the gfx1250 work.

### Follow-up

#12320 was itself a *fix* for skipped lanes falsely reporting success.
Reverting it restores that false-green behaviour. If the breakage is
traced to #10897 or #11879 alone, #12320 should be re-landed on top
rather than left reverted.

Issue IDs: AICK-2170 (#12320), AICK-2171 (#11879). #10897 had no issue
ID.

## Checklist

- [ ] I have added tests relevant to the introduced functionality, and
the unit tests are passing locally
- [ ] I have added the test to REGRESSION_TESTS list defined at the top
of CMakeLists.txt in tests/CMakeLists.txt, **IF** the test takes more
than 30 seconds to run.
- [x] I have added inline documentation which enables the maintainers
with understanding the motivation
- [ ] I have removed the stale documentation which is no longer relevant
after this pull request
- [ ] (If this change is user-facing) I have added release notes which
provide the end users with a brief summary of the improvement from this
pull request
- [ ] I have run `clang-format` on all changed files
- [x] Any dependent changes have been merged

No tests are added: this PR only restores CI pipeline configuration to a
prior state and changes no library or test source.

## Discussion

The alternative is the full revert in #12390. A full revert is tidier in
that it does not leave unexercised code behind, but it requires
hand-resolving conflicts against three gfx1250 PRs that landed on the
dispatcher files after #10897, where a naive resolution silently deletes
another author's work. This PR avoids that risk entirely by touching
only files that no one else has modified, at the cost of leaving the
dispatcher sources in place until they are cleaned up or re-landed.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
ozturkosu added a commit that referenced this pull request Sep 26, 2026
…bridge (#12544)

JIRA ID: AICK-2427
JIRA ID: AICK-2137
JIRA ID: AICK-2136
JIRA ID: AICK-2142
JIRA ID: AICK-2429

Related PRs: #12508 #10921 #10927 #10928 #11043 #11048 #10916

## Summary

This PR enables the gfx1250-only GEMM pipelines in the tile_engine (TE)
/ dispatcher bridge:
- `comp_tdm` (TDM V1) and `comp_tdm_v2` (TDM V2);
- `comp_async`;
- the TDM epilogue.

They are enabled for GEMM universal and batched GEMM, whose gfx1250
bridge enablement is already merged. GEMM multi-ABD and the grouped
quant ops only gain reject-guards: their kernels have no async/TDM
pipeline yet.

Most of the change is in `tile_engine/` and `dispatcher/`. It also fixes
two CK-Tile kernel defects in pipelines this PR enables (see Kernel
fixes), with gtest coverage. On gfx942 and gfx950 the generated bridge
code is byte-identical.

## Motivation

The gfx1250 bridge enablement PRs (#10921, #10928, #11048) exposed only
the pipelines shared with gfx942/gfx950. The TDM and async pipelines
exist in CK-Tile for gfx1250, and #10916 already uses them for mx GEMM,
but the bridge could not generate or benchmark them for the other GEMM
ops.

## Design note

- **TDM pipelines** (`comp_tdm`, `comp_tdm_v2`):
- They are intrawave only, always use the TDM epilogue, and are
non-persistent.
- They require `pad_m = pad_n = pad_k = false`. TDM clips A/B/E in
hardware, and right-padding N/K inflates the bound-clip extents, which
gives wrong results on unaligned shapes. This matches the CK-Tile TDM
gtests, which pin padding off.
- The TE builders and dispatcher codegen reject any other combination
with a reason string.
- **`comp_async` on gfx1250:**
- It is emitted only for row-major A / col-major B. Other layouts hit
the transpose-load `static_assert` with the gfx1250 WMMA tiles.
- For non-MX ops it is also emitted only fully padded (`pad_m = pad_n =
pad_k = true`). On develop, unpadded `comp_async` on gfx1250 has a
kernel defect: the async K-prefetch reads past the A/B extent, and the
`TailNumber::Two` path has no LDS fence. The kernel fix goes into
#12508; this restriction can be relaxed once it merges.
  - fp8/bf8 `comp_async` on gfx1250 requires `warp_tile_k >= 128`.
  - mx GEMM is unaffected.
- **TDM LDS budget:** the gfx1250 TDM pipelines share the `comp_async`
LDS budget (aliases in `arch_specs.json`, emitted by
`generate_arch_specs.py`). The LDS check adds the TDM
16-byte-per-row-group descriptor padding and checks that the TDM output
tile fits. The tile_engine and the dispatcher ArchFilter use the same
rule. Non-MX `comp_async` GEMM on gfx1250 uses the same base LDS
descriptor, so it gets the same row padding in both gates (without it,
fp16/bf16 64x256x256, 192x128x256 and 256x64x256 fail to link at 337856
B > 327680 B); MX `comp_async` keeps its own rule from #10916.
- **GEMM multi-ABD:** reject-guards only for `comp_async`, `comp_tdm`,
`comp_tdm_v2` and the TDM epilogue.
- **Grouped quant ops** (`aquant`, `bquant`, `abquant`, `rowcolquant`,
`tensorquant`):
- These files only gain reject-guards for `comp_async`, `comp_tdm`,
`comp_tdm_v2` and the TDM epilogue. Nothing is enabled for them.
- The quant GEMM pipelines have no TDM/async implementation yet, so an
unguarded config would fail to compile instead of being filtered.
- **Runtime reject = SKIP.** The dispatcher returns status `-1` when a
kernel's `IsSupportedArgument` rejects a problem (for example TDM
alignment at `N % 8 != 0`). It is exposed as `GemmResult.unsupported`
(also on grouped and multi-D results). The gfx1250 tests and
`test_gemm_search_space` count it as a skip, not a failure, and a lane
still fails if nothing was verified.
- **Trait parsing** now goes through one shared helper (`trait_parse.py`
/ `gemm_trait_parse.cmake`), so the CMake and Python sides split kernel
trait strings the same way.

### Deferred ops

grouped GEMM, GEMM multi-D, batched contraction and contraction
multi-ABD are not enabled here. Their kernels need the core changes in
#12508 before they can host a TDM pipeline. A follow-up PR will enable
them after #12508 merges.

## Kernel fixes

Found by a gfx1250 tuning sweep with `--verify` (bf16 `rcr`,
5608x4096x2048):

- **`comp_async` with `persistent=true` wrote only part of the output
(9-36% of tiles non-zero).**
- Cause: `GemmPipelineAgBgCrCompAsync` did not expose
`UsePersistentKernel`. The host launched a persistent-sized grid, but
the device took the one-block-per-tile entry, so most tiles were never
written and out-of-range tile indices could fault.
- Fix: the pipeline now exposes `UsePersistentKernel` for plain GEMM
problems. MX problems stay non-persistent, because their host path
launches a split-K grid that the persistent loop does not read.
- The same pipeline is used on gfx950. The gfx950 persistent comp_async
tests pass with the fix.
- **`comp_tdm` / `comp_tdm_v2` returned all-NaN when the TDM policy
picks `sub_tile_num > 1`.**
- The TDM policy picks `sub_tile_num > 1` from its VGPR estimate. In the
sweep this happened at 4 waves with 128x128x256 and 128x256x128 tiles.
- Cause: the hot loop in `block_gemm_areg_breg_creg_v1.hpp` ran over the
full `KIterPerWarp` instead of one K sub-tile (`KIterPerWarp /
KSubTileNum`), so it read past the sub-tile register window.
- Fix: the loop now runs over one K sub-tile. With `KSubTileNum == 1` it
is unchanged; the compv4 code object is byte-identical.
- **`comp_tdm_v2`:** the 4-wave `static_assert` is now written as
`NumWarps == 4`. It is equivalent to the old block-size form, and the
pipeline body hardcodes four wave roles.
- **Tests:**
  - A Persistent variant is added to the comp_async WMMA gtest list.
- New `test_ck_tile_gemm_pipeline_tdm_subtile_wmma`: fp16/bf16, V1/V2,
with sub-tile 2 cases plus sub-tile 1 controls.
  - Both tests fail on the old kernel and pass with the fix.
- **Scope:** `comp_async_v2` (sub-tiling policy) shares the block GEMM
fix but has no gtest here and is not emitted by the bridge.

## Changes

| Area | Change |
|---|---|
| `dispatcher/codegen/arch_filter.py`, `codegen_common.py`,
`unified_gemm_codegen.py` | gfx1250 pipeline/epilogue emission and
validity rules (TDM pad/scheduler/persistent, `comp_async` layout + pad)
|
| `dispatcher/codegen/arch_specs_generated.py`,
`include/.../arch_specs_generated.hpp`, `kernel_key.hpp` | New
pipeline/epilogue enum values and names |
| `dispatcher/codegen/unified_grouped_gemm_*quant_codegen.py` |
Reject-guards only |
| `dispatcher/codegen/arch_specs.json`, `generate_arch_specs.py` | TDM
LDS budget aliases (generated, not hand-edited) |
| `dispatcher/python/gemm_utils.py` | Pipeline names,
`GemmResult.unsupported` for status `-1` |
| `dispatcher/tests/CMakeLists.txt` | Registers the gfx1250 pipeline,
gate-equivalence, arch-specs regeneration and result-status tests |
| `tile_engine/ops/gemm/gemm_instance_builder.py`,
`gemm_validation_utils.py` | Shared gfx1250 validity checks |
| `tile_engine/ops/gemm/gemm_universal/*`, `batched_gemm/*`,
`gemm_multi_abd/*` | Builders, CMake, gfx1250 default and CI configs |
| `tile_engine/ops/gemm/grouped_gemm_quant/*` | Reject-guards only |
| `dispatcher/codegen/codegen_common.py`,
`tile_engine/ops/gemm/gemm_validation_utils.py` | One shared async/TDM
reject guard per layer; the grouped quant codegens and the multi-ABD /
grouped quant builders bind it to their op name with `functools.partial`
instead of keeping a local copy |
| `tile_engine/ops/gemm/trait_parse.py`, `gemm_trait_parse.cmake` |
Shared trait-string parsing |
| `include/ck_tile/ops/gemm/block/block_gemm_areg_breg_creg_v1.hpp`,
`pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp`,
`pipeline/gemm_pipeline_ag_bg_cr_comp_tdm_v2.hpp` | Kernel fixes (see
above) |
| `test/ck_tile/gemm/*` | Persistent comp_async gtest variant, new TDM
sub-tile gtest |
| `test_*.py` (dispatcher and tile_engine) | Unit tests for the gates,
configs and guards |

## Testing

| Check | Result |
|---|---|
| gfx942 / gfx950 generated-code byte identity (before vs after, all
ops) | Identical |
| TDM V1/V2 native on gfx1250, pad false | Correct (18/18) |
| `comp_async` native on gfx1250, fully padded | Correct |
| clang-format-18 on the C++ headers | Clean |
| Dispatcher unit tests (targeted gfx1250 + new tests) | 109 passed |
| tile_engine gfx1250 tests | 92/92 passed |
| ArchFilter vs tile_engine gate equivalence (gfx1250 enumeration) |
Identical |
| gfx1250 build + `--verify` of every comp_async / comp_tdm /
comp_tdm_v2 default_config instance (7323 kernels x 9 shapes) | 0 wrong
results; the LDS gate rejects exactly the 90 instances that fail to
link, 0 false rejects |
| Generated arch specs reproducible from `arch_specs.json` | Identical |
| ruff check (F, E9) on the changed Python | Clean |
| gfx1250 end-to-end bridge run (TE + dispatcher, universal / batched /
multi-ABD) | 0 failures. TE 1188 pass / 1102 skip; dispatcher universal
620/840 and batched 94/96 verified, the rest are runtime-unsupported
skips; multi-ABD 18/22 with the 4 negative cases correctly rejected;
codegen gate vs emitted kernels and ArchFilter vs codegen: 0 mismatches
|
| gfx1250 kernel-fix gtests | TDM sub-tile 24/24, comp_async WMMA 14/14
(incl. Persistent), persistent WMMA 14/14 |
| gfx1250 tile_engine re-run of the reported configs (bf16 `rcr`, 5
shapes incl. 5608x4096x2048) | comp_async 40/40 pass (20 with
`persistent=true`); comp_tdm / comp_tdm_v2 72/72 pass, including
128x128x256 and 128x256x128 at 2x2x1 |
| Mutation check (gfx1250, old kernel code restored) | Persistent cases
fail without the comp_async fix; all sub-tile 2 cases fail without the
block GEMM fix; controls pass |
| gfx950 regression (comp_async, eight-waves, persistent async, MX async
rcr/rrr/crr/ccr) | All pass |

---------

Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
ThruptiRajLakshmanaGowda added a commit that referenced this pull request Sep 28, 2026
Re-adds the CI lanes that were reverted in commit 33ba206.
Changes are CI infrastructure only (Dockerfile/Jenkinsfile/ck.groovy);
Python dispatcher source is left at develop state to preserve
Ozturk's gfx1250 work (#11043, #10934, #12216).

- Dockerfile: re-pin ml_dtypes==0.6.0 for consistent fp8 OCP/FNUZ
  reference values across gfx942 (FNUZ) and gfx950 (OCP)
- Jenkinsfile: restore RUN_DISPATCHER_CORRECTNESS_TESTS and
  RUN_DISPATCHER_PERF_TESTS boolean params, dispatcher declarative
  stage, and nightly cron entry at 22:00 UTC
- ck.groovy: restore _dispatcherPerfCmd, runDispatcherPerfTests,
  dispatcherGemmVariantsFor, dispatcherSweepDtypesFor,
  dispatcherVariantCmd, _dispatcherCorrectnessBaseCmd,
  runDispatcherCorrectnessTests, runDispatcherTests (573 lines)

Correctness lane covers 13+ operators on gfx942/gfx950 with
arch-aware fp8 encoding; bquant/mx_gemm gfx950-only; run_ok
wrapper ensures exit-77 skips do not abort the chain.

JIRA ID : AICK-2146

Co-Authored-By: Claude <noreply@anthropic.com>
illsilin pushed a commit that referenced this pull request Oct 2, 2026
ISSUE ID : AICK-2171

## Motivation

PR #10897 added dispatcher correctness and performance CI lanes for
gfx942/gfx950 but was reverted in #12392 due to unrelated gfx1250
conflicts
in other files. This re-lands the CI infrastructure only — Dockerfile,
Jenkinsfile, and ck.groovy — leaving the Python dispatcher source at the
develop state to preserve Ozturk's gfx1250 work (#11043, #10934,
#12216).

## Technical Details

**Re-landed from #10897:**
- `Dockerfile`: re-pin `ml_dtypes==0.6.0` for consistent fp8 OCP/FNUZ
  reference values across gfx942 (FNUZ) and gfx950 (OCP)
- `Jenkinsfile`: restore `RUN_DISPATCHER_CORRECTNESS_TESTS` and
  `RUN_DISPATCHER_PERF_TESTS` boolean params (default: off), the
  declarative dispatcher stage, and nightly cron entry at 22:00 UTC
- `groovy/vars/ck.groovy`: restore `runDispatcherTests`,
  `runDispatcherCorrectnessTests`, `runDispatcherPerfTests` and helpers
  (573 lines) — correctness lane covers 13+ operators on gfx942/gfx950
  with arch-aware fp8 encoding; bquant/mx_gemm gfx950-only

**New fixes over #10897:**
- `dispatcher/CMakeLists.txt`: add `HIP` to project languages and
  `set_source_files_properties(LANGUAGE HIP)` for `fmha_dispatcher.cpp`
  and `fmha_registry.cpp` — required for the standalone cmake build used
by the correctness lane (these files use `<<<>>>` kernel launch syntax)
- `ck.groovy` correctness cmd: point cmake at `-S ../dispatcher -B
  ../dispatcher/build` and ninja at `-C ../dispatcher/build` so
  `libck_tile_dispatcher.a` lands in `dispatcher/build/` where
  `dispatcher_common.get_build_dir()` expects it
- `tile_engine/sampling/op_weights.json`: register
`contraction_multi_abd`
  with weight `0.0625` — the perf lane activates it via
`CONTRACTION_MULTI_ABD_DATATYPE=fp16` and `allocate_budget()` hard-fails
  in strict mode for any unregistered active op

## Test Plan

- `RUN_DISPATCHER_CORRECTNESS_TESTS=true` on gfx942 and gfx950: cmake +
  ninja build of `ck_tile_dispatcher`, then 13+ operator sweeps via
  `test_gemm_search_space.py` and per-op correctness scripts
- `RUN_DISPATCHER_PERF_TESTS=true` on gfx942 and gfx950: tile_engine AOT
  build of all dispatcher-bridged operators, then benchmark drivers for
  each

## Test Result

CI triggered on this branch — results pending.

## Submission Checklist

- [x] CI infrastructure only; Python dispatcher source unchanged from
develop
- [x] `defaultValue: false` for both dispatcher params (opt-in, not
opt-out)
- [x] `contraction_multi_abd` added to `op_weights.json`
- [x] HIP language fix in `dispatcher/CMakeLists.txt` for standalone
build

---------

Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Muhammed  Emin Ozturk <Muhammed.Ozturk@amd.com>
Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants