Skip to content

fix(CK): Forward port convolutions update from standalone repo - #10245

Merged
geyyer merged 3 commits into
developfrom
users/rogeyyer/ck/aick-1818
Aug 6, 2026
Merged

fix(CK): Forward port convolutions update from standalone repo#10245
geyyer merged 3 commits into
developfrom
users/rogeyyer/ck/aick-1818

Conversation

@geyyer

@geyyer geyyer commented Jul 31, 2026

Copy link
Copy Markdown
Contributor

Motivation

Forward-port grouped-conv-fwd perf-regression fix (im2col GEMM-size 2GB check) to rocm-libraries mainline

Technical Details

Port this commit ROCm/composable_kernel@248f155#diff-d921922979d6cfc604769c3bc231107ae574375875ecce04bcb1065d43a5ce2e

Introduced by a mainline change for grouped-conv large-tensor / global-load-store CK enablement, which added a LargeTensors template flag to GridwiseGemmMultiD_xdl_cshuffle_v3 and reworked its CheckValidity, but did not exempt the convolution im2col-view path from the pre-existing logical-GEMM-size 2GB check

Test Plan

Regular CI

Test Result

CI should pass

Submission Checklist

JIRA ID : AICK-1818

@therock-pr-bot

therock-pr-bot Bot commented Jul 31, 2026

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

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/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp, projects/composablekernel/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp; no test file found
🔎 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

therock-pr-bot Bot commented Jul 31, 2026

Copy link
Copy Markdown

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

@geyyer
geyyer marked this pull request as ready for review August 4, 2026 14:59
@geyyer
geyyer requested a review from a team as a code owner August 4, 2026 14:59
Copilot AI lite review requested due to automatic review settings August 4, 2026 14:59

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.

Pull request overview

Forward-ports a composable-kernel fix to avoid incorrectly rejecting grouped-convolution implicit-GEMM (im2col-view) problems due to the logical GEMM matrix size exceeding the existing 2 GiB check, even when the actual underlying tensor footprint fits within int32 addressing.

Changes:

  • Add a SkipGemmSizeCheck template flag to GridwiseGemmMultiD_xdl_cshuffle_v3 and gate the existing logical GEMM-size validity check on it.
  • Update the grouped-convolution forward device op to instantiate the gridwise GEMM with SkipGemmSizeCheck=true for the convolution/implicit-GEMM path (relying on existing descriptor-span checks instead).

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.

File Description
projects/composablekernel/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp Adds SkipGemmSizeCheck and uses it to bypass the logical 2 GiB GEMM-size check when appropriate.
projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp Enables skipping the logical GEMM-size check for convolution implicit-GEMM instantiations (while keeping descriptor-size validation).
Suppressed comments (2)

projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp:616

  • If you rename the macro to a uniquely-prefixed name, the instantiation here should be updated to use the new identifier to avoid relying on a generic global macro name.
        GridwiseGemmV3TemplateParams,

projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp:627

  • When using a uniquely-prefixed macro name, make sure the corresponding #undef matches the renamed identifier; otherwise the generic macro name remains a collision risk in translation units that include this header.
#undef GridwiseGemmV3TemplateParams

@AmosLewis

Copy link
Copy Markdown
Contributor

Have you check the warning, better to add the test.

@geyyer

geyyer commented Aug 6, 2026

Copy link
Copy Markdown
Contributor Author

Have you check the warning, better to add the test.

@AmosLewis, the added functionality is already covered by existing tests.

@geyyer
geyyer merged commit f89f75f into develop Aug 6, 2026
46 checks passed
@geyyer
geyyer deleted the users/rogeyyer/ck/aick-1818 branch August 6, 2026 14:47
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 6, 2026
fix(CK): Forward port convolutions update from standalone
 repo (#10245)

## Motivation

Forward-port grouped-conv-fwd perf-regression fix (im2col GEMM-size 2GB
check) to rocm-libraries mainline

## Technical Details
Port this commit
248f155#diff-d921922979d6cfc604769c3bc231107ae574375875ecce04bcb1065d43a5ce2e

Introduced by a mainline change for grouped-conv large-tensor /
global-load-store CK enablement, which added a LargeTensors template
flag to GridwiseGemmMultiD_xdl_cshuffle_v3 and reworked its
CheckValidity, but did not exempt the convolution im2col-view path from
the pre-existing logical-GEMM-size 2GB check

## Test Plan

Regular CI

## Test Result

CI should pass

## Submission Checklist

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

JIRA ID : AICK-1818
shumway pushed a commit to ROCm/composable_kernel that referenced this pull request Aug 18, 2026
fix(CK): Forward port convolutions update from standalone repo

## Motivation

Forward-port grouped-conv-fwd perf-regression fix (im2col GEMM-size 2GB
check) to rocm-libraries mainline

## Technical Details
Port this commit
248f155#diff-d921922979d6cfc604769c3bc231107ae574375875ecce04bcb1065d43a5ce2e

Introduced by a mainline change for grouped-conv large-tensor /
global-load-store CK enablement, which added a LargeTensors template
flag to GridwiseGemmMultiD_xdl_cshuffle_v3 and reworked its
CheckValidity, but did not exempt the convolution im2col-view path from
the pre-existing logical-GEMM-size 2GB check

## Test Plan

Regular CI

## Test Result

CI should pass

## Submission Checklist

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

JIRA ID : AICK-1818
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.

4 participants