Skip to content

Enable opt-in wider gfx1250 FP8 WMMA - #14391

Merged
illsilin merged 1 commit into
developfrom
users/mbidlek/ck/fp8tdm_fix
Oct 27, 2026
Merged

illsilin merged 1 commit into
developfrom
users/mbidlek/ck/fp8tdm_fix

Conversation

@bidlekm

@bidlekm bidlekm commented Oct 26, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Enable CK Tile configurations on gfx1250 to select the native FP8/BF8 16x16x128 WMMA instruction instead of the default 16x16x64 instruction.

This must remain opt-in because K128 requires a compatible block K tile and may increase register pressure or reduce occupancy. Smaller tiles can perform better in some cases:

FP8 problem Pipeline K64 block tile K128 block tile K64 advantage
512x512x512 Basic V1 11.803 TFLOP/s 9.401 TFLOP/s +25.6%
512x512x512 TDM 12.474 TFLOP/s 11.822 TFLOP/s +5.5%
4096x4096x4096 TDM 1336.35 TFLOP/s 1210.67 TFLOP/s +10.4%

These results compare block tile sizes rather than only instruction width, demonstrating why configurations should retain explicit control over K128 selection.

Closes #14391

Technical Details

Adds a WiderInstruction template parameter to get_k_warp_tile().

  • Defaults to false, preserving existing behavior.
  • When enabled for gfx1250 FP8/BF8 with a 16x16 warp tile, selects K128 instead of K64.
  • Does not change FP16, FP32, integer, or non-gfx1250 selection.
  • Configurations must opt in only when their block K supports the K128 atom.

Example:

get_k_warp_tile<fp8_t, 16, false, true>()

Test Plan

  • Compile-time verification that default gfx1250 FP8 still selects K64.
  • Compile-time verification that opt-in FP8 and BF8 select K128.
  • Compile-time verification that opt-in FP16 remains K32.
  • Compile the selector assertions with the gfx1250 toolchain.
  • Run clang-format-18.

Benchmark FP8 RCR GEMM with physical 16x16x128 WMMA instructions. Each result is the median of three trials with 20 warmups and 100 measured iterations.

Test Result

All compile-time checks and formatting checks passed.

Shape Block / warp tile Reference K128 Gain Latency
2304x512x16384 64x64x256 / 16x16x128 464.9 550.3 +18.4% 70.2 us
512x2304x16384 64x64x256 / 16x16x128 452.2 547.0 +21.0% 70.7 us
1024x5120x8192 64x128x128 / 32x32x128 1389.5 1805.6 +29.9% 47.6 us
5120x1024x8192 64x128x128 / 32x32x128 1516.9 1920.9 +26.6% 44.7 us
4096x4096x16384 256x128x128 / 32x32x128 1942.4 2423.9 +24.8% 226.8 us
14336x2048x1024 128x64x128 / 32x32x128 1334.4 1453.7 +8.9% 41.4 us
2048x14336x1024 128x64x128 / 32x32x128 1275.0 1434.4 +12.5% 41.9 us

Submission Checklist

@bidlekm bidlekm self-assigned this Oct 26, 2026
@bidlekm
bidlekm requested a review from a team as a code owner October 26, 2026 15:21
Copilot AI balanced review requested due to automatic review settings October 26, 2026 15:21

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.

🟢 Approval recommended

The narrowly scoped change preserves default behavior and opts into existing K128 backend support, with no blocking issues identified.

No open issues 🎉

What changed in this PR

Adds opt-in wider FP8/BF8 WMMA selection for CK Tile on gfx1250 while preserving existing defaults.

Changes:

  • Adds a WiderInstruction template parameter, defaulting to false.
  • Selects K128 for eligible FP8/BF8 tiles when enabled, leaving other selection paths unchanged.
File Description
projects/​composablekernel/​include/​ck_tile/​ops/​gemm/​pipeline/​tile_gemm_shape.hpp Adds default-disabled K128 selection to get_k_warp_tile().

Review effort: Balanced

@therock-pr-bot

therock-pr-bot Bot commented Oct 26, 2026 •

Copy link
Copy Markdown

⏳ Waiting on Required Checks

⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/composablekernel/include/ck_tile/ops/gemm/pipeline/tile_gemm_shape.hpp; no test file found
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⏳ Required checks are still running or have not started yet.
This PR is not ready for review until they complete.

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

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Oct 26, 2026 •

Copy link
Copy Markdown

⏳ Waiting for required checks to start or finish:

  • ⏳ pre-commit

@illsilin illsilin added the ci:skip Skip all CI builds/tests for this PR label Oct 26, 2026
@illsilin
illsilin merged commit 2deccbd into develop Oct 27, 2026
155 of 185 checks passed
@illsilin
illsilin deleted the users/mbidlek/ck/fp8tdm_fix branch October 27, 2026 02:13
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Oct 27, 2026
Enable opt-in wider gfx1250 FP8 WMMA

## Motivation

Enable CK Tile configurations on gfx1250 to select the native FP8/BF8
`16x16x128` WMMA instruction instead of the default `16x16x64`
instruction.

This must remain opt-in because K128 requires a compatible block K tile
and may increase register pressure or reduce occupancy. Smaller tiles
can perform better in some cases:

| FP8 problem | Pipeline | K64 block tile | K128 block tile | K64
advantage |
|---|---|---:|---:|---:|
| `512x512x512` | Basic V1 | 11.803 TFLOP/s | 9.401 TFLOP/s | +25.6% |
| `512x512x512` | TDM | 12.474 TFLOP/s | 11.822 TFLOP/s | +5.5% |
| `4096x4096x4096` | TDM | 1336.35 TFLOP/s | 1210.67 TFLOP/s | +10.4% |

These results compare block tile sizes rather than only instruction
width, demonstrating why configurations should retain explicit control
over K128 selection.

## Technical Details

Adds a `WiderInstruction` template parameter to `get_k_warp_tile()`.

- Defaults to `false`, preserving existing behavior.
- When enabled for gfx1250 FP8/BF8 with a `16x16` warp tile, selects
K128 instead of K64.
- Does not change FP16, FP32, integer, or non-gfx1250 selection.
- Configurations must opt in only when their block K supports the K128
atom.

Example:

```cpp
get_k_warp_tile<fp8_t, 16, false, true>()
```

## Test Plan

- Compile-time verification that default gfx1250 FP8 still selects K64.
- Compile-time verification that opt-in FP8 and BF8 select K128.
- Compile-time verification that opt-in FP16 remains K32.
- Compile the selector assertions with the gfx1250 toolchain.
- Run `clang-format-18`.

Benchmark FP8 RCR GEMM with physical `16x16x128` WMMA instructions. Each
result is the median of three trials with 20 warmups and 100 measured
iterations.

## Test Result

All compile-time checks and formatting checks passed.

| Shape | Block / warp tile | Reference | K128 | Gain | Latency |
|---|---|---:|---:|---:|---:|
| `2304x512x16384` | `64x64x256` / `16x16x128` | 464.9 | **550.3** |
+18.4% | 70.2 us |
| `512x2304x16384` | `64x64x256` / `16x16x128` | 452.2 | **547.0** |
+21.0% | 70.7 us |
| `1024x5120x8192` | `64x128x128` / `32x32x128` | 1389.5 | **1805.6** |
+29.9% | 47.6 us |
| `5120x1024x8192` | `64x128x128` / `32x32x128` | 1516.9 | **1920.9** |
+26.6% | 44.7 us |
| `4096x4096x16384` | `256x128x128` / `32x32x128` | 1942.4 | **2423.9**
| +24.8% | 226.8 us |
| `14336x2048x1024` | `128x64x128` / `32x32x128` | 1334.4 | **1453.7** |
+8.9% | 41.4 us |
| `2048x14336x1024` | `128x64x128` / `32x32x128` | 1275.0 | **1434.4** |
+12.5% | 41.9 us |

The `32x32x128` logical tiles decompose into four physical `16x16x128`
WMMA operations. The reference column represents the previous benchmark
champion and is not a strict same-kernel K64/K128 comparison.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Co-authored-by: Marton Bidlek <marton.bidlek@streamhpc.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

shumway pushed a commit to ROCm/composable_kernel that referenced this pull request Nov 4, 2026
Enable opt-in wider gfx1250 FP8 WMMA

## Motivation

Enable CK Tile configurations on gfx1250 to select the native FP8/BF8
`16x16x128` WMMA instruction instead of the default `16x16x64`
instruction.

This must remain opt-in because K128 requires a compatible block K tile
and may increase register pressure or reduce occupancy. Smaller tiles
can perform better in some cases:

| FP8 problem | Pipeline | K64 block tile | K128 block tile | K64
advantage |
|---|---|---:|---:|---:|
| `512x512x512` | Basic V1 | 11.803 TFLOP/s | 9.401 TFLOP/s | +25.6% |
| `512x512x512` | TDM | 12.474 TFLOP/s | 11.822 TFLOP/s | +5.5% |
| `4096x4096x4096` | TDM | 1336.35 TFLOP/s | 1210.67 TFLOP/s | +10.4% |

These results compare block tile sizes rather than only instruction
width, demonstrating why configurations should retain explicit control
over K128 selection.

## Technical Details

Adds a `WiderInstruction` template parameter to `get_k_warp_tile()`.

- Defaults to `false`, preserving existing behavior.
- When enabled for gfx1250 FP8/BF8 with a `16x16` warp tile, selects
K128 instead of K64.
- Does not change FP16, FP32, integer, or non-gfx1250 selection.
- Configurations must opt in only when their block K supports the K128
atom.

Example:

```cpp
get_k_warp_tile<fp8_t, 16, false, true>()
```

## Test Plan

- Compile-time verification that default gfx1250 FP8 still selects K64.
- Compile-time verification that opt-in FP8 and BF8 select K128.
- Compile-time verification that opt-in FP16 remains K32.
- Compile the selector assertions with the gfx1250 toolchain.
- Run `clang-format-18`.

Benchmark FP8 RCR GEMM with physical `16x16x128` WMMA instructions. Each
result is the median of three trials with 20 warmups and 100 measured
iterations.

## Test Result

All compile-time checks and formatting checks passed.

| Shape | Block / warp tile | Reference | K128 | Gain | Latency |
|---|---|---:|---:|---:|---:|
| `2304x512x16384` | `64x64x256` / `16x16x128` | 464.9 | **550.3** |
+18.4% | 70.2 us |
| `512x2304x16384` | `64x64x256` / `16x16x128` | 452.2 | **547.0** |
+21.0% | 70.7 us |
| `1024x5120x8192` | `64x128x128` / `32x32x128` | 1389.5 | **1805.6** |
+29.9% | 47.6 us |
| `5120x1024x8192` | `64x128x128` / `32x32x128` | 1516.9 | **1920.9** |
+26.6% | 44.7 us |
| `4096x4096x16384` | `256x128x128` / `32x32x128` | 1942.4 | **2423.9**
| +24.8% | 226.8 us |
| `14336x2048x1024` | `128x64x128` / `32x32x128` | 1334.4 | **1453.7** |
+8.9% | 41.4 us |
| `2048x14336x1024` | `128x64x128` / `32x32x128` | 1275.0 | **1434.4** |
+12.5% | 41.9 us |

The `32x32x128` logical tiles decompose into four physical `16x16x128`
WMMA operations. The reference column represents the previous benchmark
champion and is not a strict same-kernel K64/K128 comparison.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Co-authored-by: Marton Bidlek <marton.bidlek@streamhpc.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ci:skip Skip all CI builds/tests for this PR Not ready to Review organization: streamhpc contributors from streamhpc project: composablekernel

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants