From 1764c77fb299e8f29b877111b2b04b1bd298dd09 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Fri, 26 Sep 2025 07:51:29 +0000 Subject: [PATCH] Enable running multiple GEMM batches of merged conv groups. --- ...uped_convolution_backward_weight_kernel.hpp | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp index d1c8452c5e..51d3b891f3 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp @@ -89,8 +89,8 @@ struct GroupedConvBwdWeightKernelArgs group_stride_a = args.K_ * NumGroupsPerBatch; // A: Out NWGK group_stride_b = args.C_ * NumGroupsPerBatch; // B: In NWGC - group_stride_c = args.K_ * args.C_ * // C: Wei GKXC - std::accumulate(args.filter_spatial_lengths_.begin(), + group_stride_c = args.K_ * args.C_ * NumGroupsPerBatch // C: Wei GKXC + * std::accumulate(args.filter_spatial_lengths_.begin(), args.filter_spatial_lengths_.end(), 1, std::multiplies()); @@ -98,7 +98,7 @@ struct GroupedConvBwdWeightKernelArgs GemmM = a_grid_desc_m_k.get_length(number<0>{}); GemmN = b_grid_desc_n_k.get_length(number<0>{}); GemmK = a_grid_desc_m_k.get_length(number<1>{}); - GemmBatch = std::max(static_cast(1), args.G_ / NumGroupsPerBatch); + GemmBatch = integer_divide_ceil(args.G_, NumGroupsPerBatch); if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) { @@ -175,8 +175,8 @@ struct GroupedConvBwdWeightKernelArgs group_stride_a = args.K_ * NumGroupsPerBatch; // A: Out NHWGK group_stride_b = args.C_ * NumGroupsPerBatch; // B: In NHWGC - group_stride_c = args.K_ * args.C_ * // C: Wei GKYXC - std::accumulate(args.filter_spatial_lengths_.begin(), + group_stride_c = args.K_ * args.C_ * NumGroupsPerBatch // C: Wei GKYXC + * std::accumulate(args.filter_spatial_lengths_.begin(), args.filter_spatial_lengths_.end(), 1, std::multiplies()); @@ -184,7 +184,7 @@ struct GroupedConvBwdWeightKernelArgs GemmM = a_grid_desc_m_k.get_length(number<0>{}); GemmN = b_grid_desc_n_k.get_length(number<0>{}); GemmK = a_grid_desc_m_k.get_length(number<1>{}); - GemmBatch = std::max(static_cast(1), args.G_ / NumGroupsPerBatch); + GemmBatch = integer_divide_ceil(args.G_, NumGroupsPerBatch); if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) { @@ -268,8 +268,8 @@ struct GroupedConvBwdWeightKernelArgs group_stride_a = args.K_ * NumGroupsPerBatch; // A: Out NDHWGK group_stride_b = args.C_ * NumGroupsPerBatch; // B: In NDHWGC - group_stride_c = args.K_ * args.C_ * // C: Wei GKZYXC - std::accumulate(args.filter_spatial_lengths_.begin(), + group_stride_c = args.K_ * args.C_ * NumGroupsPerBatch // C: Wei GKZYXC + * std::accumulate(args.filter_spatial_lengths_.begin(), args.filter_spatial_lengths_.end(), 1, std::multiplies()); @@ -277,7 +277,7 @@ struct GroupedConvBwdWeightKernelArgs GemmM = a_grid_desc_m_k.get_length(number<0>{}); GemmN = b_grid_desc_n_k.get_length(number<0>{}); GemmK = a_grid_desc_m_k.get_length(number<1>{}); - GemmBatch = std::max(static_cast(1), args.G_ / NumGroupsPerBatch); + GemmBatch = integer_divide_ceil(args.G_, NumGroupsPerBatch); if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) {