fix(CK): Forward port convolutions update from standalone repo - #10245
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
There was a problem hiding this comment.
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
SkipGemmSizeChecktemplate flag toGridwiseGemmMultiD_xdl_cshuffle_v3and gate the existing logical GEMM-size validity check on it. - Update the grouped-convolution forward device op to instantiate the gridwise GEMM with
SkipGemmSizeCheck=truefor 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
|
Have you check the warning, better to add the test. |
@AmosLewis, the added functionality is already covered by existing tests. |
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
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
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