ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-07 15:47:05 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add more gfx942 instances.

Browse Source
This commit is contained in:
Ville Pietilä
2025-10-22 13:16:00 +00:00
parent 60db94037e
commit fbbadd4315
2 changed files with 45 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

31
library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances.hpp
View File

@@ -37,13 +37,10 @@ using tile_grouped_conv_bwd_weight_bf16_instances = std::tuple<
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
@@ -53,19 +50,37 @@ using tile_grouped_conv_bwd_weight_bf16_instances = std::tuple<
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
#endif
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>
#endif
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 1, 1, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
// clang-format on
>;

29
library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_fp16_instances.hpp
View File

@@ -37,13 +37,10 @@ using tile_grouped_conv_bwd_weight_f16_instances = std::tuple<
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
@@ -53,19 +50,37 @@ using tile_grouped_conv_bwd_weight_f16_instances = std::tuple<
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
#endif
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 1, 1, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 1, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
// clang-format on
>;