Skip to content

Grouped conv bwd data skip B LDS #2774

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 4 commits into from
Closed

Grouped conv bwd data skip B LDS #2774

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
dd310e4
Grouped conv bwd data skip B LDS
bartekxk Sep 1, 2025
0f39363
Merge branch 'develop' of github.com:ROCm/composable_kernel into bark…
bartekxk Sep 1, 2025
957639a
instances
bartekxk Sep 2, 2025
e688dcf
fixes
bartekxk Sep 2, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
223 changes: 173 additions & 50 deletions ...sor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp
View file
Open in desktop

Large diffs are not rendered by default.

791 changes: 791 additions & 0 deletions include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_skip_b_lds_multiple_d_cshuffle.hpp
View file
Open in desktop

Large diffs are not rendered by default.

297 changes: 150 additions & 147 deletions ...peration_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_instance.hpp
View file
Open in desktop

Large diffs are not rendered by default.

24 changes: 24 additions & 0 deletions ...ry/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -111,6 +111,10 @@ struct DeviceOperationInstanceFactory<
add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f16_instances(op_ptrs);
add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f16_16_16_instances(
op_ptrs);
add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f16_instances(
op_ptrs);
add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f16_16_16_instances(
op_ptrs);
}
#endif
#ifdef CK_ENABLE_FP32
Expand All @@ -121,6 +125,10 @@ struct DeviceOperationInstanceFactory<
add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f32_instances(op_ptrs);
add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f32_16_16_instances(
op_ptrs);
add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f32_instances(
op_ptrs);
add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f32_16_16_instances(
op_ptrs);
}
#endif
#ifdef CK_ENABLE_BF16
Expand All @@ -132,6 +140,10 @@ struct DeviceOperationInstanceFactory<
op_ptrs);
add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_bf16_16_16_instances(
op_ptrs);
add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_bf16_instances(
op_ptrs);
add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_bf16_16_16_instances(
op_ptrs);
}
#endif
}
Expand Down Expand Up @@ -251,6 +263,10 @@ struct DeviceOperationInstanceFactory<
op_ptrs);
add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_f16_16_16_instances(
op_ptrs);
add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f16_instances(
op_ptrs);
add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f16_16_16_instances(
op_ptrs);
}
#endif
#if defined CK_ENABLE_FP16 && defined CK_ENABLE_FP8 && defined CK_ENABLE_BF8
Expand All @@ -271,6 +287,10 @@ struct DeviceOperationInstanceFactory<
op_ptrs);
add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_f32_16_16_instances(
op_ptrs);
add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f32_instances(
op_ptrs);
add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f32_16_16_instances(
op_ptrs);
}
#endif
#ifdef CK_ENABLE_BF16
Expand All @@ -282,6 +302,10 @@ struct DeviceOperationInstanceFactory<
op_ptrs);
add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_16_16_instances(
op_ptrs);
add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_instances(
op_ptrs);
add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_16_16_instances(
op_ptrs);
}
#endif
}
Expand Down
168 changes: 168 additions & 0 deletions ...nclude/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data_xdl.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,6 +83,34 @@ void add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f16_16_16_instance
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<2,
NHWGK,
GKYXC,
Empty_Tuple,
NHWGC,
F16,
F16,
Empty_Tuple,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f16_16_16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<2,
NHWGK,
GKYXC,
Empty_Tuple,
NHWGC,
F16,
F16,
Empty_Tuple,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances);
#endif
#ifdef CK_ENABLE_FP32
void add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f32_instances(
Expand Down Expand Up @@ -112,6 +140,34 @@ void add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_f32_16_16_instance
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f32_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<2,
NHWGK,
GKYXC,
Empty_Tuple,
NHWGC,
F32,
F32,
Empty_Tuple,
F32,
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_f32_16_16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<2,
NHWGK,
GKYXC,
Empty_Tuple,
NHWGC,
F32,
F32,
Empty_Tuple,
F32,
PassThrough,
PassThrough,
PassThrough>>>& instances);
#endif
#ifdef CK_ENABLE_BF16
void add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_bf16_instances(
Expand Down Expand Up @@ -141,6 +197,34 @@ void add_device_grouped_conv2d_bwd_data_xdl_nhwgk_gkyxc_nhwgc_bf16_16_16_instanc
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_bf16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<2,
NHWGK,
GKYXC,
Empty_Tuple,
NHWGC,
BF16,
BF16,
Empty_Tuple,
BF16,
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv2d_bwd_data_skip_b_lds_xdl_nhwgk_gkyxc_nhwgc_bf16_16_16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<2,
NHWGK,
GKYXC,
Empty_Tuple,
NHWGC,
BF16,
BF16,
Empty_Tuple,
BF16,
PassThrough,
PassThrough,
PassThrough>>>& instances);
#endif

#ifdef CK_ENABLE_FP16
Expand Down Expand Up @@ -393,6 +477,34 @@ void add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_f16_16_16_insta
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<3,
NDHWGK,
GKZYXC,
Empty_Tuple,
NDHWGC,
F16,
F16,
Empty_Tuple,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f16_16_16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<3,
NDHWGK,
GKZYXC,
Empty_Tuple,
NDHWGC,
F16,
F16,
Empty_Tuple,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances);
#endif
#ifdef CK_ENABLE_FP32
void add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_f32_instances(
Expand Down Expand Up @@ -422,6 +534,34 @@ void add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_f32_16_16_insta
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f32_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<3,
NDHWGK,
GKZYXC,
Empty_Tuple,
NDHWGC,
F32,
F32,
Empty_Tuple,
F32,
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_f32_16_16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<3,
NDHWGK,
GKZYXC,
Empty_Tuple,
NDHWGC,
F32,
F32,
Empty_Tuple,
F32,
PassThrough,
PassThrough,
PassThrough>>>& instances);
#endif
#ifdef CK_ENABLE_BF16
void add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_instances(
Expand Down Expand Up @@ -451,6 +591,34 @@ void add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_16_16_inst
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<3,
NDHWGK,
GKZYXC,
Empty_Tuple,
NDHWGC,
BF16,
BF16,
Empty_Tuple,
BF16,
PassThrough,
PassThrough,
PassThrough>>>& instances);

void add_device_grouped_conv3d_bwd_data_skip_b_lds_xdl_ndhwgk_gkzyxc_ndhwgc_bf16_16_16_instances(
std::vector<std::unique_ptr<DeviceGroupedConvBwdDataMultipleD<3,
NDHWGK,
GKZYXC,
Empty_Tuple,
NDHWGC,
BF16,
BF16,
Empty_Tuple,
BF16,
PassThrough,
PassThrough,
PassThrough>>>& instances);
#endif
#if defined CK_ENABLE_FP16 && defined CK_ENABLE_FP8 && defined CK_ENABLE_BF8
void add_device_grouped_conv3d_bwd_data_xdl_ndhwgk_gkzyxc_ndhwgc_input_f16_comp_bf8f8_instances(
Expand Down
Loading
Loading