From 7b4d9879daef8082232fd3379b7dfb90a22aa80a Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Bart=C5=82omiej=20Kocot?= <barkocot@amd.com>
Date: Tue, 4 Nov 2025 15:04:53 +0100
Subject: [PATCH] [CK TILE] Refactor Conv configs and Conv Elementwise (#3151)

* [CK TILE] Refactor Conv configs and Conv Elementwise

* fix

[ROCm/composable_kernel commit: 8681ced9629f6e952afa5b77c5f3549d60920efa]
---
 .../20_grouped_convolution/conv_configs.hpp   |   8 +-
 .../grouped_convolution_backward_data.cpp     |   6 +-
 ...uped_convolution_backward_data_invoker.hpp |  53 ++--
 ...ed_convolution_backward_weight_invoker.hpp |   9 +-
 ...tion_backward_weight_two_stage_invoker.hpp |   9 +-
 .../grouped_convolution_forward.cpp           |   6 +-
 ...grouped_convolution_forward_bias_clamp.cpp |   6 +-
 .../grouped_convolution_forward_invoker.hpp   |  57 ++--
 ...ouped_convolution_forward_large_tensor.cpp |   6 +-
 ...nvolution_forward_large_tensor_invoker.hpp | 248 ++++++++++--------
 ...n_grouped_convolution_bwd_data_example.inc |  16 +-
 .../run_grouped_convolution_fwd_example.inc   |  16 +-
 .../grouped_convolution_forward_kernel.hpp    |   7 +-
 .../utils/grouped_convolution_utils.hpp       |   2 -
 14 files changed, 230 insertions(+), 219 deletions(-)

diff --git a/example/ck_tile/20_grouped_convolution/conv_configs.hpp b/example/ck_tile/20_grouped_convolution/conv_configs.hpp
index c688215280..b9edf247cc 100644
--- a/example/ck_tile/20_grouped_convolution/conv_configs.hpp
+++ b/example/ck_tile/20_grouped_convolution/conv_configs.hpp
@@ -18,11 +18,7 @@ struct ConvConfigBase
     static constexpr bool kPadN = true;
     static constexpr bool kPadK = true;
 
-    static constexpr bool PermuteA = false;
-    static constexpr bool PermuteB = false;
-
-    static constexpr bool TransposeC            = false;
-    static constexpr bool UseStructuredSparsity = false;
+    static constexpr bool TransposeC = false;
 
     static constexpr ck_tile::index_t VectorSizeA = 4;
     static constexpr ck_tile::index_t VectorSizeB = 8;
@@ -34,8 +30,6 @@ struct ConvConfigBase
     static constexpr auto Scheduler                 = ck_tile::GemmPipelineScheduler::Intrawave;
     static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V3;
     static constexpr ck_tile::index_t NumWaveGroups = 1;
-    static constexpr bool Preshuffle                = false;
-    static constexpr bool TiledMMAPermuteN          = false;
 
     static constexpr ck_tile::index_t NumGroupsToMerge = 1;
 };
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp
index ad593b1418..8ea892a215 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp
@@ -14,7 +14,7 @@
 #include "grouped_convolution_backward_data_invoker.hpp"
 #include "run_grouped_convolution_bwd_data_example.inc"
 
-template <template <typename PrecType> typename GemmConfig>
+template <template <typename PrecType> typename ConvConfig>
 int run_grouped_conv_bwd_data_example(int argc, char* argv[])
 {
     using Invoker = GroupedConvolutionBackwardDataInvoker;
@@ -31,14 +31,14 @@ int run_grouped_conv_bwd_data_example(int argc, char* argv[])
     if(data_type == "fp16")
     {
         return run_grouped_conv_bwd_data_example_prec_type<Invoker,
-                                                           GemmConfig<ck_tile::half_t>,
+                                                           ConvConfig<ck_tile::half_t>,
                                                            ck_tile::half_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
     else if(data_type == "bf16")
     {
         return run_grouped_conv_bwd_data_example_prec_type<Invoker,
-                                                           GemmConfig<ck_tile::bf16_t>,
+                                                           ConvConfig<ck_tile::bf16_t>,
                                                            ck_tile::bf16_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp
index f6d20c3d3a..14a533ffc9 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp
@@ -8,7 +8,7 @@ struct GroupedConvolutionBackwardDataInvoker
 {
 
     template <ck_tile::index_t NDimSpatial,
-              typename GemmConfig,
+              typename ConvConfig,
               typename InDataType,
               typename WeiDataType,
               typename AccDataType,
@@ -26,12 +26,11 @@ struct GroupedConvolutionBackwardDataInvoker
 
         // Implicit GEMM Traits
         using GemmShape = ck_tile::TileGemmShape<
-            ck_tile::sequence<GemmConfig::M_Tile, GemmConfig::N_Tile, GemmConfig::K_Tile>,
-            ck_tile::sequence<GemmConfig::M_Warp, GemmConfig::N_Warp, GemmConfig::K_Warp>,
-            ck_tile::
-                sequence<GemmConfig::M_Warp_Tile, GemmConfig::N_Warp_Tile, GemmConfig::K_Warp_Tile>,
-            GemmConfig::PermuteA,
-            GemmConfig::PermuteB>;
+            ck_tile::sequence<ConvConfig::M_Tile, ConvConfig::N_Tile, ConvConfig::K_Tile>,
+            ck_tile::sequence<ConvConfig::M_Warp, ConvConfig::N_Warp, ConvConfig::K_Warp>,
+            ck_tile::sequence<ConvConfig::M_Warp_Tile,
+                              ConvConfig::N_Warp_Tile,
+                              ConvConfig::K_Warp_Tile>>;
 
         constexpr ck_tile::index_t VectorSizeA = 8;
         constexpr ck_tile::index_t VectorSizeB = 8;
@@ -40,8 +39,8 @@ struct GroupedConvolutionBackwardDataInvoker
         constexpr auto ConvSpec = ck_tile::ConvolutionSpecialization::Default;
         using TilePartitioner =
             ck_tile::GemmSpatiallyLocalTilePartitioner<GemmShape,
-                                                       GemmConfig::TileParitionerGroupNum,
-                                                       GemmConfig::TileParitionerM01>;
+                                                       ConvConfig::TileParitionerGroupNum,
+                                                       ConvConfig::TileParitionerM01>;
         using GroupedConvTraitsType = ck_tile::GroupedConvTraits<NDimSpatial,
                                                                  ConvSpec,
                                                                  InLayout,
@@ -53,17 +52,17 @@ struct GroupedConvolutionBackwardDataInvoker
                                                                  VectorSizeC>;
 
         using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits<
-            GemmConfig::kPadM,
-            GemmConfig::kPadN,
-            GemmConfig::kPadK,
-            GemmConfig::DoubleSmemBuffer,
+            ConvConfig::kPadM,
+            ConvConfig::kPadN,
+            ConvConfig::kPadK,
+            ConvConfig::DoubleSmemBuffer,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData::AsLayout,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData::BsLayout,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData::CLayout,
-            GemmConfig::TransposeC,
-            GemmConfig::UseStructuredSparsity,
+            ConvConfig::TransposeC,
+            false,
             false, // Persistent,
-            GemmConfig::NumWaveGroups>;
+            ConvConfig::NumWaveGroups>;
 
         using GemmPipelineProblem = ck_tile::GemmPipelineProblem<
             OutDataType,
@@ -79,7 +78,7 @@ struct GroupedConvolutionBackwardDataInvoker
             VectorSizeB>;
 
         using BaseGemmPipeline = typename PipelineTypeTraits<
-            GemmConfig::Pipeline>::template UniversalGemmPipeline<GemmPipelineProblem>;
+            ConvConfig::Pipeline>::template UniversalGemmPipeline<GemmPipelineProblem>;
 
         const ck_tile::index_t gemm_k =
             args.K_ * std::accumulate(args.filter_spatial_lengths_.begin(),
@@ -87,8 +86,8 @@ struct GroupedConvolutionBackwardDataInvoker
                                       1,
                                       std::multiplies<ck_tile::index_t>());
 
-        const ck_tile::index_t k_grain     = args.k_batch * GemmConfig::K_Tile;
-        const ck_tile::index_t K_split     = (gemm_k + k_grain - 1) / k_grain * GemmConfig::K_Tile;
+        const ck_tile::index_t k_grain     = args.k_batch * ConvConfig::K_Tile;
+        const ck_tile::index_t K_split     = (gemm_k + k_grain - 1) / k_grain * ConvConfig::K_Tile;
         const ck_tile::index_t num_loop    = TilePartitioner::GetLoopNum(K_split);
         const bool has_hot_loop            = BaseGemmPipeline::BlockHasHotloop(num_loop);
         const ck_tile::TailNumber tail_num = BaseGemmPipeline::GetBlockLoopTailNum(num_loop);
@@ -98,7 +97,7 @@ struct GroupedConvolutionBackwardDataInvoker
             [&](const auto has_hot_loop_, const auto tail_number_, const auto memory_operation_) {
                 constexpr bool has_hot_loop_v   = has_hot_loop_.value;
                 constexpr auto tail_number_v    = tail_number_.value;
-                constexpr auto scheduler        = GemmConfig::Scheduler;
+                constexpr auto scheduler        = ConvConfig::Scheduler;
                 constexpr auto memory_operation = memory_operation_.value;
 
                 using UniversalGemmProblem =
@@ -118,7 +117,7 @@ struct GroupedConvolutionBackwardDataInvoker
                                                           VectorSizeB>;
 
                 using GemmPipeline = typename PipelineTypeTraits<
-                    GemmConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
+                    ConvConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
 
                 using ConvEpilogue = ck_tile::CShuffleEpilogue<ck_tile::CShuffleEpilogueProblem<
                     OutDataType,
@@ -131,12 +130,12 @@ struct GroupedConvolutionBackwardDataInvoker
                     CDEElementWise,
                     TilePartitioner::MPerBlock,
                     TilePartitioner::NPerBlock,
-                    GemmConfig::M_Warp,
-                    GemmConfig::N_Warp,
-                    GemmConfig::M_Warp_Tile,
-                    GemmConfig::N_Warp_Tile,
-                    GemmConfig::K_Warp_Tile,
-                    GemmConfig::TransposeC,
+                    ConvConfig::M_Warp,
+                    ConvConfig::N_Warp,
+                    ConvConfig::M_Warp_Tile,
+                    ConvConfig::N_Warp_Tile,
+                    ConvConfig::K_Warp_Tile,
+                    ConvConfig::TransposeC,
                     memory_operation,
                     1,
                     true,
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp
index 0c00bb78e1..0e777c5f8a 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp
@@ -27,10 +27,9 @@ struct GroupedConvolutionBackwardWeightInvoker
         using GemmShape = ck_tile::TileGemmShape<
             ck_tile::sequence<ConvConfig::M_Tile, ConvConfig::N_Tile, ConvConfig::K_Tile>,
             ck_tile::sequence<ConvConfig::M_Warp, ConvConfig::N_Warp, ConvConfig::K_Warp>,
-            ck_tile::
-                sequence<ConvConfig::M_Warp_Tile, ConvConfig::N_Warp_Tile, ConvConfig::K_Warp_Tile>,
-            ConvConfig::PermuteA,
-            ConvConfig::PermuteB>;
+            ck_tile::sequence<ConvConfig::M_Warp_Tile,
+                              ConvConfig::N_Warp_Tile,
+                              ConvConfig::K_Warp_Tile>>;
 
         constexpr ck_tile::index_t VectorSizeA = ConvConfig::VectorSizeA;
         constexpr ck_tile::index_t VectorSizeB = ConvConfig::VectorSizeB;
@@ -61,7 +60,7 @@ struct GroupedConvolutionBackwardWeightInvoker
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdWeight::BsLayout,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdWeight::CLayout,
             ConvConfig::TransposeC,
-            ConvConfig::UseStructuredSparsity,
+            false,
             false, // Persistent,
             ConvConfig::NumWaveGroups>;
 
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp
index 047d143553..a8e41438c8 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp
@@ -29,10 +29,9 @@ struct GroupedConvolutionBackwardWeightTwoStageInvoker
         using GemmShape = ck_tile::TileGemmShape<
             ck_tile::sequence<ConvConfig::M_Tile, ConvConfig::N_Tile, ConvConfig::K_Tile>,
             ck_tile::sequence<ConvConfig::M_Warp, ConvConfig::N_Warp, ConvConfig::K_Warp>,
-            ck_tile::
-                sequence<ConvConfig::M_Warp_Tile, ConvConfig::N_Warp_Tile, ConvConfig::K_Warp_Tile>,
-            ConvConfig::PermuteA,
-            ConvConfig::PermuteB>;
+            ck_tile::sequence<ConvConfig::M_Warp_Tile,
+                              ConvConfig::N_Warp_Tile,
+                              ConvConfig::K_Warp_Tile>>;
 
         constexpr ck_tile::index_t VectorSizeA = 4;
         constexpr ck_tile::index_t VectorSizeB = 8;
@@ -62,7 +61,7 @@ struct GroupedConvolutionBackwardWeightTwoStageInvoker
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdWeight::BsLayout,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdWeight::CLayout,
             ConvConfig::TransposeC,
-            ConvConfig::UseStructuredSparsity,
+            false,
             false, // Persistent,
             ConvConfig::NumWaveGroups>;
 
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward.cpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward.cpp
index b979d4feb3..bef404b53a 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward.cpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward.cpp
@@ -14,7 +14,7 @@
 #include "grouped_convolution_forward_invoker.hpp"
 #include "run_grouped_convolution_fwd_example.inc"
 
-template <template <typename PrecType> typename GemmConfig>
+template <template <typename PrecType> typename ConvConfig>
 int run_grouped_conv_fwd_example(int argc, char* argv[])
 {
     using Invoker = GroupedConvolutionForwardInvoker;
@@ -31,14 +31,14 @@ int run_grouped_conv_fwd_example(int argc, char* argv[])
     if(data_type == "fp16")
     {
         return run_grouped_conv_fwd_example_prec_type<Invoker,
-                                                      GemmConfig<ck_tile::half_t>,
+                                                      ConvConfig<ck_tile::half_t>,
                                                       ck_tile::half_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
     else if(data_type == "bf16")
     {
         return run_grouped_conv_fwd_example_prec_type<Invoker,
-                                                      GemmConfig<ck_tile::bf16_t>,
+                                                      ConvConfig<ck_tile::bf16_t>,
                                                       ck_tile::bf16_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp
index bd31a99a7e..47395275b2 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp
@@ -14,7 +14,7 @@
 #include "grouped_convolution_forward_invoker.hpp"
 #include "run_grouped_convolution_fwd_bias_clamp_example.inc"
 
-template <template <typename PrecType> typename GemmConfig>
+template <template <typename PrecType> typename ConvConfig>
 int run_grouped_conv_fwd_bias_clamp_example(int argc, char* argv[])
 {
     using Invoker = GroupedConvolutionForwardInvoker;
@@ -31,14 +31,14 @@ int run_grouped_conv_fwd_bias_clamp_example(int argc, char* argv[])
     if(data_type == "fp16")
     {
         return run_grouped_conv_fwd_bias_clamp_example_prec_type<Invoker,
-                                                                 GemmConfig<ck_tile::half_t>,
+                                                                 ConvConfig<ck_tile::half_t>,
                                                                  ck_tile::half_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
     else if(data_type == "bf16")
     {
         return run_grouped_conv_fwd_bias_clamp_example_prec_type<Invoker,
-                                                                 GemmConfig<ck_tile::bf16_t>,
+                                                                 ConvConfig<ck_tile::bf16_t>,
                                                                  ck_tile::bf16_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_invoker.hpp
index d9a65b9639..2290f60d1f 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_invoker.hpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_invoker.hpp
@@ -14,7 +14,7 @@
 struct GroupedConvolutionForwardInvoker
 {
     template <ck_tile::index_t NDimSpatial,
-              typename GemmConfig,
+              typename ConvConfig,
               typename InDataType,
               typename WeiDataType,
               typename AccDataType,
@@ -36,12 +36,11 @@ struct GroupedConvolutionForwardInvoker
 
         // Implicit GEMM Traits
         using GemmShape = ck_tile::TileGemmShape<
-            ck_tile::sequence<GemmConfig::M_Tile, GemmConfig::N_Tile, GemmConfig::K_Tile>,
-            ck_tile::sequence<GemmConfig::M_Warp, GemmConfig::N_Warp, GemmConfig::K_Warp>,
-            ck_tile::
-                sequence<GemmConfig::M_Warp_Tile, GemmConfig::N_Warp_Tile, GemmConfig::K_Warp_Tile>,
-            GemmConfig::PermuteA,
-            GemmConfig::PermuteB>;
+            ck_tile::sequence<ConvConfig::M_Tile, ConvConfig::N_Tile, ConvConfig::K_Tile>,
+            ck_tile::sequence<ConvConfig::M_Warp, ConvConfig::N_Warp, ConvConfig::K_Warp>,
+            ck_tile::sequence<ConvConfig::M_Warp_Tile,
+                              ConvConfig::N_Warp_Tile,
+                              ConvConfig::K_Warp_Tile>>;
 
         constexpr ck_tile::index_t VectorSizeA      = 8;
         constexpr ck_tile::index_t VectorSizeB      = 8;
@@ -51,8 +50,8 @@ struct GroupedConvolutionForwardInvoker
         constexpr auto ConvSpec = ck_tile::ConvolutionSpecialization::Default;
         using TilePartitioner =
             ck_tile::GemmSpatiallyLocalTilePartitioner<GemmShape,
-                                                       GemmConfig::TileParitionerGroupNum,
-                                                       GemmConfig::TileParitionerM01>;
+                                                       ConvConfig::TileParitionerGroupNum,
+                                                       ConvConfig::TileParitionerM01>;
         using GroupedConvTraitsType = ck_tile::GroupedConvTraits<NDimSpatial,
                                                                  ConvSpec,
                                                                  InLayout,
@@ -62,22 +61,20 @@ struct GroupedConvolutionForwardInvoker
                                                                  VectorSizeA,
                                                                  VectorSizeB,
                                                                  VectorSizeC,
-                                                                 NumGroupsToMerge,
-                                                                 CDElementWise>;
+                                                                 NumGroupsToMerge>;
 
         using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits<
-            GemmConfig::kPadM,
-            GemmConfig::kPadN,
-            GemmConfig::kPadK,
-            GemmConfig::DoubleSmemBuffer,
+            ConvConfig::kPadM,
+            ConvConfig::kPadN,
+            ConvConfig::kPadK,
+            ConvConfig::DoubleSmemBuffer,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd::AsLayout,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd::BsLayout,
             typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd::CLayout,
-            GemmConfig::TransposeC,
-            GemmConfig::UseStructuredSparsity,
+            ConvConfig::TransposeC,
+            false,
             false, // Persistent,
-            GemmConfig::NumWaveGroups,
-            GemmConfig::Preshuffle>;
+            ConvConfig::NumWaveGroups>;
 
         using GemmPipelineProblem = ck_tile::GemmPipelineProblem<
             InDataType,
@@ -93,7 +90,7 @@ struct GroupedConvolutionForwardInvoker
             VectorSizeB>;
 
         using BaseGemmPipeline = typename PipelineTypeTraits<
-            GemmConfig::Pipeline>::template UniversalGemmPipeline<GemmPipelineProblem>;
+            ConvConfig::Pipeline>::template UniversalGemmPipeline<GemmPipelineProblem>;
 
         const ck_tile::index_t gemm_k =
             args.C_ * std::accumulate(args.filter_spatial_lengths_.begin(),
@@ -102,8 +99,8 @@ struct GroupedConvolutionForwardInvoker
                                       std::multiplies<ck_tile::index_t>());
 
         // Split-K parameters
-        const ck_tile::index_t k_grain     = args.k_batch * GemmConfig::K_Tile;
-        const ck_tile::index_t K_split     = (gemm_k + k_grain - 1) / k_grain * GemmConfig::K_Tile;
+        const ck_tile::index_t k_grain     = args.k_batch * ConvConfig::K_Tile;
+        const ck_tile::index_t K_split     = (gemm_k + k_grain - 1) / k_grain * ConvConfig::K_Tile;
         const ck_tile::index_t num_loop    = TilePartitioner::GetLoopNum(K_split);
         const bool has_hot_loop            = BaseGemmPipeline::BlockHasHotloop(num_loop);
         const ck_tile::TailNumber tail_num = BaseGemmPipeline::GetBlockLoopTailNum(num_loop);
@@ -116,7 +113,7 @@ struct GroupedConvolutionForwardInvoker
             [&](const auto has_hot_loop_, const auto tail_number_, const auto memory_operation_) {
                 constexpr bool has_hot_loop_v   = has_hot_loop_.value;
                 constexpr auto tail_number_v    = tail_number_.value;
-                constexpr auto scheduler        = GemmConfig::Scheduler;
+                constexpr auto scheduler        = ConvConfig::Scheduler;
                 constexpr auto memory_operation = memory_operation_.value;
 
                 using UniversalGemmProblem =
@@ -136,7 +133,7 @@ struct GroupedConvolutionForwardInvoker
                                                           VectorSizeB>;
 
                 using GemmPipeline = typename PipelineTypeTraits<
-                    GemmConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
+                    ConvConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
 
                 using ConvEpilogue = ck_tile::CShuffleEpilogue<ck_tile::CShuffleEpilogueProblem<
                     InDataType,
@@ -149,12 +146,12 @@ struct GroupedConvolutionForwardInvoker
                     CDElementWise,
                     TilePartitioner::MPerBlock,
                     TilePartitioner::NPerBlock,
-                    GemmConfig::M_Warp,
-                    GemmConfig::N_Warp,
-                    GemmConfig::M_Warp_Tile,
-                    GemmConfig::N_Warp_Tile,
-                    GemmConfig::K_Warp_Tile,
-                    GemmConfig::TransposeC,
+                    ConvConfig::M_Warp,
+                    ConvConfig::N_Warp,
+                    ConvConfig::M_Warp_Tile,
+                    ConvConfig::N_Warp_Tile,
+                    ConvConfig::K_Warp_Tile,
+                    ConvConfig::TransposeC,
                     memory_operation,
                     1,
                     true,
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor.cpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor.cpp
index 9534f19711..a76d0874ee 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor.cpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor.cpp
@@ -19,7 +19,7 @@
 #include "grouped_convolution_forward_large_tensor_invoker.hpp"
 #include "run_grouped_convolution_fwd_example.inc"
 
-template <template <typename PrecType> typename GemmConfig>
+template <template <typename PrecType> typename ConvConfig>
 int run_grouped_conv_fwd_example(int argc, char* argv[])
 {
     using Invoker = GroupedConvolutionForwardInvoker;
@@ -36,14 +36,14 @@ int run_grouped_conv_fwd_example(int argc, char* argv[])
     if(data_type == "fp16")
     {
         return run_grouped_conv_fwd_example_prec_type<Invoker,
-                                                      GemmConfig<ck_tile::half_t>,
+                                                      ConvConfig<ck_tile::half_t>,
                                                       ck_tile::half_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
     else if(data_type == "bf16")
     {
         return run_grouped_conv_fwd_example_prec_type<Invoker,
-                                                      GemmConfig<ck_tile::bf16_t>,
+                                                      ConvConfig<ck_tile::bf16_t>,
                                                       ck_tile::bf16_t>(
             in_layout, wei_layout, out_layout, argc, argv);
     }
diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor_invoker.hpp
index 2e98c0863b..4d983baac5 100644
--- a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor_invoker.hpp
+++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_large_tensor_invoker.hpp
@@ -7,7 +7,7 @@
 struct GroupedConvolutionForwardInvoker
 {
     template <ck_tile::index_t NDimSpatial,
-              typename GemmConfig,
+              typename ConvConfig,
               typename InDataType,
               typename WeiDataType,
               typename AccDataType,
@@ -29,12 +29,11 @@ struct GroupedConvolutionForwardInvoker
 
         // Implicit GEMM Traits
         using GemmShape = ck_tile::TileGemmShape<
-            ck_tile::sequence<GemmConfig::M_Tile, GemmConfig::N_Tile, GemmConfig::K_Tile>,
-            ck_tile::sequence<GemmConfig::M_Warp, GemmConfig::N_Warp, GemmConfig::K_Warp>,
-            ck_tile::
-                sequence<GemmConfig::M_Warp_Tile, GemmConfig::N_Warp_Tile, GemmConfig::K_Warp_Tile>,
-            GemmConfig::PermuteA,
-            GemmConfig::PermuteB>;
+            ck_tile::sequence<ConvConfig::M_Tile, ConvConfig::N_Tile, ConvConfig::K_Tile>,
+            ck_tile::sequence<ConvConfig::M_Warp, ConvConfig::N_Warp, ConvConfig::K_Warp>,
+            ck_tile::sequence<ConvConfig::M_Warp_Tile,
+                              ConvConfig::N_Warp_Tile,
+                              ConvConfig::K_Warp_Tile>>;
 
         constexpr ck_tile::index_t VectorSizeA = 8;
         constexpr ck_tile::index_t VectorSizeB = 8;
@@ -43,41 +42,53 @@ struct GroupedConvolutionForwardInvoker
         constexpr auto ConvSpec = ck_tile::ConvolutionSpecialization::Default;
         using TilePartitioner =
             ck_tile::GemmSpatiallyLocalTilePartitioner<GemmShape,
-                                                       GemmConfig::TileParitionerGroupNum,
-                                                       GemmConfig::TileParitionerM01>;
-        using GroupedConvTraitsType = ck_tile::GroupedConvTraits<NDimSpatial,
-                                                                 ConvSpec,
-                                                                 InLayout,
-                                                                 WeiLayout,
-                                                                 DsLayout,
-                                                                 OutLayout,
-                                                                 VectorSizeA,
-                                                                 VectorSizeB,
-                                                                 VectorSizeC,
-                                                                 1, /*NumGroupsToMerge*/
-                                                                 ck_tile::element_wise::PassThrough,
-                                                                 true /*EnableSplitImage*/>;
+                                                       ConvConfig::TileParitionerGroupNum,
+                                                       ConvConfig::TileParitionerM01>;
+
+        using GroupedConvTraitsTypeDefault = ck_tile::GroupedConvTraits<NDimSpatial,
+                                                                        ConvSpec,
+                                                                        InLayout,
+                                                                        WeiLayout,
+                                                                        DsLayout,
+                                                                        OutLayout,
+                                                                        VectorSizeA,
+                                                                        VectorSizeB,
+                                                                        VectorSizeC,
+                                                                        1, /*NumGroupsToMerge*/
+                                                                        false /*EnableSplitImage*/>;
+
+        using GroupedConvTraitsTypeLargeTensor =
+            ck_tile::GroupedConvTraits<NDimSpatial,
+                                       ConvSpec,
+                                       InLayout,
+                                       WeiLayout,
+                                       DsLayout,
+                                       OutLayout,
+                                       VectorSizeA,
+                                       VectorSizeB,
+                                       VectorSizeC,
+                                       1, /*NumGroupsToMerge*/
+                                       true /*EnableSplitImage*/>;
 
         using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits<
-            GemmConfig::kPadM,
-            GemmConfig::kPadN,
-            GemmConfig::kPadK,
-            GemmConfig::DoubleSmemBuffer,
-            typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd::AsLayout,
-            typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd::BsLayout,
-            typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd::CLayout,
-            GemmConfig::TransposeC,
-            GemmConfig::UseStructuredSparsity,
+            ConvConfig::kPadM,
+            ConvConfig::kPadN,
+            ConvConfig::kPadK,
+            ConvConfig::DoubleSmemBuffer,
+            typename GroupedConvTraitsTypeDefault::GroupedConvImplicitGemmTraitsFwd::AsLayout,
+            typename GroupedConvTraitsTypeDefault::GroupedConvImplicitGemmTraitsFwd::BsLayout,
+            typename GroupedConvTraitsTypeDefault::GroupedConvImplicitGemmTraitsFwd::CLayout,
+            ConvConfig::TransposeC,
+            false,
             false, // Persistent,
-            GemmConfig::NumWaveGroups,
-            GemmConfig::Preshuffle>;
+            ConvConfig::NumWaveGroups>;
 
         using GemmPipelineProblem = ck_tile::GemmPipelineProblem<
             InDataType,
             WeiDataType,
             AccDataType,
             GemmShape,
-            typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsFwd,
+            typename GroupedConvTraitsTypeDefault::GroupedConvImplicitGemmTraitsFwd,
             ck_tile::element_wise::PassThrough,
             ck_tile::element_wise::PassThrough,
             OutDataType,
@@ -86,7 +97,7 @@ struct GroupedConvolutionForwardInvoker
             VectorSizeB>;
 
         using BaseGemmPipeline = typename PipelineTypeTraits<
-            GemmConfig::Pipeline>::template UniversalGemmPipeline<GemmPipelineProblem>;
+            ConvConfig::Pipeline>::template UniversalGemmPipeline<GemmPipelineProblem>;
 
         const ck_tile::index_t gemm_k =
             args.C_ * std::accumulate(args.filter_spatial_lengths_.begin(),
@@ -95,8 +106,8 @@ struct GroupedConvolutionForwardInvoker
                                       std::multiplies<ck_tile::index_t>());
 
         // Split-K parameters
-        const ck_tile::index_t k_grain     = args.k_batch * GemmConfig::K_Tile;
-        const ck_tile::index_t K_split     = (gemm_k + k_grain - 1) / k_grain * GemmConfig::K_Tile;
+        const ck_tile::index_t k_grain     = args.k_batch * ConvConfig::K_Tile;
+        const ck_tile::index_t K_split     = (gemm_k + k_grain - 1) / k_grain * ConvConfig::K_Tile;
         const ck_tile::index_t num_loop    = TilePartitioner::GetLoopNum(K_split);
         const bool has_hot_loop            = BaseGemmPipeline::BlockHasHotloop(num_loop);
         const ck_tile::TailNumber tail_num = BaseGemmPipeline::GetBlockLoopTailNum(num_loop);
@@ -241,64 +252,71 @@ struct GroupedConvolutionForwardInvoker
         // =====================================================================
         // Kernel launch lambda: Uses EnableSplitImage based on layout support
         // =====================================================================
-        const auto Run =
-            [&](const auto has_hot_loop_, const auto tail_number_, const auto memory_operation_) {
-                constexpr bool has_hot_loop_v   = has_hot_loop_.value;
-                constexpr auto tail_number_v    = tail_number_.value;
-                constexpr auto scheduler        = GemmConfig::Scheduler;
-                constexpr auto memory_operation = memory_operation_.value;
+        const auto Run = [&]<bool EnableSplitImage>(const auto has_hot_loop_,
+                                                    const auto tail_number_,
+                                                    const auto memory_operation_) {
+            constexpr bool has_hot_loop_v   = has_hot_loop_.value;
+            constexpr auto tail_number_v    = tail_number_.value;
+            constexpr auto scheduler        = ConvConfig::Scheduler;
+            constexpr auto memory_operation = memory_operation_.value;
 
-                using UniversalGemmProblem =
-                    ck_tile::UniversalGemmPipelineProblem<InDataType,
-                                                          WeiDataType,
-                                                          AccDataType,
-                                                          GemmShape,
-                                                          GemmUniversalTraits,
-                                                          scheduler,
-                                                          has_hot_loop_v,
-                                                          tail_number_v,
-                                                          ck_tile::element_wise::PassThrough,
-                                                          ck_tile::element_wise::PassThrough,
-                                                          OutDataType,
-                                                          true,
-                                                          VectorSizeA,
-                                                          VectorSizeB>;
+            using GroupedConvTraitsType = std::conditional_t<EnableSplitImage,
+                                                             GroupedConvTraitsTypeLargeTensor,
+                                                             GroupedConvTraitsTypeDefault>;
 
-                using GemmPipeline = typename PipelineTypeTraits<
-                    GemmConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
+            using UniversalGemmProblem =
+                ck_tile::UniversalGemmPipelineProblem<InDataType,
+                                                      WeiDataType,
+                                                      AccDataType,
+                                                      GemmShape,
+                                                      GemmUniversalTraits,
+                                                      scheduler,
+                                                      has_hot_loop_v,
+                                                      tail_number_v,
+                                                      ck_tile::element_wise::PassThrough,
+                                                      ck_tile::element_wise::PassThrough,
+                                                      OutDataType,
+                                                      true,
+                                                      VectorSizeA,
+                                                      VectorSizeB>;
 
-                using ConvEpilogue = ck_tile::CShuffleEpilogue<ck_tile::CShuffleEpilogueProblem<
-                    InDataType,
-                    WeiDataType,
-                    DsDataType,
-                    AccDataType,
-                    OutDataType,
-                    typename GroupedConvTraitsType::ImplicitGemmDsLayout,
-                    ck_tile::tensor_layout::gemm::RowMajor,
-                    CDEElementWise,
-                    TilePartitioner::MPerBlock,
-                    TilePartitioner::NPerBlock,
-                    GemmConfig::M_Warp,
-                    GemmConfig::N_Warp,
-                    GemmConfig::M_Warp_Tile,
-                    GemmConfig::N_Warp_Tile,
-                    GemmConfig::K_Warp_Tile,
-                    GemmConfig::TransposeC,
-                    memory_operation,
-                    1,
-                    true,
-                    GroupedConvTraitsType::VectorSizeC>>;
+            using GemmPipeline = typename PipelineTypeTraits<
+                ConvConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
 
-                // Use split-image kernel if layout supports it, otherwise use regular kernel
-                using Kernel = ck_tile::GroupedConvolutionForwardKernel<GroupedConvTraitsType,
-                                                                        TilePartitioner,
-                                                                        GemmPipeline,
-                                                                        ConvEpilogue>;
+            using ConvEpilogue = ck_tile::CShuffleEpilogue<ck_tile::CShuffleEpilogueProblem<
+                InDataType,
+                WeiDataType,
+                DsDataType,
+                AccDataType,
+                OutDataType,
+                typename GroupedConvTraitsType::ImplicitGemmDsLayout,
+                ck_tile::tensor_layout::gemm::RowMajor,
+                CDEElementWise,
+                TilePartitioner::MPerBlock,
+                TilePartitioner::NPerBlock,
+                ConvConfig::M_Warp,
+                ConvConfig::N_Warp,
+                ConvConfig::M_Warp_Tile,
+                ConvConfig::N_Warp_Tile,
+                ConvConfig::K_Warp_Tile,
+                ConvConfig::TransposeC,
+                memory_operation,
+                1,
+                true,
+                GroupedConvTraitsType::VectorSizeC>>;
 
-                // Create kargs
-                auto kargs = Kernel::MakeKernelArgs(args);
+            // Use split-image kernel if layout supports it, otherwise use regular kernel
+            using Kernel = ck_tile::GroupedConvolutionForwardKernel<GroupedConvTraitsType,
+                                                                    TilePartitioner,
+                                                                    GemmPipeline,
+                                                                    ConvEpilogue>;
 
-                // Populate split-image metadata ONLY if using split-image kernel
+            // Create kargs
+            auto kargs = Kernel::MakeKernelArgs(args);
+
+            // Populate split-image metadata ONLY if using split-image kernel
+            if constexpr(EnableSplitImage)
+            {
                 kargs.num_spatial_pieces        = total_pieces;
                 kargs.split_image.total_d       = total_d;
                 kargs.split_image.total_h       = total_h;
@@ -319,35 +337,41 @@ struct GroupedConvolutionForwardInvoker
                                                    temp_pieces[i].h_size,
                                                    temp_pieces[i].w_size};
                 }
+            }
 
-                // Calculate grid: use total_blocks for split-image, or normal GridSize for regular
-                const dim3 grids  = dim3(total_blocks, kargs.GemmBatch, kargs.n_splits);
-                const dim3 blocks = Kernel::BlockSize();
+            // Calculate grid: use total_blocks for split-image, or normal GridSize for regular
+            const dim3 grids = [&]() {
+                if constexpr(EnableSplitImage)
+                    return dim3(total_blocks, kargs.GemmBatch, kargs.n_splits);
+                else
+                    return Kernel::GridSize(kargs);
+            }();
+            const dim3 blocks = Kernel::BlockSize();
 
-                if(!Kernel::IsSupportedArgument(kargs))
-                {
-                    throw std::runtime_error("Wrong! Arguments not supported! Skipping conv!\n");
-                }
+            if(!Kernel::IsSupportedArgument(kargs))
+            {
+                throw std::runtime_error("Wrong! Arguments not supported! Skipping conv!\n");
+            }
 
-                if(s.log_level_ > 0)
-                {
-                    std::cout << "Launching kernel with args: " << Kernel::GetName() << '\n'
-                              << "shape: " << GemmShape::GetName() << '\n'
-                              << "problem: " << UniversalGemmProblem::GetName() << '\n'
-                              << "pipeline: " << GemmPipeline::GetName() << '\n'
-                              << "grid: {" << grids.x << ", " << grids.y << ", " << grids.z << "}"
-                              << ", blocks: {" << blocks.x << ", " << blocks.y << ", " << blocks.z
-                              << "}" << '\n'
-                              << "Vector size A: " << GemmPipeline::GetVectorSizeA()
-                              << ", Vector size B: " << GemmPipeline::GetVectorSizeB()
-                              << ", Vector size C: " << ConvEpilogue::GetVectorSizeC() << std::endl;
-                }
+            if(s.log_level_ > 0)
+            {
+                std::cout << "Launching kernel with args: " << Kernel::GetName() << '\n'
+                          << "shape: " << GemmShape::GetName() << '\n'
+                          << "problem: " << UniversalGemmProblem::GetName() << '\n'
+                          << "pipeline: " << GemmPipeline::GetName() << '\n'
+                          << "grid: {" << grids.x << ", " << grids.y << ", " << grids.z << "}"
+                          << ", blocks: {" << blocks.x << ", " << blocks.y << ", " << blocks.z
+                          << "}" << '\n'
+                          << "Vector size A: " << GemmPipeline::GetVectorSizeA()
+                          << ", Vector size B: " << GemmPipeline::GetVectorSizeB()
+                          << ", Vector size C: " << ConvEpilogue::GetVectorSizeC() << std::endl;
+            }
 
-                ave_time = ck_tile::launch_kernel(
-                    s, ck_tile::make_kernel<kBlockPerCu>(Kernel{}, grids, blocks, 0, kargs));
+            ave_time = ck_tile::launch_kernel(
+                s, ck_tile::make_kernel<kBlockPerCu>(Kernel{}, grids, blocks, 0, kargs));
 
-                return ave_time;
-            };
+            return ave_time;
+        };
 
         // =====================================================================
         // Step 4: Dispatch kernel (split-image or regular based on decision)
diff --git a/example/ck_tile/20_grouped_convolution/run_grouped_convolution_bwd_data_example.inc b/example/ck_tile/20_grouped_convolution/run_grouped_convolution_bwd_data_example.inc
index c7be53f2df..e2ebfd6f8e 100644
--- a/example/ck_tile/20_grouped_convolution/run_grouped_convolution_bwd_data_example.inc
+++ b/example/ck_tile/20_grouped_convolution/run_grouped_convolution_bwd_data_example.inc
@@ -3,7 +3,7 @@
 #pragma once
 
 template <ck_tile::index_t NDimSpatial,
-          typename GemmConfig,
+          typename ConvConfig,
           typename Invoker,
           typename InDataType,
           typename WeiDataType,
@@ -17,7 +17,7 @@ float invoke_grouped_conv_bwd_data(ck_tile::GroupedConvBwdDataHostArgs& args,
                                    int n_repeat)
 {
     float ave_time = Invoker::template grouped_conv_bwd_data<NDimSpatial,
-                                                             GemmConfig,
+                                                             ConvConfig,
                                                              InDataType,
                                                              WeiDataType,
                                                              AccDataType,
@@ -39,7 +39,7 @@ float invoke_grouped_conv_bwd_data(ck_tile::GroupedConvBwdDataHostArgs& args,
 }
 
 template <ck_tile::index_t NDimSpatial,
-          typename GemmConfig,
+          typename ConvConfig,
           typename Invoker,
           typename InDataType,
           typename WeiDataType = InDataType,
@@ -141,7 +141,7 @@ int run_grouped_conv_bwd_data_example_with_layouts(
     std::cout << "output: " << output.mDesc << std::endl;
 
     invoke_grouped_conv_bwd_data<NDimSpatial,
-                                 GemmConfig,
+                                 ConvConfig,
                                  Invoker,
                                  InDataType,
                                  WeiDataType,
@@ -193,7 +193,7 @@ int run_grouped_conv_bwd_data_example_with_layouts(
 }
 
 template <typename Invoker,
-          typename GemmConfig,
+          typename ConvConfig,
           typename InPrecType,
           typename WeiPrecType = InPrecType,
           typename OutPrecType = InPrecType>
@@ -215,7 +215,7 @@ int run_grouped_conv_bwd_data_example_prec_type(
     if(in_layout == "NWGC" && wei_layout == "GKXC" && out_layout == "NWGK")
     {
         return run_grouped_conv_bwd_data_example_with_layouts<ck_tile::number<1>{},
-                                                              GemmConfig,
+                                                              ConvConfig,
                                                               Invoker,
                                                               InPrecType,
                                                               WeiPrecType,
@@ -225,7 +225,7 @@ int run_grouped_conv_bwd_data_example_prec_type(
     else if(in_layout == "NHWGC" && wei_layout == "GKYXC" && out_layout == "NHWGK")
     {
         return run_grouped_conv_bwd_data_example_with_layouts<ck_tile::number<2>{},
-                                                              GemmConfig,
+                                                              ConvConfig,
                                                               Invoker,
                                                               InPrecType,
                                                               WeiPrecType,
@@ -235,7 +235,7 @@ int run_grouped_conv_bwd_data_example_prec_type(
     else if(in_layout == "NDHWGC" && wei_layout == "GKZYXC" && out_layout == "NDHWGK")
     {
         return run_grouped_conv_bwd_data_example_with_layouts<ck_tile::number<3>{},
-                                                              GemmConfig,
+                                                              ConvConfig,
                                                               Invoker,
                                                               InPrecType,
                                                               WeiPrecType,
diff --git a/example/ck_tile/20_grouped_convolution/run_grouped_convolution_fwd_example.inc b/example/ck_tile/20_grouped_convolution/run_grouped_convolution_fwd_example.inc
index a6b1334d9e..7175a85ba7 100644
--- a/example/ck_tile/20_grouped_convolution/run_grouped_convolution_fwd_example.inc
+++ b/example/ck_tile/20_grouped_convolution/run_grouped_convolution_fwd_example.inc
@@ -3,7 +3,7 @@
 #pragma once
 
 template <ck_tile::index_t NDimSpatial,
-          typename GemmConfig,
+          typename ConvConfig,
           typename Invoker,
           typename InDataType,
           typename WeiDataType,
@@ -17,7 +17,7 @@ float invoke_grouped_conv_fwd(const ck_tile::GroupedConvFwdHostArgs<>& args,
                               int n_repeat)
 {
     float ave_time = Invoker::template grouped_conv_fwd<NDimSpatial,
-                                                        GemmConfig,
+                                                        ConvConfig,
                                                         InDataType,
                                                         WeiDataType,
                                                         AccDataType,
@@ -39,7 +39,7 @@ float invoke_grouped_conv_fwd(const ck_tile::GroupedConvFwdHostArgs<>& args,
 }
 
 template <ck_tile::index_t NDimSpatial,
-          typename GemmConfig,
+          typename ConvConfig,
           typename Invoker,
           typename InDataType,
           typename WeiDataType = InDataType,
@@ -141,7 +141,7 @@ int run_grouped_conv_fwd_example_with_layouts(
     std::cout << "output: " << output.mDesc << std::endl;
 
     invoke_grouped_conv_fwd<NDimSpatial,
-                            GemmConfig,
+                            ConvConfig,
                             Invoker,
                             InDataType,
                             WeiDataType,
@@ -193,7 +193,7 @@ int run_grouped_conv_fwd_example_with_layouts(
 }
 
 template <typename Invoker,
-          typename GemmConfig,
+          typename ConvConfig,
           typename InPrecType,
           typename WeiPrecType = InPrecType,
           typename OutPrecType = InPrecType>
@@ -215,7 +215,7 @@ int run_grouped_conv_fwd_example_prec_type(
     if(in_layout == "NWGC" && wei_layout == "GKXC" && out_layout == "NWGK")
     {
         return run_grouped_conv_fwd_example_with_layouts<ck_tile::number<1>{},
-                                                         GemmConfig,
+                                                         ConvConfig,
                                                          Invoker,
                                                          InPrecType,
                                                          WeiPrecType,
@@ -225,7 +225,7 @@ int run_grouped_conv_fwd_example_prec_type(
     else if(in_layout == "NHWGC" && wei_layout == "GKYXC" && out_layout == "NHWGK")
     {
         return run_grouped_conv_fwd_example_with_layouts<ck_tile::number<2>{},
-                                                         GemmConfig,
+                                                         ConvConfig,
                                                          Invoker,
                                                          InPrecType,
                                                          WeiPrecType,
@@ -235,7 +235,7 @@ int run_grouped_conv_fwd_example_prec_type(
     else if(in_layout == "NDHWGC" && wei_layout == "GKZYXC" && out_layout == "NDHWGK")
     {
         return run_grouped_conv_fwd_example_with_layouts<ck_tile::number<3>{},
-                                                         GemmConfig,
+                                                         ConvConfig,
                                                          Invoker,
                                                          InPrecType,
                                                          WeiPrecType,
diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp
index 8676530d87..7e70d2b422 100644
--- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp
+++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp
@@ -19,7 +19,7 @@
 namespace ck_tile {
 
 /// @brief The Grouped Convolution kernel device arguments.
-template <typename GroupedConvTraitsType_>
+template <typename GroupedConvTraitsType_, typename CDElementwise_>
 struct GroupedConvFwdKernelArgs
 {
 
@@ -31,7 +31,7 @@ struct GroupedConvFwdKernelArgs
                                GroupedConvTraitsType_::VectorSizeC,
                                GroupedConvTraitsType_::NumGroupsToMerge,
                                true>; // Split N enabled
-    using CDElementwise                 = typename GroupedConvTraitsType_::CDElementwise;
+    using CDElementwise                 = CDElementwise_;
     static constexpr index_t NumDTensor = GroupedConvTraitsType_::NumDTensor;
 
     template <
@@ -469,7 +469,8 @@ struct GroupedConvolutionForwardKernel
 
     using CDElementwise = typename EpiloguePipeline::CDElementwise;
 
-    using GroupedConvFwdKernelArgsSpecialized = GroupedConvFwdKernelArgs<GroupedConvTraitsType_>;
+    using GroupedConvFwdKernelArgsSpecialized =
+        GroupedConvFwdKernelArgs<GroupedConvTraitsType_, CDElementwise>;
 
     static constexpr bool IsSplitKSupported = false;
 
diff --git a/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp b/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp
index 703205fd6e..9b5a60ee1f 100644
--- a/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp
+++ b/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp
@@ -63,7 +63,6 @@ template <index_t NDimSpatial_,
           index_t VectorSizeB_      = 1,
           index_t VectorSizeC_      = 1,
           index_t NumGroupsToMerge_ = 1,
-          typename CDElementwise_   = PassThrough,
           bool EnableSplitImage_    = false>
 struct GroupedConvTraits
 {
@@ -83,7 +82,6 @@ struct GroupedConvTraits
     using WeiLayout                                               = WeiLayout_;
     using DsLayout                                                = DsLayout_;
     using OutLayout                                               = OutLayout_;
-    using CDElementwise                                           = CDElementwise_;
     using GroupedConvImplicitGemmTraitsFwd =
         TileGemmTraits<true,
                        true,