Skip to content

[CK Tile] Fix V6 pipeline applicability and split-image initialization - #7936

Merged
vpietila-amd merged 7 commits into
developfrom
users/vpietila/ck/fix-v6-instances-in-ck-tile-dispatcher
Jun 3, 2026
Merged

[CK Tile] Fix V6 pipeline applicability and split-image initialization#7936
vpietila-amd merged 7 commits into
developfrom
users/vpietila/ck/fix-v6-instances-in-ck-tile-dispatcher

Conversation

@vpietila-amd

Copy link
Copy Markdown
Contributor

Motivation

After adding code generation via CK Tile Dispatcher, some fwd and bwd weight tests for CK Tile convolutions are failing. This PR introduced correct applicability checks and fixes the split-image parameter initialization such that non-applicable instances are not invoked during test execution and split-image instances are correctly initialized.

Technical Details

Investigation revealed two distinct problems

  1. For bwd weight, the compute V3 uses prefetch of 3 distinct tiles, which works incorrectly when the number of K-slices addressed by the workgroup is 1. This occurs when a large split-K value is used for a problem that results in a small Gemm-K value.
  2. For fwd direction, the current CK Profiler/test infrastructure doesn't initialize the split-image parameters for instance where split-image is enable. Uninitialized split-image values result in non-deterministic behavior where the tests might randomly fail.

Fixed problem 1. by adding a check in IsSupportedArgument that marks the instance invalid if the num_loops = ceil(GemmK / (k_batch * KPerBlock)) < 4 for V6 pipeline kernel instances. The check is compile-time eliminated for other kernels.

Fixed problem 2. by adding initialization of split-image parameters when split-image is enabled. The default initialization corresponds to full image with no split, i.e., the number of splits is 1 and it has the size of the full image.

Added unit tests for the added logic.

Test Plan

Running the following test suites cover the logic added in this PR

  • test_grouped_convnd_fwd_tile
  • test_ck_tile_grouped_conv_fwd
  • test_grouped_convnd_bwd_weight_tile
  • test_ck_tile_grouped_conv_bwd_weight

All test suites above are included in the automated test runs.

Test Result

Submission Checklist

@vpietila-amd
vpietila-amd enabled auto-merge (squash) June 2, 2026 11:09
@vpietila-amd
vpietila-amd merged commit 3dc91e6 into develop Jun 3, 2026
50 of 52 checks passed
@vpietila-amd
vpietila-amd deleted the users/vpietila/ck/fix-v6-instances-in-ck-tile-dispatcher branch June 3, 2026 08:38
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Jun 3, 2026
[CK Tile] Fix V6 pipeline applicability and split-image
 initialization (#7936)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Motivation

After adding code generation via CK Tile Dispatcher, some fwd and bwd
weight tests for CK Tile convolutions are failing. This PR introduced
correct applicability checks and fixes the split-image parameter
initialization such that non-applicable instances are not invoked during
test execution and split-image instances are correctly initialized.

## Technical Details

Investigation revealed two distinct problems

1. For bwd weight, the compute V3 uses prefetch of 3 distinct tiles,
which works incorrectly when the number of K-slices addressed by the
workgroup is 1. This occurs when a large split-K value is used for a
problem that results in a small Gemm-K value.
2. For fwd direction, the current CK Profiler/test infrastructure
doesn't initialize the split-image parameters for instance where
split-image is enable. Uninitialized split-image values result in
non-deterministic behavior where the tests might randomly fail.

Fixed problem 1. by adding a check in `IsSupportedArgument` that marks
the instance invalid if the `num_loops = ceil(GemmK / (k_batch *
KPerBlock)) < 4` for V6 pipeline kernel instances. The check is
compile-time eliminated for other kernels.

Fixed problem 2. by adding initialization of split-image parameters when
split-image is enabled. The default initialization corresponds to full
image with no split, i.e., the number of splits is 1 and it has the size
of the full image.

Added unit tests for the added logic.

## Test Plan

Running the following test suites cover the logic added in this PR
- test_grouped_convnd_fwd_tile
- test_ck_tile_grouped_conv_fwd
- test_grouped_convnd_bwd_weight_tile
- test_ck_tile_grouped_conv_bwd_weight

All test suites above are included in the automated test runs.

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
shumway pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 18, 2026
[CK Tile] Fix V6 pipeline applicability and split-image initialization

## Motivation

After adding code generation via CK Tile Dispatcher, some fwd and bwd
weight tests for CK Tile convolutions are failing. This PR introduced
correct applicability checks and fixes the split-image parameter
initialization such that non-applicable instances are not invoked during
test execution and split-image instances are correctly initialized.

## Technical Details

Investigation revealed two distinct problems

1. For bwd weight, the compute V3 uses prefetch of 3 distinct tiles,
which works incorrectly when the number of K-slices addressed by the
workgroup is 1. This occurs when a large split-K value is used for a
problem that results in a small Gemm-K value.
2. For fwd direction, the current CK Profiler/test infrastructure
doesn't initialize the split-image parameters for instance where
split-image is enable. Uninitialized split-image values result in
non-deterministic behavior where the tests might randomly fail.

Fixed problem 1. by adding a check in `IsSupportedArgument` that marks
the instance invalid if the `num_loops = ceil(GemmK / (k_batch *
KPerBlock)) < 4` for V6 pipeline kernel instances. The check is
compile-time eliminated for other kernels.

Fixed problem 2. by adding initialization of split-image parameters when
split-image is enabled. The default initialization corresponds to full
image with no split, i.e., the number of splits is 1 and it has the size
of the full image.

Added unit tests for the added logic.

## Test Plan

Running the following test suites cover the logic added in this PR
- test_grouped_convnd_fwd_tile
- test_ck_tile_grouped_conv_fwd 
- test_grouped_convnd_bwd_weight_tile 
- test_ck_tile_grouped_conv_bwd_weight

All test suites above are included in the automated test runs.  

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---------

Co-authored-by: Ville Pietilä <>
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.

2 participants