ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 17:51:40 +00:00
Code Issues Packages Projects Releases Wiki Activity

Enable running multiple GEMM batches of merged conv groups.

Browse Source
This commit is contained in:
Ville Pietilä
2025-09-26 07:51:29 +00:00
parent b864c077ed
commit 1764c77fb2
1 changed files with 9 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp
View File

@@ -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<index_t>());
@@ -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<long>(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<index_t>());
@@ -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<long>(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<index_t>());
@@ -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<long>(1), args.G_ / NumGroupsPerBatch);
GemmBatch = integer_divide_ceil(args.G_, NumGroupsPerBatch);
if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{