ROCm / ROCm/composable_kernel

[CK Tile] UniversalGemmKernel silently rejects D tensors whose layout differs from CLayout

Open
#3,766 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

Dominant language
C++
Stars
549
Forks
312
PR merge metrics
No merged PRs in 30d

Description

universal_gemm_kernel.hpp:581-584 rejects any D tensor whose layout is not CLayout — and does so silently: no log line, no assert message, the kernel simply does not run. Localizing it took a CK_TILE_LOGGING run plus a source read.

This blocks a natural and otherwise-supported composition: a weight-preshuffled GEMM plus a per-token/per-channel scale epilogue, expressed as two broadcast D tensors —

  • D0 = per-token A scale: an M-vector, ColumnMajor, stride 0 in N
  • D1 = per-channel B scale: an N-vector, RowMajor, stride 0 in M

D1 is accepted; D0 is not, because its layout differs from CLayout. Independently, cshuffle_epilogue.hpp:782 wraps every D window with the (M, N) distribution and has no handling for a column-layout D.

This is exactly the architecture CK's own legacy XDL path ships (GridwiseGemmMultiD with row/col scale Ds), so the operation is well-precedented; only the CK Tile universal path can't express it.

Asks, in increasing order of effort:

  1. make the rejection loud (a static_assert message or a log line naming the offending D index and layout);
  2. relax the check and add column-layout D distribution handling in the CShuffle epilogue;
  3. or, if RowColQuant is the intended route for this shape of problem, document that — it works (see the QuantGemmKernel RowColQuant path) and it is not obvious from the examples.

Context: found while wiring per-token/per-channel FP8 GEMMs on gfx950; the RowColQuant + weight-preshuffle route works and measured +8-11% over the flatmm path at three GEMM shapes, so the composition is worth having.

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with universal_gemm_kernel.hpp:581-584 and cshuffle_epilogue.hpp:782, then compare the legacy GridwiseGemmMultiD row/column scale handling and the QuantGemmKernel RowColQuant path. Use a CK_TILE_LOGGING run to observe the current rejection. Done means either the D layout is supported in the CShuffle epilogue, the rejection identifies the D index and layout, or the supported RowColQuant route is documented.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
backend, performance
Issue type
Feature
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
38/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.