Skip to content

feat(dispatcher/codegen): fixed A/B vector widths for the default epilogue and comp_async - #13610

Merged
andriy-ca merged 2 commits into
developfrom
users/andriy/ck/gemm-fixed-widths
Oct 14, 2026
Merged

andriy-ca merged 2 commits into
developfrom
users/andriy/ck/gemm-fixed-widths

Conversation

@andriy-ca

@andriy-ca andriy-ca commented Oct 14, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : AICK-2558

Why

Problems whose contiguous extents are not multiples of a kernel's native vector width are served by kernels compiled with narrower fixed A/B/C widths. Two pipelines' worth of those kernels were never generated:

  • Default epilogue. resolve_gemm_vector_sizes refused every fixed-width request unless the epilogue was cshuffle. The fixed A/B widths live in the pipeline problem, so they apply with any epilogue; only storing C narrower than its native width needs CShuffle, because the default epilogue's C width is fixed by the warp-tile layout.
  • comp_async. It was missing from VECTOR_SIZE_PIPELINES. Its pipeline problem honours fixed A/B widths like the other universal pipelines, with one extra limit: A and B are copied global to LDS by amd_async_global_load_to_lds, which moves only 4, 8, 12 or 16 bytes, or 1 byte for 8-bit types.

What changes

dispatcher/codegen/codegen_common.py:

  • Fixed widths are accepted with the default epilogue when C keeps its native width; a narrowed C is refused there, and the CShuffle split check is scoped to the cshuffle epilogue, the only one that shuffles C.
  • comp_async is added to VECTOR_SIZE_PIPELINES, and A/B widths whose copy size that instruction cannot move are rejected. comp_async's existing gfx1250 rules (row-major A, column-major B, full padding) are unchanged.

Testing

  • New cases in dispatcher/tests/test_gemm_vector_sizes.py. The driver and codegen suites (test_gemm_vector_sizes.py, test_gfx1250_tdm_pipelines.py, test_gemm_driver_verify.py, test_codegen_common.py, and tile_engine's gfx1250 pipeline tests) give 353 passed, 1 skipped on develop plus this PR (develop: 351 passed, the same skip).
  • Validated on gfx1250 hardware against an fp32 host reference: default-epilogue kernels with narrowed A/B widths, and comp_async kernels with A/B widths 2 and 4.
  • ruff 0.14.0 (check and format) is clean on the changed lines.

Merge order and related PRs

One of five independent PRs for gfx1250 GEMM support in the dispatcher and its heuristics. Suggested merge order:

  1. fix(dispatcher/heuristics): --capture_hw refused gfx950 and gfx1250 nodes; gfx1250 pipeline names misparsed #13614 heuristics: kernel-name and rocminfo parsing (also fixes --capture_hw on gfx950)
  2. fix(dispatcher): gfx1250 GEMM sweep gates refused valid TDM problems and 8-wave compv3, and offered compv5 #13609 gfx1250 GEMM sweep gates
  3. feat(dispatcher/codegen): fixed A/B vector widths for the default epilogue and comp_async #13610 fixed vector widths for the default epilogue and comp_async (this PR)
  4. perf(tile_engine): compute the GEMM verify reference once per shape, not per kernel #13611 GEMM verify reference computed once per shape
  5. feat(dispatcher/heuristics): distinct encodings for the gfx1250 GEMM pipelines and the tdm epilogue #13612 heuristics: encodings for the gfx1250 GEMM pipelines

This PR is number 3. Each of the five applies to develop on its own and merges cleanly with each of the others in either order (checked with three-way merges), so the order is a suggestion, not a requirement. This PR and #13609 both edit codegen_common.py and test_gemm_vector_sizes.py, in different places.

Related PRs from earlier work on the dispatcher heuristics and the GEMM sweep driver:

🤖 Generated with Claude Code

@therock-pr-bot

therock-pr-bot Bot commented Oct 14, 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?

@therock-pr-bot

Copy link
Copy Markdown

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

@andriy-ca andriy-ca added the ci:skip Skip all CI builds/tests for this PR label Oct 14, 2026
@andriy-ca
andriy-ca requested a review from JiaLuo-CAN October 14, 2026 19:33
@andriy-ca
andriy-ca marked this pull request as ready for review October 14, 2026 19:34
@andriy-ca
andriy-ca requested a review from a team as a code owner October 14, 2026 19:34

@JiaLuo-CAN JiaLuo-CAN 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.

LGTM.

@andriy-ca
andriy-ca merged commit a0a5156 into develop Oct 14, 2026
87 of 88 checks passed
@andriy-ca
andriy-ca deleted the users/andriy/ck/gemm-fixed-widths branch October 14, 2026 21:20
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Oct 14, 2026
feat(dispatcher/codegen): fixed A/B vector widths for the
 default epilogue and comp_async (#13610)

JIRA ID : AICK-2558

## Why

Problems whose contiguous extents are not multiples of a kernel's native
vector width are served by kernels compiled with narrower fixed A/B/C
widths. Two pipelines' worth of those kernels were never generated:

- **Default epilogue.** `resolve_gemm_vector_sizes` refused every
fixed-width request unless the epilogue was cshuffle. The fixed A/B
widths live in the pipeline problem, so they apply with any epilogue;
only storing C narrower than its native width needs CShuffle, because
the default epilogue's C width is fixed by the warp-tile layout.
- **comp_async.** It was missing from `VECTOR_SIZE_PIPELINES`. Its
pipeline problem honours fixed A/B widths like the other universal
pipelines, with one extra limit: A and B are copied global to LDS by
`amd_async_global_load_to_lds`, which moves only 4, 8, 12 or 16 bytes,
or 1 byte for 8-bit types.

## What changes

`dispatcher/codegen/codegen_common.py`:

- Fixed widths are accepted with the default epilogue when C keeps its
native width; a narrowed C is refused there, and the CShuffle split
check is scoped to the cshuffle epilogue, the only one that shuffles C.
- comp_async is added to `VECTOR_SIZE_PIPELINES`, and A/B widths whose
copy size that instruction cannot move are rejected. comp_async's
existing gfx1250 rules (row-major A, column-major B, full padding) are
unchanged.
@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
feat(dispatcher/codegen): fixed A/B vector widths for the
 default epilogue and comp_async (#13610)

JIRA ID : AICK-2558

## Why

Problems whose contiguous extents are not multiples of a kernel's native
vector width are served by kernels compiled with narrower fixed A/B/C
widths. Two pipelines' worth of those kernels were never generated:

- **Default epilogue.** `resolve_gemm_vector_sizes` refused every
fixed-width request unless the epilogue was cshuffle. The fixed A/B
widths live in the pipeline problem, so they apply with any epilogue;
only storing C narrower than its native width needs CShuffle, because
the default epilogue's C width is fixed by the warp-tile layout.
- **comp_async.** It was missing from `VECTOR_SIZE_PIPELINES`. Its
pipeline problem honours fixed A/B widths like the other universal
pipelines, with one extra limit: A and B are copied global to LDS by
`amd_async_global_load_to_lds`, which moves only 4, 8, 12 or 16 bytes,
or 1 byte for 8-bit types.

## What changes

`dispatcher/codegen/codegen_common.py`:

- Fixed widths are accepted with the default epilogue when C keeps its
native width; a narrowed C is refused there, and the CShuffle split
check is scoped to the cshuffle epilogue, the only one that shuffles C.
- comp_async is added to `VECTOR_SIZE_PIPELINES`, and A/B widths whose
copy size that instruction cannot move are rejected. comp_async's
existing gfx1250 rules (row-major A, column-major B, full padding) are
unchanged.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants