diff --git a/Jenkinsfile b/Jenkinsfile
index aa4045186e..80db99684a 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -1836,10 +1836,11 @@ pipeline {
                     }
                     agent{ label rocmnode("gfx90a") }
                     environment{
-                        setup_args = """ -DCMAKE_INSTALL_PREFIX=../install -DGPU_TARGETS="gfx90a" -DCMAKE_CXX_FLAGS=" -O3 " """
+                        setup_args = """ -DCMAKE_INSTALL_PREFIX=../install -DGPU_TARGETS="gfx90a" -DCK_CXX_STANDARD="17" -DCMAKE_CXX_FLAGS=" -O3 " """
                         execute_args = """ cd ../client_example && rm -rf build && mkdir build && cd build && \
                                            cmake -DCMAKE_PREFIX_PATH="${env.WORKSPACE}/install;/opt/rocm" \
                                            -DGPU_TARGETS="gfx90a" \
+                                           -DCK_CXX_STANDARD="17" \
                                            -DCMAKE_CXX_COMPILER="${build_compiler()}" \
                                            -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \
                                            -DCMAKE_CXX_FLAGS=" -O3 " .. && make -j """
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/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp b/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp
index aec9dfdde4..589caf88f4 100644
--- a/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp
+++ b/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp
@@ -12,12 +12,11 @@
 #include "ck_tile/ops/gemm_quant.hpp"
 
 #define CK_TILE_SUPPORTED_QUANT_GROUPS(X) \
-    X(1, 1, 64)    /* 1D */               \
-    X(1, 1, 128)   /* 1D */               \
-    X(1, 8, 128)   /* 2D N=8  */          \
-    X(1, 32, 128)  /* 2D N=32 */          \
-    X(1, 64, 128)  /* 2D N=64 */          \
-    X(1, 128, 128) /* 2D N=128 */
+    X(1, 1, 64)   /* 1D */                \
+    X(1, 1, 128)  /* 1D */                \
+    X(1, 8, 128)  /* 2D N=8  */           \
+    X(1, 32, 128) /* 2D N=32 */           \
+    X(1, 64, 128) /* 2D N=64 */
 
 template <typename PrecType, ck_tile::index_t M_Warp_Tile>
 constexpr ck_tile::index_t get_k_warp_tile()
diff --git a/example/ck_tile/40_streamk_gemm/README.md b/example/ck_tile/40_streamk_gemm/README.md
index fe9eb0c4f8..0272b1fe97 100644
--- a/example/ck_tile/40_streamk_gemm/README.md
+++ b/example/ck_tile/40_streamk_gemm/README.md
@@ -22,8 +22,8 @@ args:
           -a_layout    tensor A data layout (default: R)
           -b_layout    tensor B data layout (default: C)
           -c_layout    tensor C data layout (default: R)
-     -num_sk_blocks    number of Stream-K blocks. -1: chosen by algorithm, or user selected (default:-1)
 -reduction_strategy    strategy for storing results in C tensor. atomic/reduction (default:atomic)
+     -persistent_dp    persistent strategy for data-parallel section. Set to 0 for non-persistent or to 1 for persistent. (default:0)
           -stride_a    tensor A stride (default:0)
           -stride_b    tensor B stride (default:0)
           -stride_c    tensor C stride (default:0)
diff --git a/example/ck_tile/40_streamk_gemm/gemm_utils.hpp b/example/ck_tile/40_streamk_gemm/gemm_utils.hpp
index abcca7eaec..69095ca3d7 100644
--- a/example/ck_tile/40_streamk_gemm/gemm_utils.hpp
+++ b/example/ck_tile/40_streamk_gemm/gemm_utils.hpp
@@ -18,7 +18,6 @@ struct GemmConfigBase
 
     static constexpr bool TransposeC            = false;
     static constexpr bool UseStructuredSparsity = false;
-    static constexpr bool Persistent            = false;
 
     static constexpr int kBlockPerCu                = 1;
     static constexpr auto Scheduler                 = ck_tile::GemmPipelineScheduler::Intrawave;
@@ -27,12 +26,12 @@ struct GemmConfigBase
     static constexpr bool DoubleSmemBuffer          = false;
 };
 
-template <typename PrecType>
+template <typename PrecType, bool Persistent_>
 struct GemmConfigMemoryInterwave : public GemmConfigBase
 {
-    static constexpr ck_tile::index_t M_Tile = 128;
-    static constexpr ck_tile::index_t N_Tile = 128;
-    static constexpr ck_tile::index_t K_Tile = 32;
+    static constexpr ck_tile::index_t M_Tile = 256;
+    static constexpr ck_tile::index_t N_Tile = 256;
+    static constexpr ck_tile::index_t K_Tile = 16;
 
     static constexpr ck_tile::index_t M_Warp = 2;
     static constexpr ck_tile::index_t N_Warp = 2;
@@ -42,7 +41,8 @@ struct GemmConfigMemoryInterwave : public GemmConfigBase
     static constexpr ck_tile::index_t N_Warp_Tile = 32;
     static constexpr ck_tile::index_t K_Warp_Tile = sizeof(PrecType) == 2 ? 8 : 16;
 
-    static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave;
+    static constexpr bool Persistent = Persistent_;
+    static constexpr auto Scheduler  = ck_tile::GemmPipelineScheduler::Intrawave;
 };
 
 template <typename ADataType_, typename BDataType_ = ADataType_, typename CDataType_ = ADataType_>
@@ -96,12 +96,12 @@ auto create_args(int argc, char* argv[])
         .insert("a_layout", "R", "A tensor data layout - Row by default")
         .insert("b_layout", "C", "B tensor data layout - Column by default")
         .insert("c_layout", "R", "C tensor data layout - Row by default")
-        .insert("num_sk_blocks",
-                "-1",
-                "number of Stream-K blocks. -1: chosen by algorithm, or user selected")
         .insert("reduction_strategy",
                 "atomic",
                 "strategy for storing results in C tensor - atomic/reduction")
+        .insert("persistent_dp",
+                "0",
+                "0. Non-persistent data-parallel section, 1 Fully persistent kernel.")
         .insert("stride_a", "0", "Tensor A stride")
         .insert("stride_b", "0", "Tensor B stride")
         .insert("stride_c", "0", "Tensor C stride")
diff --git a/example/ck_tile/40_streamk_gemm/run_gemm_example.inc b/example/ck_tile/40_streamk_gemm/run_gemm_example.inc
index 6dd054ee11..17182d87dc 100644
--- a/example/ck_tile/40_streamk_gemm/run_gemm_example.inc
+++ b/example/ck_tile/40_streamk_gemm/run_gemm_example.inc
@@ -69,20 +69,18 @@ invoke_gemm(ck_tile::DeviceMem& a_m_k_dev_buf,
             int n_warmup,
             int n_repeat,
             bool flush_cache,
-            ck_tile::StreamKReductionStrategy reduction_strategy,
-            uint32_t num_sk_blocks)
+            ck_tile::StreamKReductionStrategy reduction_strategy)
 {
-    ck_tile::StreamKHostArgs args{a_m_k_dev_buf.GetDeviceBuffer(),
-                                  b_k_n_dev_buf.GetDeviceBuffer(),
-                                  c_m_n_dev_buf.GetDeviceBuffer(),
-                                  M,
-                                  N,
-                                  K,
-                                  stride_A,
-                                  stride_B,
-                                  stride_C,
-                                  reduction_strategy,
-                                  num_sk_blocks};
+    ck_tile::reboot::StreamKHostArgs args{a_m_k_dev_buf.GetDeviceBuffer(),
+                                          b_k_n_dev_buf.GetDeviceBuffer(),
+                                          c_m_n_dev_buf.GetDeviceBuffer(),
+                                          M,
+                                          N,
+                                          K,
+                                          stride_A,
+                                          stride_B,
+                                          stride_C,
+                                          reduction_strategy};
 
     std::tuple<float, ck_tile::index_t> ave_time_and_batch;
 
@@ -197,7 +195,6 @@ int run_gemm_example_with_layouts(int argc,
 
     ck_tile::StreamKReductionStrategy reduction_strategy =
         get_reduction_strategy_value(arg_parser.get_str("reduction_strategy"));
-    uint32_t num_sk_blocks = static_cast<uint32_t>(arg_parser.get_int("num_sk_blocks"));
 
     stride_A = ck_tile::get_default_stride(M, K, stride_A, is_row_major(a_layout));
     stride_B = ck_tile::get_default_stride(K, N, stride_B, is_row_major(b_layout));
@@ -261,8 +258,7 @@ int run_gemm_example_with_layouts(int argc,
                                                              n_warmup,
                                                              n_repeat,
                                                              flush_cache,
-                                                             reduction_strategy,
-                                                             num_sk_blocks);
+                                                             reduction_strategy);
 
     c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data());
 
@@ -279,10 +275,10 @@ int run_gemm_example_with_layouts(int argc,
               << " B_Type=" << DataTypeTraits<BDataType>::name
               << " C_Type=" << DataTypeTraits<CDataType>::name
               << " reduction_strategy=" << arg_parser.get_str("reduction_strategy") << " "
-              << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, "
-              << std::endl;
+              << " persistent_dp=" << arg_parser.get_str("persistent_dp") << " " << ave_time
+              << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " << std::endl;
 
-    bool pass = true;
+    bool pass = false;
 
     // Memory on host to store gpu reference result
     ck_tile::HostTensor<CDataType> c_m_n_ref(
diff --git a/example/ck_tile/40_streamk_gemm/streamk_gemm_basic.cpp b/example/ck_tile/40_streamk_gemm/streamk_gemm_basic.cpp
index 8ec409023d..e04cb00379 100644
--- a/example/ck_tile/40_streamk_gemm/streamk_gemm_basic.cpp
+++ b/example/ck_tile/40_streamk_gemm/streamk_gemm_basic.cpp
@@ -2,7 +2,6 @@
 // SPDX-License-Identifier:  MIT
 
 #include "gemm_utils.hpp"
-#include "run_gemm_example.inc"
 #include "ck_tile/ops/common.hpp"
 
 template <typename GemmConfig,
@@ -17,9 +16,8 @@ template <typename GemmConfig,
           typename ELayout,
           typename CDEElementWise,
           ck_tile::StreamKReductionStrategy ReductionStrategy>
-std::tuple<float, ck_tile::index_t> gemm(const ck_tile::StreamKHostArgs& args,
+std::tuple<float, ck_tile::index_t> gemm(const ck_tile::reboot::StreamKHostArgs& args,
                                          const ck_tile::stream_config& s)
-
 {
     using GemmShape = ck_tile::TileGemmShape<
         ck_tile::sequence<GemmConfig::M_Tile, GemmConfig::N_Tile, GemmConfig::K_Tile>,
@@ -29,7 +27,8 @@ std::tuple<float, ck_tile::index_t> gemm(const ck_tile::StreamKHostArgs& args,
         GemmConfig::PermuteA,
         GemmConfig::PermuteB>;
 
-    using TilePartitioner = ck_tile::StreamKTilePartitioner<GemmShape, ReductionStrategy>;
+    using TilePartitioner =
+        ck_tile::StreamKTilePartitioner_v2<GemmShape, ReductionStrategy, GemmConfig::Persistent>;
 
     using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits<GemmConfig::kPadM,
                                                                  GemmConfig::kPadN,
@@ -78,9 +77,13 @@ std::tuple<float, ck_tile::index_t> gemm(const ck_tile::StreamKHostArgs& args,
                                              memory_operation.value,
                                              GemmConfig::NumWaveGroups>>;
 
-        using Kernel = ck_tile::StreamKKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
+        using Kernel = ck_tile::reboot::StreamKKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
 
-        auto kargs = Kernel::MakeKernelArgs(args);
+        auto kargs                = Kernel::MakeKernelArgs(args);
+        const auto workspace_size = Kernel::GetWorkSpaceSize(kargs);
+        ck_tile::DeviceMem workspace_data(workspace_size);
+        workspace_data.SetZero();
+        kargs.workspace_ptr = workspace_data.GetDeviceBuffer();
 
         dim3 grids  = Kernel::GridSize(kargs.tile_partitioner);
         dim3 blocks = Kernel::BlockSize();
@@ -101,28 +104,28 @@ std::tuple<float, ck_tile::index_t> gemm(const ck_tile::StreamKHostArgs& args,
                       << std::endl;
         }
 
-        // Function to clear the output C tensor results after each repetition of the kernel
-        auto clear_gemm_output = [&]() {
+        auto reset_data_buffers = [&]() {
             if(ReductionStrategy == ck_tile::StreamKReductionStrategy::Atomic)
+            {
+                // Clear the output C tensor results after each repetition of the kernel
                 hipGetErrorString(hipMemsetAsync(
                     args.e_ptr, 0, args.M * args.N * sizeof(CDataType), s.stream_id_));
+            }
+            else if(ReductionStrategy == ck_tile::StreamKReductionStrategy::Reduction)
+            {
+                // Reset sk flags to zero before each repetition of the kernel
+                workspace_data.SetZero();
+            }
         };
 
-        std::function<void()> preprocess = clear_gemm_output;
+        std::function<void()> preprocess = reset_data_buffers;
 
         float ave_time = ck_tile::launch_kernel_time_mask(
             s,
             preprocess,
             ck_tile::make_kernel<GemmConfig::kBlockPerCu>(Kernel{}, grids, blocks, 0, kargs));
 
-        ck_tile::index_t num_wgs_per_tile = ck_tile::estimate_num_wgs_per_tile<ReductionStrategy>(
-            kargs.tile_partitioner.sk_num_blocks,
-            // k_iters_per_big_block could be 1, which indicates that all Stream-K workgroups are
-            // big and each does one iteration. Thus, we ensure the value passed in is at least 1 to
-            // avoid division by zero errors.
-            ck_tile::max(kargs.tile_partitioner.k_iters_per_big_block - 1, 1u),
-            kargs.tile_partitioner.k_iters_per_tile.get());
-
+        ck_tile::index_t num_wgs_per_tile = kargs.tile_partitioner.estimate_num_wgs_per_tile();
         return std::tuple{ave_time, num_wgs_per_tile};
     };
 
@@ -145,6 +148,8 @@ std::tuple<float, ck_tile::index_t> gemm(const ck_tile::StreamKHostArgs& args,
     }
 }
 
+#include "run_gemm_example.inc"
+
 template <typename GemmConfig, typename TypeConfig>
 int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int argc, char* argv[])
 {
@@ -164,7 +169,7 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a
     return 0;
 }
 
-template <template <typename PreType> typename GemmConfig>
+template <template <typename PreType, bool Persistent_> typename GemmConfig>
 int run_gemm_example(int argc, char* argv[])
 {
     auto [result, arg_parser] = create_args(argc, argv);
@@ -174,30 +179,63 @@ int run_gemm_example(int argc, char* argv[])
     std::string data_type = arg_parser.get_str("prec");
     std::string a_layout  = arg_parser.get_str("a_layout");
     std::string b_layout  = arg_parser.get_str("b_layout");
+    auto persistent_dp    = arg_parser.get_bool("persistent_dp");
 
     if(data_type == "bf16")
     {
         using TypeConfig = StreamKGemmTypeConfig<ck_tile::bf16_t>;
-        return run_gemm_example_prec_type<GemmConfig<ck_tile::bf16_t>, TypeConfig>(
-            a_layout, b_layout, argc, argv);
+        if(persistent_dp)
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::bf16_t, true>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
+        else
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::bf16_t, false>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
     }
     else if(data_type == "fp16")
     {
         using TypeConfig = StreamKGemmTypeConfig<ck_tile::half_t>;
-        return run_gemm_example_prec_type<GemmConfig<ck_tile::half_t>, TypeConfig>(
-            a_layout, b_layout, argc, argv);
+        if(persistent_dp)
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::half_t, true>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
+        else
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::half_t, false>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
     }
     else if(data_type == "fp8")
     {
         using TypeConfig = StreamKGemmTypeConfig<ck_tile::fp8_t, ck_tile::fp8_t, ck_tile::half_t>;
-        return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t>, TypeConfig>(
-            a_layout, b_layout, argc, argv);
+        if(persistent_dp)
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t, true>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
+        else
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t, false>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
     }
     else if(data_type == "bf8")
     {
         using TypeConfig = StreamKGemmTypeConfig<ck_tile::bf8_t, ck_tile::bf8_t, ck_tile::half_t>;
-        return run_gemm_example_prec_type<GemmConfig<ck_tile::bf8_t>, TypeConfig>(
-            a_layout, b_layout, argc, argv);
+        if(persistent_dp)
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::bf8_t, true>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
+        else
+        {
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::bf8_t, false>, TypeConfig>(
+                a_layout, b_layout, argc, argv);
+        }
     }
     else
     {
diff --git a/experimental/builder/include/ck_tile/builder/builder_utils.hpp b/experimental/builder/include/ck_tile/builder/builder_utils.hpp
index a69471c9ed..5b4981c630 100644
--- a/experimental/builder/include/ck_tile/builder/builder_utils.hpp
+++ b/experimental/builder/include/ck_tile/builder/builder_utils.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
index 586a119c75..5c397b1162 100644
--- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
+++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -95,7 +95,7 @@ concept AccessOrderDescriptor = requires(T t) {
     { t.order } -> std::convertible_to<std::array<size_t, 3>>;
 };
 
-// No requirements yet for a ConvAlgorithm concept.
+// Base requirement for all ConvAlgorithm concepts, i.e., all conv algorithm concepts must meet this concept.
 template <typename T>
 concept ConvAlgorithmDescriptor = std::is_class_v<T>;
 
@@ -183,4 +183,49 @@ concept SpecifiesLoopScheduler = requires {
     { T::loop_scheduler } -> std::convertible_to<LoopScheduler>;
 };
 
+/******************************************** */
+/* Concepts for the different device ops */
+/******************************************** */
+
+template <typename T>
+concept DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 = 
+    ConvAlgorithmDescriptor<T> &&
+    SpecifiesThreadBlock<T> && 
+    SpecifiesGridwiseXdlGemm<T> &&
+    SpecifiesBlockTransfer<T> && 
+    SpecifiesLdsTransfer<T> &&
+    SpecifiesThreadClusterAccessOrder<T> && 
+    SpecifiesSourceAccessOrder<T> &&
+    SpecifiesBlockGemm<T> && 
+    SpecifiesGemmSpecialization<T>;
+
+template <typename T>
+concept DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle =
+    ConvAlgorithmDescriptor<T> &&
+    SpecifiesThreadBlock<T> && 
+    SpecifiesGridwiseXdlGemm<T> &&
+    SpecifiesBlockTransfer<T> && 
+    SpecifiesLdsTransfer<T> &&
+    SpecifiesThreadClusterAccessOrder<T> && 
+    SpecifiesSourceAccessOrder<T> &&
+    SpecifiesFwdConcSpecialization<T> &&
+    SpecifiesGemmSpecialization<T> &&
+    SpecifiesNumPrefetchStages<T> &&
+    SpecifiesNumGroupsToMerge<T> &&
+    SpecifiesLoopScheduler<T>;
+
+template <typename T>
+concept DeviceGroupedConvFwdMultipleABD_Wmma_CShuffle = 
+    ConvAlgorithmDescriptor<T> &&
+    SpecifiesThreadBlock<T> && 
+    SpecifiesGridwiseWmmaGemm<T> &&
+    SpecifiesBlockTransfer<T> && 
+    SpecifiesLdsTransfer<T> &&
+    SpecifiesThreadClusterAccessOrder<T> && 
+    SpecifiesSourceAccessOrder<T> &&
+    SpecifiesFwdConcSpecialization<T> &&
+    SpecifiesGemmSpecialization<T> &&
+    SpecifiesNumPrefetchStages<T> &&
+    SpecifiesLoopScheduler<T>;
+
 } // namespace ck_tile::builder
diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp
index 68d5ec5a83..b24d0f47e9 100644
--- a/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp
+++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
diff --git a/experimental/builder/include/ck_tile/builder/conv_builder.hpp b/experimental/builder/include/ck_tile/builder/conv_builder.hpp
index d74948709b..bfb7386ea1 100644
--- a/experimental/builder/include/ck_tile/builder/conv_builder.hpp
+++ b/experimental/builder/include/ck_tile/builder/conv_builder.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
diff --git a/experimental/builder/include/ck_tile/builder/conv_factory.hpp b/experimental/builder/include/ck_tile/builder/conv_factory.hpp
index 8ea3e18d65..c48228fa37 100644
--- a/experimental/builder/include/ck_tile/builder/conv_factory.hpp
+++ b/experimental/builder/include/ck_tile/builder/conv_factory.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // A factory for instantiating CK convolution kernels.
 //
@@ -290,49 +290,32 @@ struct BlockGemmSpec
 };
 
 template <ConvAlgorithmDescriptor auto ALGORITHM>
-constexpr BlockGemmSpec SetBlockGemm()
+consteval BlockGemmSpec SetBlockGemm()
 {
     constexpr auto& BG = ALGORITHM.block_gemm;
 
     ck::BlockGemmPipelineScheduler scheduler;
     ck::BlockGemmPipelineVersion version;
 
-    if constexpr(BG.scheduler == BlockGemmPipelineScheduler::INTRAWAVE)
+    switch(BG.scheduler)
     {
+    case BlockGemmPipelineScheduler::INTRAWAVE:
         scheduler = ck::BlockGemmPipelineScheduler::Intrawave;
-    }
-    else if constexpr(BG.scheduler == BlockGemmPipelineScheduler::INTERWAVE)
-    {
+        break;
+    case BlockGemmPipelineScheduler::INTERWAVE:
         scheduler = ck::BlockGemmPipelineScheduler::Interwave;
-    }
-    else
-    {
-        static_assert(false, "Unknown BlockGemmPipelineScheduler");
+        break;
+    default: throw "Unknown BlockGemmPipelineScheduler";
     }
 
-    if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V1)
+    switch(BG.pipeline_version)
     {
-        version = ck::BlockGemmPipelineVersion::v1;
-    }
-    else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V2)
-    {
-        version = ck::BlockGemmPipelineVersion::v2;
-    }
-    else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V3)
-    {
-        version = ck::BlockGemmPipelineVersion::v3;
-    }
-    else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V4)
-    {
-        version = ck::BlockGemmPipelineVersion::v4;
-    }
-    else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V5)
-    {
-        version = ck::BlockGemmPipelineVersion::v5;
-    }
-    else
-    {
-        static_assert(false, "Unknown BlockGemmPipelineVersion");
+    case BlockGemmPipelineVersion::V1: version = ck::BlockGemmPipelineVersion::v1; break;
+    case BlockGemmPipelineVersion::V2: version = ck::BlockGemmPipelineVersion::v2; break;
+    case BlockGemmPipelineVersion::V3: version = ck::BlockGemmPipelineVersion::v3; break;
+    case BlockGemmPipelineVersion::V4: version = ck::BlockGemmPipelineVersion::v4; break;
+    case BlockGemmPipelineVersion::V5: version = ck::BlockGemmPipelineVersion::v5; break;
+    default: throw "Unknown BlockGemmPipelineVersion";
     }
 
     return BlockGemmSpec{.pipeline_version = version, .scheduler = scheduler};
@@ -441,18 +424,12 @@ template <ConvAlgorithmDescriptor auto ALGORITHM>
 consteval ck::LoopScheduler SetLoopScheduler()
 {
     constexpr auto loop_scheduler = ALGORITHM.loop_scheduler;
-
-    if constexpr(loop_scheduler == LoopScheduler::DEFAULT)
+    using ck_loop_sched           = ck::LoopScheduler;
+    switch(loop_scheduler)
     {
-        return ck::LoopScheduler::Default;
-    }
-    else if constexpr(loop_scheduler == LoopScheduler::INTERWAVE)
-    {
-        return ck::LoopScheduler::Interwave;
-    }
-    else
-    {
-        static_assert(false, "Unknown LoopScheduler");
+    case LoopScheduler::DEFAULT: return ck_loop_sched::Default;
+    case LoopScheduler::INTERWAVE: return ck_loop_sched::Interwave;
+    default: throw "Unknown LoopScheduler";
     }
 }
 
@@ -460,29 +437,16 @@ template <ConvAlgorithmDescriptor auto ALGORITHM>
 consteval ck::PipelineVersion SetGridwiseGemmPipelineVersion()
 {
     constexpr auto pipeline_version = ALGORITHM.gridwise_gemm.pipeline_version;
-    if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V1)
+    using ck_pipeline               = ck::PipelineVersion;
+    switch(pipeline_version)
     {
-        return ck::PipelineVersion::v1;
-    }
-    else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V2)
-    {
-        return ck::PipelineVersion::v2;
-    }
-    else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V3)
-    {
-        static_assert(false, "V3 is used only for stream-K.");
-    }
-    else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V4)
-    {
-        return ck::PipelineVersion::v4;
-    }
-    else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::WEIGHT_ONLY)
-    {
-        return ck::PipelineVersion::weight_only;
-    }
-    else
-    {
-        static_assert(false, "Unknown GridwiseGemmPipelineVersion");
+    case GridwiseGemmPipelineVersion::V1: return ck_pipeline::v1;
+    case GridwiseGemmPipelineVersion::V2: return ck_pipeline::v2;
+    case GridwiseGemmPipelineVersion::V4: return ck_pipeline::v4;
+    case GridwiseGemmPipelineVersion::WEIGHT_ONLY: return ck_pipeline::weight_only;
+    case GridwiseGemmPipelineVersion::V3:
+        throw "GridwiseGemmPipelineVersion::V3 is used only for stream-K.";
+    default: throw "Unknown GridwiseGemmPipelineVersion";
     }
 }
 
@@ -490,74 +454,27 @@ template <ConvAlgorithmDescriptor auto ALGORITHM>
 consteval ck::tensor_operation::device::GemmSpecialization SetGemmSpecialization()
 {
     constexpr auto gemm_spec = ALGORITHM.gemm_specialization;
+    using ck_gemm_spec       = ck::tensor_operation::device::GemmSpecialization;
 
-    if constexpr(gemm_spec == GemmSpecialization::Default)
+    switch(gemm_spec)
     {
-        return ck::tensor_operation::device::GemmSpecialization::Default;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::NPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::NPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::KPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::KPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MNPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MNPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MKPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MKPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::NKPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::NKPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MNKPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MNKPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::OPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::OPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MOPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::NOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::NOPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::KOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::KOPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MNOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MNOPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MKOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MKOPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::NKOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::NKOPadding;
-    }
-    else if constexpr(gemm_spec == GemmSpecialization::MNKOPadding)
-    {
-        return ck::tensor_operation::device::GemmSpecialization::MNKOPadding;
-    }
-    else
-    {
-        static_assert(false, "Unknown GemmSpecialization");
+    case GemmSpecialization::Default: return ck_gemm_spec::Default;
+    case GemmSpecialization::MPadding: return ck_gemm_spec::MPadding;
+    case GemmSpecialization::NPadding: return ck_gemm_spec::NPadding;
+    case GemmSpecialization::KPadding: return ck_gemm_spec::KPadding;
+    case GemmSpecialization::MNPadding: return ck_gemm_spec::MNPadding;
+    case GemmSpecialization::MKPadding: return ck_gemm_spec::MKPadding;
+    case GemmSpecialization::NKPadding: return ck_gemm_spec::NKPadding;
+    case GemmSpecialization::MNKPadding: return ck_gemm_spec::MNKPadding;
+    case GemmSpecialization::OPadding: return ck_gemm_spec::OPadding;
+    case GemmSpecialization::MOPadding: return ck_gemm_spec::MOPadding;
+    case GemmSpecialization::NOPadding: return ck_gemm_spec::NOPadding;
+    case GemmSpecialization::KOPadding: return ck_gemm_spec::KOPadding;
+    case GemmSpecialization::MNOPadding: return ck_gemm_spec::MNOPadding;
+    case GemmSpecialization::MKOPadding: return ck_gemm_spec::MKOPadding;
+    case GemmSpecialization::NKOPadding: return ck_gemm_spec::NKOPadding;
+    case GemmSpecialization::MNKOPadding: return ck_gemm_spec::MNKOPadding;
+    default: throw "Unknown GemmSpecialization";
     }
 }
 
@@ -565,30 +482,15 @@ template <ConvAlgorithmDescriptor auto ALGORITHM>
 consteval ck::BlockGemmPipelineVersion SetBlockGemmPipelineVersion()
 {
     constexpr auto version = ALGORITHM.pipeline_version;
-
-    if constexpr(version == BlockGemmPipelineVersion::V1)
+    using ck_block_gemm    = ck::BlockGemmPipelineVersion;
+    switch(version)
     {
-        return ck::BlockGemmPipelineVersion::v1;
-    }
-    else if constexpr(version == BlockGemmPipelineVersion::V2)
-    {
-        return ck::BlockGemmPipelineVersion::v2;
-    }
-    else if constexpr(version == BlockGemmPipelineVersion::V3)
-    {
-        return ck::BlockGemmPipelineVersion::v3;
-    }
-    else if constexpr(version == BlockGemmPipelineVersion::V4)
-    {
-        return ck::BlockGemmPipelineVersion::v4;
-    }
-    else if constexpr(version == BlockGemmPipelineVersion::V5)
-    {
-        return ck::BlockGemmPipelineVersion::v5;
-    }
-    else
-    {
-        static_assert(false, "Unknown BlockGemmPipelineVersion");
+    case BlockGemmPipelineVersion::V1: return ck_block_gemm::v1;
+    case BlockGemmPipelineVersion::V2: return ck_block_gemm::v2;
+    case BlockGemmPipelineVersion::V3: return ck_block_gemm::v3;
+    case BlockGemmPipelineVersion::V4: return ck_block_gemm::v4;
+    case BlockGemmPipelineVersion::V5: return ck_block_gemm::v5;
+    default: throw "Unknown BlockGemmPipelineVersion";
     }
 }
 
@@ -596,26 +498,14 @@ template <ConvAlgorithmDescriptor auto ALGORITHM>
 consteval ck::tensor_operation::device::ConvolutionForwardSpecialization SetFwdConvSpecialization()
 {
     constexpr auto specialization = ALGORITHM.fwd_specialization;
-
-    if constexpr(specialization == ConvFwdSpecialization::DEFAULT)
+    using ck_conv_spec            = ck::tensor_operation::device::ConvolutionForwardSpecialization;
+    switch(specialization)
     {
-        return ck::tensor_operation::device::ConvolutionForwardSpecialization::Default;
-    }
-    else if constexpr(specialization == ConvFwdSpecialization::FILTER_1X1_PAD0)
-    {
-        return ck::tensor_operation::device::ConvolutionForwardSpecialization::Filter1x1Pad0;
-    }
-    else if constexpr(specialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0)
-    {
-        return ck::tensor_operation::device::ConvolutionForwardSpecialization::Filter1x1Stride1Pad0;
-    }
-    else if constexpr(specialization == ConvFwdSpecialization::FILTER_3x3)
-    {
-        return ck::tensor_operation::device::ConvolutionForwardSpecialization::Filter3x3;
-    }
-    else
-    {
-        static_assert(false, "Unknown ConvFwdSpecialization");
+    case ConvFwdSpecialization::DEFAULT: return ck_conv_spec::Default;
+    case ConvFwdSpecialization::FILTER_1X1_PAD0: return ck_conv_spec::Filter1x1Pad0;
+    case ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0: return ck_conv_spec::Filter1x1Stride1Pad0;
+    case ConvFwdSpecialization::FILTER_3x3: return ck_conv_spec::Filter3x3;
+    default: throw "Unknown ConvFwdSpecialization";
     }
 }
 
@@ -627,7 +517,12 @@ namespace ck_tile::builder {
 template <ConvSignatureDescriptor auto SIGNATURE,
           ConvAlgorithmDescriptor auto ALGORITHM,
           auto VERSION>
-struct ConvFactory;
+struct ConvFactory
+{
+    // This will trigger if a specialization for the given convolution direction is not found.
+    // We should always catch this in an earlier validation check.
+    static_assert(false, "Unsupported device operation.");
+};
 
 // Factory specialization for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 instance
 // of a grouped forward convolution kernel.
@@ -635,7 +530,7 @@ template <ConvSignatureDescriptor auto SIGNATURE,
           ConvAlgorithmDescriptor auto ALGORITHM,
           StringLiteral VERSION>
     requires ConvDirectionIsForward<SIGNATURE> &&
-             ConvDeviceOpIs_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<SIGNATURE>
+             DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<decltype(ALGORITHM)>
 struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
 {
     static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim;
@@ -646,26 +541,6 @@ struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
     using Ops           = factory_internal::ElementwiseOps<SIGNATURE.elementwise_operation>;
     using AlgorithmType = decltype(ALGORITHM);
 
-    static_assert(SpecifiesThreadBlock<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify thread block info.");
-    static_assert(SpecifiesGridwiseXdlGemm<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify gridwise GEMM info.");
-    static_assert(SpecifiesBlockTransfer<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify block transfer info.");
-    static_assert(SpecifiesLdsTransfer<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify LDS transfer info.");
-    static_assert(
-        SpecifiesThreadClusterAccessOrder<AlgorithmType>,
-        "The convolution algorithm descriptor must specify thread cluster access order info.");
-    static_assert(SpecifiesSourceAccessOrder<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify source access order info.");
-    static_assert(SpecifiesBlockGemm<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify block gemm pipeline.");
-    static_assert(SpecifiesFwdConcSpecialization<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify forward convolution "
-                  "specialization.");
-    static_assert(SpecifiesGemmSpecialization<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify gemm specialization.");
     static_assert(ALGORITHM.block_transfer.lds_transfer_a.is_direct_load ==
                       ALGORITHM.block_transfer.lds_transfer_b.is_direct_load,
                   "A and B block transfers must both be direct load or not.");
@@ -757,7 +632,7 @@ template <ConvSignatureDescriptor auto SIGNATURE,
           ConvAlgorithmDescriptor auto ALGORITHM,
           StringLiteral VERSION>
     requires ConvDirectionIsForward<SIGNATURE> &&
-             ConvDeviceOpIs_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<SIGNATURE>
+             DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<decltype(ALGORITHM)>
 struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
 {
     static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim;
@@ -768,31 +643,6 @@ struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
     using Ops           = factory_internal::ElementwiseOps<SIGNATURE.elementwise_operation>;
     using AlgorithmType = decltype(ALGORITHM);
 
-    static_assert(SpecifiesThreadBlock<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify thread block info.");
-    static_assert(SpecifiesGridwiseXdlGemm<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify gridwise GEMM info.");
-    static_assert(SpecifiesBlockTransfer<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify block transfer info.");
-    static_assert(SpecifiesLdsTransfer<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify LDS transfer info.");
-    static_assert(
-        SpecifiesThreadClusterAccessOrder<AlgorithmType>,
-        "The convolution algorithm descriptor must specify thread cluster access order info.");
-    static_assert(SpecifiesSourceAccessOrder<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify source access order info.");
-    static_assert(SpecifiesFwdConcSpecialization<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify forward convolution "
-                  "specialization.");
-    static_assert(SpecifiesGemmSpecialization<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify gemm specialization.");
-    static_assert(SpecifiesNumPrefetchStages<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify number of prefetch stages.");
-    static_assert(SpecifiesLoopScheduler<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify loop scheduler.");
-    static_assert(SpecifiesNumGroupsToMerge<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify number of groups to merge.");
-
     static constexpr auto FWD_CONV_SPECIALIZATION =
         factory_internal::SetFwdConvSpecialization<ALGORITHM>();
     static constexpr auto GEMM_SPECIALIZATION =
@@ -879,7 +729,7 @@ template <ConvSignatureDescriptor auto SIGNATURE,
           ConvAlgorithmDescriptor auto ALGORITHM,
           StringLiteral VERSION>
     requires ConvDirectionIsForward<SIGNATURE> &&
-             ConvDeviceOpIs_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle<SIGNATURE>
+             DeviceGroupedConvFwdMultipleABD_Wmma_CShuffle<decltype(ALGORITHM)>
 struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
 {
     static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim;
@@ -890,27 +740,6 @@ struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
     using Ops           = factory_internal::ElementwiseOps<SIGNATURE.elementwise_operation>;
     using AlgorithmType = decltype(ALGORITHM);
 
-    static_assert(SpecifiesThreadBlock<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify thread block info.");
-    static_assert(SpecifiesGridwiseWmmaGemm<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify gridwise GEMM info.");
-    static_assert(SpecifiesBlockTransfer<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify block transfer info.");
-    static_assert(SpecifiesLdsTransfer<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify LDS transfer info.");
-    static_assert(
-        SpecifiesThreadClusterAccessOrder<AlgorithmType>,
-        "The convolution algorithm descriptor must specify thread cluster access order info.");
-    static_assert(SpecifiesSourceAccessOrder<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify source access order info.");
-    static_assert(SpecifiesFwdConcSpecialization<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify forward convolution "
-                  "specialization.");
-    static_assert(SpecifiesNumPrefetchStages<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify number of prefetch stages.");
-    static_assert(SpecifiesLoopScheduler<AlgorithmType>,
-                  "The convolution algorithm descriptor must specify loop scheduler.");
-
     static constexpr auto FWD_CONV_SPECIALIZATION =
         factory_internal::SetFwdConvSpecialization<ALGORITHM>();
     static constexpr auto GEMM_SPECIALIZATION =
diff --git a/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp
index 370e7b6521..983273b439 100644
--- a/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp
+++ b/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // This file defines the compile-time "signature" for grouped convolution operations.
 // A signature is a collection of properties that fully describe a convolution kernel's
@@ -21,7 +21,6 @@
 #include <type_traits>
 
 #include "ck_tile/builder/types.hpp"
-#include "ck_tile/builder/conv_signature_predicates.hpp"
 
 namespace ck_tile::builder {
 
@@ -41,9 +40,6 @@ template <DataType T>
 concept ConvDataType = (T == DataType::FP32) || (T == DataType::FP16) || (T == DataType::BF16) ||
                        (T == DataType::FP8) || (T == DataType::I8) || (T == DataType::U8);
 
-template <typename T>
-concept ConvDeviceOp = std::same_as<std::remove_cvref_t<T>, GroupConvDeviceOp>;
-
 template <typename T>
 concept ConvLayout = std::same_as<std::remove_cvref_t<T>, GroupConvLayout>;
 
@@ -55,7 +51,6 @@ concept ConvSignatureDescriptor = requires(T t) {
     { t.layout } -> ConvLayout;
     { t.data_type } -> std::convertible_to<DataType>;
     { t.elementwise_operation } -> std::convertible_to<ElementwiseOperation>;
-    { t.device_operation } -> ConvDeviceOp;
 };
 
 // Concept to validate a convolution signature's values.
@@ -63,7 +58,18 @@ template <auto Sig>
 concept ValidConvSignature = requires {
     requires ConvSpatialDim<Sig.spatial_dim>;
     requires ConvDataType<Sig.data_type>;
-    requires IsValidConvDeviceOp<Sig>;
 };
 
+// Predicate for forward convolution.
+template <auto Sig>
+concept ConvDirectionIsForward = (Sig.direction == ConvDirection::FORWARD);
+
+// Predicate for backward data convolution.
+template <auto Sig>
+concept ConvDirectionIsBackwardData = (Sig.direction == ConvDirection::BACKWARD_DATA);
+
+// Predicate for backward weight convolution.
+template <auto Sig>
+concept ConvDirectionIsBackwardWeight = (Sig.direction == ConvDirection::BACKWARD_WEIGHT);
+
 } // namespace ck_tile::builder
diff --git a/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp b/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp
deleted file mode 100644
index f947c7e329..0000000000
--- a/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp
+++ /dev/null
@@ -1,174 +0,0 @@
-// SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
-
-#pragma once
-
-#include <concepts>
-#include <type_traits>
-
-#include "ck_tile/builder/types.hpp"
-
-namespace ck_tile::builder {
-
-/**********************************************
- * Conv Direction Predicates
- **********************************************/
-
-// Predicate for forward convolution.
-template <auto Sig>
-concept ConvDirectionIsForward = (Sig.direction == ConvDirection::FORWARD);
-
-// Predicate for backward data convolution.
-template <auto Sig>
-concept ConvDirectionIsBackwardData = (Sig.direction == ConvDirection::BACKWARD_DATA);
-
-// Predicate for backward weight convolution.
-template <auto Sig>
-concept ConvDirectionIsBackwardWeight = (Sig.direction == ConvDirection::BACKWARD_WEIGHT);
-
-/**********************************************
- * Conv Fwd Device Op Predicates
- **********************************************/
-
-// Predicate for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 =
-    (Sig.device_operation._fwd ==
-     FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3);
-
-// Predicate for DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK =
-    (Sig.device_operation._fwd ==
-     FwdGroupConvDeviceOperation::DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK);
-
-// Predicate for DeviceGroupedConvFwdMultipleD_Wmma_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle =
-    (Sig.device_operation._fwd ==
-     FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleD_Wmma_CShuffle);
-
-// Predicate for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle =
-    (Sig.device_operation._fwd ==
-     FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle);
-
-// Predicate for DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor =
-    (Sig.device_operation._fwd ==
-     FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor);
-
-// Generic predicate to check if signature uses any forward convolution device operation.
-template <auto Sig>
-concept ConvDeviceOpIsForward =
-    ConvDeviceOpIs_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor<Sig>;
-
-/**********************************************
- * Conv Bwd Weight Device Op Predicates
- **********************************************/
-
-// Predicate for DeviceGroupedConvBwdWeight operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeight =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeight);
-
-// Predicate for DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle);
-
-// Predicate for DeviceGroupedConvBwdWeight_Xdl_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeight_Xdl_CShuffle);
-
-// Predicate for DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle);
-
-// Predicate for DeviceGroupedConvBwdWeight_Wmma_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Wmma_CShuffle =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeight_Wmma_CShuffle);
-
-// Predicate for DeviceGroupedConvBwdWeight_Xdl_CShuffleV3 operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Xdl_CShuffleV3 =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeight_Xdl_CShuffleV3);
-
-// Predicate for DeviceGroupedConvBwdWeightMultipleD operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeightMultipleD =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeightMultipleD);
-
-// Predicate for DeviceGroupedConvBwdWeight_Dl operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Dl =
-    (Sig.device_operation._bwd_weight ==
-     BwdWeightGroupConvDeviceOperation::DeviceGroupedConvBwdWeight_Dl);
-
-// Generic predicate to check if signature uses any backward weight convolution device operation.
-template <auto Sig>
-concept ConvDeviceOpIsBackwardWeight =
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeight<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Xdl_CShuffle<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Wmma_CShuffle<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Xdl_CShuffleV3<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeightMultipleD<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdWeight_Dl<Sig>;
-
-/**********************************************
- * Conv Bwd Data Device Op Predicates
- **********************************************/
-
-// Predicate for DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 =
-    (Sig.device_operation._bwd_data ==
-     BwdDataGroupConvDeviceOperation::DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1);
-
-// Predicate for DeviceGroupedConvBwdDataMultipleD operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdDataMultipleD =
-    (Sig.device_operation._bwd_data ==
-     BwdDataGroupConvDeviceOperation::DeviceGroupedConvBwdDataMultipleD);
-
-// Predicate for DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle operation.
-template <auto Sig>
-concept ConvDeviceOpIs_DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle =
-    (Sig.device_operation._bwd_data ==
-     BwdDataGroupConvDeviceOperation::DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle);
-
-// Generic predicate to check if signature uses any backward data convolution device operation.
-template <auto Sig>
-concept ConvDeviceOpIsBackwardData =
-    ConvDeviceOpIs_DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdDataMultipleD<Sig> ||
-    ConvDeviceOpIs_DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle<Sig>;
-
-/**********************************************
- * Generic Device Op Predicates
- **********************************************/
-
-// Generic predicate to check if signature uses any device operation.
-template <auto Sig>
-concept IsValidConvDeviceOp = ConvDeviceOpIsForward<Sig> || ConvDeviceOpIsBackwardData<Sig> ||
-                              ConvDeviceOpIsBackwardWeight<Sig>;
-
-} // namespace ck_tile::builder
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp
index a47ad0ef57..c9b45691cc 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // Compile-time reflection for CK device kernel instances.
 //
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
new file mode 100644
index 0000000000..a0def3e5d9
--- /dev/null
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
@@ -0,0 +1,285 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include "instance_traits.hpp"
+#include "ck/tensor_operation/gpu/device/convolution_backward_weight_specialization.hpp"
+
+// Forward declaration to avoid circular dependency
+namespace ck::tensor_operation::device {
+
+template <ck::index_t NDimSpatial,
+          typename InLayout,
+          typename WeiLayout,
+          typename OutLayout,
+          typename InDataType,
+          typename WeiDataType,
+          typename OutDataType,
+          typename AccDataType,
+          typename InElementwiseOperation,
+          typename WeiElementwiseOperation,
+          typename OutElementwiseOperation,
+          ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization
+              ConvBackwardWeightSpecialization,
+          ck::index_t BlockSize,
+          ck::index_t MPerBlock,
+          ck::index_t NPerBlock,
+          ck::index_t K0PerBlock,
+          ck::index_t K1,
+          ck::index_t MPerXDL,
+          ck::index_t NPerXDL,
+          ck::index_t MXdlPerWave,
+          ck::index_t NXdlPerWave,
+          typename ABlockTransferThreadClusterLengths_K0_M_K1,
+          typename ABlockTransferThreadClusterArrangeOrder,
+          typename ABlockTransferSrcAccessOrder,
+          ck::index_t ABlockTransferSrcVectorDim,
+          ck::index_t ABlockTransferSrcScalarPerVector,
+          ck::index_t ABlockTransferDstScalarPerVector_K1,
+          bool ABlockLdsAddExtraM,
+          typename BBlockTransferThreadClusterLengths_K0_N_K1,
+          typename BBlockTransferThreadClusterArrangeOrder,
+          typename BBlockTransferSrcAccessOrder,
+          ck::index_t BBlockTransferSrcVectorDim,
+          ck::index_t BBlockTransferSrcScalarPerVector,
+          ck::index_t BBlockTransferDstScalarPerVector_K1,
+          bool BBlockLdsAddExtraN,
+          ck::index_t CShuffleMXdlPerWavePerShuffle,
+          ck::index_t CShuffleNXdlPerWavePerShuffle,
+          typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
+          ck::index_t CBlockTransferScalarPerVector_NWaveNPerXdl,
+          typename ComputeTypeA,
+          typename ComputeTypeB,
+          ck::index_t MaxTransposeTransferSrcScalarPerVector,
+          ck::index_t MaxTransposeTransferDstScalarPerVector>
+struct DeviceGroupedConvBwdWeight_Xdl_CShuffle;
+
+} // namespace ck::tensor_operation::device
+
+namespace ck_tile {
+namespace reflect {
+
+template <ck::index_t NDimSpatial,
+          typename InLayout_,
+          typename WeiLayout_,
+          typename OutLayout_,
+          typename InDataType_,
+          typename WeiDataType_,
+          typename OutDataType_,
+          typename AccDataType_,
+          typename InElementwiseOperation_,
+          typename WeiElementwiseOperation_,
+          typename OutElementwiseOperation_,
+          ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization
+              ConvBackwardWeightSpecialization,
+          ck::index_t BlockSize,
+          ck::index_t MPerBlock,
+          ck::index_t NPerBlock,
+          ck::index_t K0PerBlock,
+          ck::index_t K1,
+          ck::index_t MPerXDL,
+          ck::index_t NPerXDL,
+          ck::index_t MXdlPerWave,
+          ck::index_t NXdlPerWave,
+          typename ABlockTransferThreadClusterLengths_K0_M_K1_,
+          typename ABlockTransferThreadClusterArrangeOrder_,
+          typename ABlockTransferSrcAccessOrder_,
+          ck::index_t ABlockTransferSrcVectorDim,
+          ck::index_t ABlockTransferSrcScalarPerVector,
+          ck::index_t ABlockTransferDstScalarPerVector_K1,
+          bool ABlockLdsAddExtraM,
+          typename BBlockTransferThreadClusterLengths_K0_N_K1_,
+          typename BBlockTransferThreadClusterArrangeOrder_,
+          typename BBlockTransferSrcAccessOrder_,
+          ck::index_t BBlockTransferSrcVectorDim,
+          ck::index_t BBlockTransferSrcScalarPerVector,
+          ck::index_t BBlockTransferDstScalarPerVector_K1,
+          bool BBlockLdsAddExtraN,
+          ck::index_t CShuffleMXdlPerWavePerShuffle,
+          ck::index_t CShuffleNXdlPerWavePerShuffle,
+          typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_,
+          ck::index_t CBlockTransferScalarPerVector_NWaveNPerXdl,
+          typename ComputeTypeA_,
+          typename ComputeTypeB_,
+          ck::index_t MaxTransposeTransferSrcScalarPerVector,
+          ck::index_t MaxTransposeTransferDstScalarPerVector>
+struct InstanceTraits<ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle<
+    NDimSpatial,
+    InLayout_,
+    WeiLayout_,
+    OutLayout_,
+    InDataType_,
+    WeiDataType_,
+    OutDataType_,
+    AccDataType_,
+    InElementwiseOperation_,
+    WeiElementwiseOperation_,
+    OutElementwiseOperation_,
+    ConvBackwardWeightSpecialization,
+    BlockSize,
+    MPerBlock,
+    NPerBlock,
+    K0PerBlock,
+    K1,
+    MPerXDL,
+    NPerXDL,
+    MXdlPerWave,
+    NXdlPerWave,
+    ABlockTransferThreadClusterLengths_K0_M_K1_,
+    ABlockTransferThreadClusterArrangeOrder_,
+    ABlockTransferSrcAccessOrder_,
+    ABlockTransferSrcVectorDim,
+    ABlockTransferSrcScalarPerVector,
+    ABlockTransferDstScalarPerVector_K1,
+    ABlockLdsAddExtraM,
+    BBlockTransferThreadClusterLengths_K0_N_K1_,
+    BBlockTransferThreadClusterArrangeOrder_,
+    BBlockTransferSrcAccessOrder_,
+    BBlockTransferSrcVectorDim,
+    BBlockTransferSrcScalarPerVector,
+    BBlockTransferDstScalarPerVector_K1,
+    BBlockLdsAddExtraN,
+    CShuffleMXdlPerWavePerShuffle,
+    CShuffleNXdlPerWavePerShuffle,
+    CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_,
+    CBlockTransferScalarPerVector_NWaveNPerXdl,
+    ComputeTypeA_,
+    ComputeTypeB_,
+    MaxTransposeTransferSrcScalarPerVector,
+    MaxTransposeTransferDstScalarPerVector>>
+{
+    static constexpr auto kTensorOpName = "DeviceGroupedConvBwdWeight_Xdl_CShuffle";
+
+    static constexpr ck::index_t kNDimSpatial = NDimSpatial;
+
+    using InLayout  = InLayout_;
+    using WeiLayout = WeiLayout_;
+    using OutLayout = OutLayout_;
+
+    using InDataType  = InDataType_;
+    using WeiDataType = WeiDataType_;
+    using OutDataType = OutDataType_;
+    using AccDataType = AccDataType_;
+
+    using InElementwiseOperation  = InElementwiseOperation_;
+    using WeiElementwiseOperation = WeiElementwiseOperation_;
+    using OutElementwiseOperation = OutElementwiseOperation_;
+
+    static constexpr auto kConvBackwardWeightSpecialization = ConvBackwardWeightSpecialization;
+
+    static constexpr ck::index_t kBlockSize   = BlockSize;
+    static constexpr ck::index_t kMPerBlock   = MPerBlock;
+    static constexpr ck::index_t kNPerBlock   = NPerBlock;
+    static constexpr ck::index_t kK0PerBlock  = K0PerBlock;
+    static constexpr ck::index_t kK1          = K1;
+    static constexpr ck::index_t kMPerXDL     = MPerXDL;
+    static constexpr ck::index_t kNPerXDL     = NPerXDL;
+    static constexpr ck::index_t kMXdlPerWave = MXdlPerWave;
+    static constexpr ck::index_t kNXdlPerWave = NXdlPerWave;
+
+    using ABlockTransferThreadClusterLengths_K0_M_K1 = ABlockTransferThreadClusterLengths_K0_M_K1_;
+    using ABlockTransferThreadClusterArrangeOrder    = ABlockTransferThreadClusterArrangeOrder_;
+    using ABlockTransferSrcAccessOrder               = ABlockTransferSrcAccessOrder_;
+    static constexpr ck::index_t kABlockTransferSrcVectorDim = ABlockTransferSrcVectorDim;
+    static constexpr ck::index_t kABlockTransferSrcScalarPerVector =
+        ABlockTransferSrcScalarPerVector;
+    static constexpr ck::index_t kABlockTransferDstScalarPerVector_K1 =
+        ABlockTransferDstScalarPerVector_K1;
+    static constexpr bool kABlockLdsAddExtraM = ABlockLdsAddExtraM;
+
+    using BBlockTransferThreadClusterLengths_K0_N_K1 = BBlockTransferThreadClusterLengths_K0_N_K1_;
+    using BBlockTransferThreadClusterArrangeOrder    = BBlockTransferThreadClusterArrangeOrder_;
+    using BBlockTransferSrcAccessOrder               = BBlockTransferSrcAccessOrder_;
+    static constexpr ck::index_t kBBlockTransferSrcVectorDim = BBlockTransferSrcVectorDim;
+    static constexpr ck::index_t kBBlockTransferSrcScalarPerVector =
+        BBlockTransferSrcScalarPerVector;
+    static constexpr ck::index_t kBBlockTransferDstScalarPerVector_K1 =
+        BBlockTransferDstScalarPerVector_K1;
+    static constexpr bool kBBlockLdsAddExtraN = BBlockLdsAddExtraN;
+
+    static constexpr ck::index_t kCShuffleMXdlPerWavePerShuffle = CShuffleMXdlPerWavePerShuffle;
+    static constexpr ck::index_t kCShuffleNXdlPerWavePerShuffle = CShuffleNXdlPerWavePerShuffle;
+
+    using CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock =
+        CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_;
+    static constexpr ck::index_t kCBlockTransferScalarPerVector_NWaveNPerXdl =
+        CBlockTransferScalarPerVector_NWaveNPerXdl;
+
+    using ComputeTypeA = ComputeTypeA_;
+    using ComputeTypeB = ComputeTypeB_;
+
+    static constexpr ck::index_t kMaxTransposeTransferSrcScalarPerVector =
+        MaxTransposeTransferSrcScalarPerVector;
+    static constexpr ck::index_t kMaxTransposeTransferDstScalarPerVector =
+        MaxTransposeTransferDstScalarPerVector;
+
+    // Static member function to generate instance string
+    static std::string instance_string()
+    {
+        std::ostringstream oss;
+
+        // Kernel type name
+        oss << "DeviceGroupedConvBwdWeight_Xdl_CShuffle";
+
+        // Template parameters in exact order
+        oss << "<" << kNDimSpatial;                     // 1. NDimSpatial
+        oss << "," << detail::layout_name<InLayout>();  // 2. InLayout
+        oss << "," << detail::layout_name<WeiLayout>(); // 3. WeiLayout
+        oss << "," << detail::layout_name<OutLayout>(); // 4. OutLayout
+        oss << "," << detail::type_name<InDataType>();  // 5. InDataType
+        oss << "," << detail::type_name<WeiDataType>(); // 6. WeiDataType
+        oss << "," << detail::type_name<OutDataType>(); // 7. OutDataType
+        oss << "," << detail::type_name<AccDataType>(); // 8. AccDataType
+        oss << ","
+            << detail::elementwise_op_name<InElementwiseOperation>(); // 9. InElementwiseOperation
+        oss << ","
+            << detail::elementwise_op_name<WeiElementwiseOperation>(); // 10.
+                                                                       // WeiElementwiseOperation
+        oss << ","
+            << detail::elementwise_op_name<OutElementwiseOperation>(); // 11.
+                                                                       // OutElementwiseOperation
+        oss << ","
+            << detail::conv_bwd_weight_spec_name(
+                   kConvBackwardWeightSpecialization); // 12. ConvBackwardWeightSpecialization
+        oss << "," << kBlockSize;                      // 13. BlockSize
+        oss << "," << kMPerBlock;                      // 14. MPerBlock
+        oss << "," << kNPerBlock;                      // 15. NPerBlock
+        oss << "," << kK0PerBlock;                     // 16. K0PerBlock
+        oss << "," << kK1;                             // 17. K1
+        oss << "," << kMPerXDL;                        // 18. MPerXDL
+        oss << "," << kNPerXDL;                        // 19. NPerXDL
+        oss << "," << kMXdlPerWave;                    // 20. MXdlPerWave
+        oss << "," << kNXdlPerWave;                    // 21. NXdlPerWave
+        oss << "," << detail::sequence_name<ABlockTransferThreadClusterLengths_K0_M_K1>(); // 22.
+        oss << "," << detail::sequence_name<ABlockTransferThreadClusterArrangeOrder>();    // 23.
+        oss << "," << detail::sequence_name<ABlockTransferSrcAccessOrder>();               // 24.
+        oss << "," << kABlockTransferSrcVectorDim;                                         // 25.
+        oss << "," << kABlockTransferSrcScalarPerVector;                                   // 26.
+        oss << "," << kABlockTransferDstScalarPerVector_K1;                                // 27.
+        oss << "," << (kABlockLdsAddExtraM ? "true" : "false");                            // 28.
+        oss << "," << detail::sequence_name<BBlockTransferThreadClusterLengths_K0_N_K1>(); // 29.
+        oss << "," << detail::sequence_name<BBlockTransferThreadClusterArrangeOrder>();    // 30.
+        oss << "," << detail::sequence_name<BBlockTransferSrcAccessOrder>();               // 31.
+        oss << "," << kBBlockTransferSrcVectorDim;                                         // 32.
+        oss << "," << kBBlockTransferSrcScalarPerVector;                                   // 33.
+        oss << "," << kBBlockTransferDstScalarPerVector_K1;                                // 34.
+        oss << "," << (kBBlockLdsAddExtraN ? "true" : "false");                            // 35.
+        oss << "," << kCShuffleMXdlPerWavePerShuffle;                                      // 36.
+        oss << "," << kCShuffleNXdlPerWavePerShuffle;                                      // 37.
+        oss << ","
+            << detail::sequence_name<
+                   CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock>(); // 38.
+        oss << "," << kCBlockTransferScalarPerVector_NWaveNPerXdl;                    // 39.
+        oss << "," << detail::type_name<ComputeTypeA>();                              // 40.
+        oss << "," << detail::type_name<ComputeTypeB>();                              // 41.
+        oss << "," << kMaxTransposeTransferSrcScalarPerVector;                        // 42.
+        oss << "," << kMaxTransposeTransferDstScalarPerVector;                        // 43.
+        oss << ">";
+
+        return oss.str();
+    }
+};
+
+} // namespace reflect
+} // namespace ck_tile
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp
index 54ba224f7f..8756825c3f 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // InstanceTraits specialization for DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK
 //
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
index 346b731c33..572b04e75b 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // InstanceTraits specialization for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
 //
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp
index 5784938fc6..7eb4c883b0 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // InstanceTraits specialization for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3
 //
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp
index 9f64c7d6e8..60f991b1fc 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // InstanceTraits specialization for DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
 //
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp
index b41aed784c..f60d5e7b68 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // InstanceTraits specialization for DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor
 //
diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp
index b13675a7b9..95d1c94de4 100644
--- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp
+++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 // Utility functions and helpers for instance_traits.hpp
 // Contains helper functions to convert types, enums, and sequences to string representations.
@@ -21,6 +21,7 @@
 #include <ck_tile/ops/common/tensor_layout.hpp>
 #include <ck/tensor_operation/gpu/element/element_wise_operation.hpp>
 #include <ck/tensor_operation/gpu/device/convolution_forward_specialization.hpp>
+#include <ck/tensor_operation/gpu/device/convolution_backward_weight_specialization.hpp>
 #include <ck/tensor_operation/gpu/device/gemm_specialization.hpp>
 
 namespace ck_tile::reflect::detail {
@@ -112,6 +113,20 @@ conv_fwd_spec_name(ck::tensor_operation::device::ConvolutionForwardSpecializatio
     }
 }
 
+// Convert ConvolutionBackwardWeightSpecialization enum to string
+constexpr std::string_view conv_bwd_weight_spec_name(
+    ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization spec)
+{
+    using enum ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization;
+    switch(spec)
+    {
+    case Default: return "Default";
+    case Filter1x1Stride1Pad0: return "Filter1x1Stride1Pad0";
+    case Filter1x1Pad0: return "Filter1x1Pad0";
+    case OddC: return "OddC";
+    }
+}
+
 // Convert GemmSpecialization enum to string
 constexpr std::string_view gemm_spec_name(ck::tensor_operation::device::GemmSpecialization spec)
 {
diff --git a/experimental/builder/include/ck_tile/builder/types.hpp b/experimental/builder/include/ck_tile/builder/types.hpp
index 5c1b5f9958..553bf99bfb 100644
--- a/experimental/builder/include/ck_tile/builder/types.hpp
+++ b/experimental/builder/include/ck_tile/builder/types.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -68,50 +68,6 @@ enum class ConvDirection
     BACKWARD_WEIGHT
 };
 
-// Forward convolution device operations.
-enum class FwdGroupConvDeviceOperation
-{
-    DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK,
-    DeviceGroupedConvFwdMultipleD_Wmma_CShuffle,
-    DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle,
-    DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3,
-    DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor
-};
-
-// Backward data convolution device operations.
-enum class BwdDataGroupConvDeviceOperation
-{
-    DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1,
-    DeviceGroupedConvBwdDataMultipleD,
-    DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle
-};
-
-// Backward weight convolution device operations.
-enum class BwdWeightGroupConvDeviceOperation
-{
-    DeviceGroupedConvBwdWeight,
-    DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle,
-    DeviceGroupedConvBwdWeight_Xdl_CShuffle,
-    DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle,
-    DeviceGroupedConvBwdWeight_Wmma_CShuffle,
-    DeviceGroupedConvBwdWeight_Xdl_CShuffleV3,
-    DeviceGroupedConvBwdWeightMultipleD,
-    DeviceGroupedConvBwdWeight_Dl
-};
-
-// Structural type for device operation
-struct GroupConvDeviceOp {
-    union {
-        FwdGroupConvDeviceOperation _fwd;
-        BwdDataGroupConvDeviceOperation _bwd_data;
-        BwdWeightGroupConvDeviceOperation _bwd_weight;
-    };
-    
-    constexpr GroupConvDeviceOp(FwdGroupConvDeviceOperation op) : _fwd(op) {}
-    constexpr GroupConvDeviceOp(BwdDataGroupConvDeviceOperation op) : _bwd_data(op) {}
-    constexpr GroupConvDeviceOp(BwdWeightGroupConvDeviceOperation op) : _bwd_weight(op) {}
-};
-
 // Fused element-wise operations.
 enum class ElementwiseOperation
 {
diff --git a/experimental/builder/include/ck_tile/builder/versions.hpp b/experimental/builder/include/ck_tile/builder/versions.hpp
index e8fb2fe4de..4ee45b730b 100644
--- a/experimental/builder/include/ck_tile/builder/versions.hpp
+++ b/experimental/builder/include/ck_tile/builder/versions.hpp
@@ -1,3 +1,6 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
 #pragma once
 
 #include <concepts>
diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt
index 26a666a805..90a5de6d66 100644
--- a/experimental/builder/test/CMakeLists.txt
+++ b/experimental/builder/test/CMakeLists.txt
@@ -20,6 +20,7 @@ endfunction()
 add_ck_builder_test(test_ckb_conv_builder
     test_conv_builder.cpp
     test_fwd_instance_traits.cpp
+    test_bwd_weight_instance_traits.cpp
     test_instance_traits_util.cpp)
 
 add_ck_builder_test(test_ckb_inline_diff test_inline_diff.cpp)
@@ -30,7 +31,8 @@ add_ck_builder_test(test_ckb_get_instance_string
     test_get_instance_string_fwd_grp_conv.cpp
     test_get_instance_string_fwd_grp_conv_large_tensor.cpp
     test_get_instance_string_fwd_grp_conv_wmma.cpp
-    test_get_instance_string_fwd_grp_conv_dl.cpp)
+    test_get_instance_string_fwd_grp_conv_dl.cpp
+    test_get_instance_string_bwd_weight_grp_conv_xdl.cpp)
 
 # Testing the fwd convolution builder requires kernel compilation.
 # To enable parallel compilation, the individual tests are split into separate files.
@@ -43,7 +45,8 @@ add_ck_builder_test(test_ckb_build_fwd_instances
     conv/test_ckb_conv_fwd_2d_fp32.cpp
     conv/test_ckb_conv_fwd_3d_bf16.cpp
     conv/test_ckb_conv_fwd_3d_fp16.cpp
-    conv/test_ckb_conv_fwd_3d_fp32.cpp)
+    conv/test_ckb_conv_fwd_3d_fp32.cpp
+    )
 
 function(add_ck_factory_test test_name)
     add_ck_builder_test(${test_name} ${ARGN})
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp
index 472c43438d..d6cda1f427 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 1D BF16 (channels-first) with Pipeline V2 and FILTER_1X1_STRIDE1_PAD0 specialization and SCALE
 // elementwise op
 TEST(FwdConvInstances,
@@ -14,18 +18,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout1D::NGCW_GKXC_NGKW,
         .data_type             = DataType::BF16,
-        .elementwise_operation = ElementwiseOperation::SCALE,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::SCALE};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 256, .n = 256, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_256x256x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v2_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<
-        FwdConvSignature,
-        FwdThreadBlock,
-        BlockGemmPipelineVersion::V2,
-        ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 256, 256, 32",
+                       "Filter1x1Stride1Pad0",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v2"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp
index 3f840ba2b0..330db8d457 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 1D FP16 (channels-last) with DEFAULT specialization
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_Instance_1D_FP16_ChannelsFirst_scale)
@@ -13,16 +17,21 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout1D::NWGC_GKXC_NWGK,
         .data_type             = DataType::FP16,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 64,
-                                         .tile_size  = {.m = 64, .n = 32, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle FwdConvAlgorithm{
+        .thread_block               = FwdThreadBlock_64x32x32,
+        .gridwise_gemm              = FwdGemmParams_Xdl_2x1_per_wave,
+        .block_transfer             = FwdBlockTransfer_4x16x1,
+        .fwd_specialization         = ConvFwdSpecialization::DEFAULT,
+        .gemm_specialization        = GemmSpecialization::MNKPadding,
+        .num_gemm_k_prefetch_stages = 1,
+        .num_groups_to_merge        = 2,
+        .loop_scheduler             = LoopScheduler::DEFAULT};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<FwdConvSignature,
-                                                          FwdThreadBlock,
-                                                          ConvFwdSpecialization::DEFAULT>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>(
+        {"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle", "64, 64, 32, 32", "Default"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp
index 1819cca728..1ec5bbb349 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 1D I8 (channels-last) with and DEFAULT specialization
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle_Instance_1D_FP32_ChannelsFirst_scale)
@@ -13,16 +17,20 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout1D::GNWC_GKXC_GNWK,
         .data_type             = DataType::I8,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleD_Wmma_CShuffle};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 128,
-                                         .tile_size  = {.m = 64, .n = 64, .k = 64}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle FwdConvAlgorithm{
+        .thread_block               = FwdThreadBlock_64x64x64,
+        .gridwise_gemm              = FwdGemmParams_Wmma_2x1_per_wave,
+        .block_transfer             = FwdBlockTransfer_4x32x1,
+        .fwd_specialization         = ConvFwdSpecialization::DEFAULT,
+        .gemm_specialization        = GemmSpecialization::MNKPadding,
+        .num_gemm_k_prefetch_stages = 1,
+        .loop_scheduler             = LoopScheduler::DEFAULT};
 
-    run_test_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle<FwdConvSignature,
-                                                         FwdThreadBlock,
-                                                         ConvFwdSpecialization::DEFAULT>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>(
+        {"DeviceGroupedConvFwdMultipleD_Wmma_CShuffle", "128, 64, 64, 64", "Default"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp
index b9969f7e95..31f2976fd0 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 2D BF16 NHWGC (channels-last) with Pipeline V1 and DEFAULT
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_2D_BF16_ChannelsLast)
@@ -13,17 +17,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout2D::NHWGC_GKYXC_NHWGK,
         .data_type             = DataType::BF16,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 256, .n = 256, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_256x256x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::DEFAULT,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v1_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<FwdConvSignature,
-                                                             FwdThreadBlock,
-                                                             BlockGemmPipelineVersion::V1,
-                                                             ConvFwdSpecialization::DEFAULT>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 256, 256, 32",
+                       "Default",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v1"});
 }
 
 // 2D BF16 NHWGC (channels-last) with Pipeline V5 and FILTER_3x3
@@ -35,17 +44,20 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout2D::NHWGC_GKYXC_NHWGK,
         .data_type             = DataType::BF16,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 256, .n = 256, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_256x256x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::FILTER_3x3,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v5_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<FwdConvSignature,
-                                                             FwdThreadBlock,
-                                                             BlockGemmPipelineVersion::V5,
-                                                             ConvFwdSpecialization::FILTER_3x3>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "Filter3x3",
+                       "BlkGemmPipelineVersion: v5"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp
index cd5186cc10..6276424a77 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_2D_FP16_GNHWC)
 {
@@ -12,18 +16,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout2D::GNHWC_GKYXC_GNHWK,
         .data_type             = DataType::FP16,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 256, .n = 256, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_256x256x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::FILTER_1X1_PAD0,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v3_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<
-        FwdConvSignature,
-        FwdThreadBlock,
-        BlockGemmPipelineVersion::V3,
-        ConvFwdSpecialization::FILTER_1X1_PAD0>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 256, 256, 32",
+                       "Filter1x1Pad0",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v3"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp
index 584e0ab182..a390510199 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_2D_FP32_NGCHW_GKCYX)
 {
@@ -12,18 +16,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout2D::NGCHW_GKCYX_NGKHW,
         .data_type             = DataType::FP32,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 128, .n = 128, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_128x128x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v4_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<
-        FwdConvSignature,
-        FwdThreadBlock,
-        BlockGemmPipelineVersion::V4,
-        ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 128, 128, 32",
+                       "Filter1x1Stride1Pad0",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v4"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp
index 17caf98457..3c59ae24fb 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 3D BF16 GNDHWC (group-first, channels-last) with Pipeline V3 and DEFAULT
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_3D_BF16_GNDHWC)
@@ -13,17 +17,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout3D::GNDHWC_GKZYXC_GNDHWK,
         .data_type             = DataType::BF16,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 256, .n = 256, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_256x256x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::DEFAULT,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v3_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<FwdConvSignature,
-                                                             FwdThreadBlock,
-                                                             BlockGemmPipelineVersion::V3,
-                                                             ConvFwdSpecialization::DEFAULT>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 256, 256, 32",
+                       "Default",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v3"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp
index ec4649a6ff..14d2811918 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 3D FP16 NDHWGC (channels-last) with Pipeline V4 and FILTER_1X1_PAD0
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_3D_FP16_NDHWGC_ChannelsLast)
@@ -13,18 +17,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout3D::NDHWGC_GKZYXC_NDHWGK,
         .data_type             = DataType::FP16,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 128, .n = 128, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_128x128x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::FILTER_1X1_PAD0,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v4_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<
-        FwdConvSignature,
-        FwdThreadBlock,
-        BlockGemmPipelineVersion::V4,
-        ConvFwdSpecialization::FILTER_1X1_PAD0>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 128, 128, 32",
+                       "Filter1x1Pad0",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v4"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp
index 393ea9206d..bce092d5f6 100644
--- a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp
+++ b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp
@@ -1,9 +1,13 @@
-#include "utils/ckb_conv_test_common.hpp"
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "utils/ckb_conv_test_configs.hpp"
+#include "utils/ckb_conv_test_utils.hpp"
+
+namespace {
 
 using namespace ck_tile::builder::test_utils;
 
-namespace ck_tile::builder::testing {
-
 // 3D FP32 NGCDHW (channels-first) with Pipeline V1 and FILTER_1X1_PAD0
 TEST(FwdConvInstances,
      Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_3D_FP32_ChannelsFirst)
@@ -13,18 +17,22 @@ TEST(FwdConvInstances,
         .direction             = ConvDirection::FORWARD,
         .layout                = GroupConvLayout3D::NGCDHW_GKCZYX_NGKDHW,
         .data_type             = DataType::FP32,
-        .elementwise_operation = ElementwiseOperation::PASS_THROUGH,
-        .device_operation =
-            FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3};
+        .elementwise_operation = ElementwiseOperation::PASS_THROUGH};
 
-    constexpr ThreadBlock FwdThreadBlock{.block_size = 256,
-                                         .tile_size  = {.m = 128, .n = 128, .k = 32}};
+    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
+        .thread_block        = FwdThreadBlock_256x256x32,
+        .gridwise_gemm       = FwdGemmParams_Xdl_4x4_per_wave,
+        .block_transfer      = FwdBlockTransfer_4x64_1,
+        .fwd_specialization  = ConvFwdSpecialization::FILTER_1X1_PAD0,
+        .gemm_specialization = GemmSpecialization::MNKPadding,
+        .block_gemm          = BlockGemmDesc_v1_intrawave};
 
-    run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<
-        FwdConvSignature,
-        FwdThreadBlock,
-        BlockGemmPipelineVersion::V1,
-        ConvFwdSpecialization::FILTER_1X1_PAD0>();
+    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
+    run_test<Builder>({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3",
+                       "256, 256, 256, 32",
+                       "Filter1x1Pad0",
+                       "BlkGemmPipelineScheduler: Intrawave",
+                       "BlkGemmPipelineVersion: v1"});
 }
 
-} // namespace ck_tile::builder::testing
+} // namespace
diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp
index 9c5ca9b97b..e719db89ed 100644
--- a/experimental/builder/test/impl/conv_algorithm_types.hpp
+++ b/experimental/builder/test/impl/conv_algorithm_types.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -126,26 +126,6 @@ struct ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3
     GemmSpecialization gemm_specialization;
     BlockGemm block_gemm;
 };
-static_assert(
-    ckb::ConvAlgorithmDescriptor<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(
-    ckb::SpecifiesThreadBlock<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(
-    ckb::SpecifiesGridwiseXdlGemm<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(
-    ckb::SpecifiesBlockTransfer<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(
-    ckb::SpecifiesLdsTransfer<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(ckb::SpecifiesThreadClusterAccessOrder<
-              ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(
-    ckb::SpecifiesSourceAccessOrder<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(ckb::SpecifiesFwdConcSpecialization<
-              ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(
-    ckb::SpecifiesBlockGemm<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
-static_assert(ckb::SpecifiesGemmSpecialization<
-              ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>);
 
 struct ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
 {
@@ -158,30 +138,6 @@ struct ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
     size_t num_groups_to_merge;
     LoopScheduler loop_scheduler;
 };
-static_assert(
-    ckb::ConvAlgorithmDescriptor<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesThreadBlock<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesGridwiseXdlGemm<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesBlockTransfer<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesLdsTransfer<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(ckb::SpecifiesThreadClusterAccessOrder<
-              ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesSourceAccessOrder<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(ckb::SpecifiesFwdConcSpecialization<
-              ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesNumPrefetchStages<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesGemmSpecialization<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesLoopScheduler<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
-static_assert(
-    ckb::SpecifiesNumGroupsToMerge<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>);
 
 struct ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
 {
@@ -193,25 +149,5 @@ struct ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
     size_t num_gemm_k_prefetch_stages;
     LoopScheduler loop_scheduler;
 };
-static_assert(
-    ckb::ConvAlgorithmDescriptor<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(ckb::SpecifiesThreadBlock<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesGridwiseWmmaGemm<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesBlockTransfer<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(ckb::SpecifiesLdsTransfer<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(ckb::SpecifiesThreadClusterAccessOrder<
-              ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesSourceAccessOrder<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesFwdConcSpecialization<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesNumPrefetchStages<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesGemmSpecialization<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
-static_assert(
-    ckb::SpecifiesLoopScheduler<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>);
 
 } // namespace ck_tile::builder::test
diff --git a/experimental/builder/test/impl/conv_signature_types.hpp b/experimental/builder/test/impl/conv_signature_types.hpp
index cc5490c711..71f16aefbe 100644
--- a/experimental/builder/test/impl/conv_signature_types.hpp
+++ b/experimental/builder/test/impl/conv_signature_types.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -17,7 +17,6 @@ struct ConvSignature
     GroupConvLayout layout;
     DataType data_type;
     ElementwiseOperation elementwise_operation;
-    GroupConvDeviceOp device_operation;
 };
 static_assert(ConvSignatureDescriptor<ConvSignature>);
 
diff --git a/experimental/builder/test/test_bwd_weight_instance_traits.cpp b/experimental/builder/test/test_bwd_weight_instance_traits.cpp
new file mode 100644
index 0000000000..24c28c2b9d
--- /dev/null
+++ b/experimental/builder/test/test_bwd_weight_instance_traits.cpp
@@ -0,0 +1,112 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <gtest/gtest.h>
+#include <ck/ck.hpp>
+#include <ck_tile/builder/reflect/instance_traits.hpp>
+#include <ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp>
+
+namespace {
+
+TEST(InstanceTraits, BwdWeightXdlCShuffleInstanceStringReturnsCorrectFormat)
+{
+    using DeviceInstance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle<
+        2,                                               // NDimSpatial
+        ck::tensor_layout::convolution::GNHWC,           // InLayout
+        ck::tensor_layout::convolution::GKYXC,           // WeiLayout
+        ck::tensor_layout::convolution::GNHWK,           // OutLayout
+        ck::half_t,                                      // InDataType
+        ck::half_t,                                      // WeiDataType
+        ck::half_t,                                      // OutDataType
+        float,                                           // AccDataType
+        ck::tensor_operation::element_wise::PassThrough, // InElementwiseOperation
+        ck::tensor_operation::element_wise::PassThrough, // WeiElementwiseOperation
+        ck::tensor_operation::element_wise::PassThrough, // OutElementwiseOperation
+        ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization::
+            Default,            // ConvBackwardWeightSpecialization
+        256,                    // BlockSize
+        128,                    // MPerBlock
+        128,                    // NPerBlock
+        4,                      // K0PerBlock
+        8,                      // K1
+        32,                     // MPerXDL
+        32,                     // NPerXDL
+        2,                      // MXdlPerWave
+        2,                      // NXdlPerWave
+        ck::Sequence<4, 64, 1>, // ABlockTransferThreadClusterLengths_K0_M_K1
+        ck::Sequence<1, 0, 2>,  // ABlockTransferThreadClusterArrangeOrder
+        ck::Sequence<1, 0, 2>,  // ABlockTransferSrcAccessOrder
+        2,                      // ABlockTransferSrcVectorDim
+        8,                      // ABlockTransferSrcScalarPerVector
+        8,                      // ABlockTransferDstScalarPerVector_K1
+        false,                  // ABlockLdsAddExtraM
+        ck::Sequence<4, 64, 1>, // BBlockTransferThreadClusterLengths_K0_N_K1
+        ck::Sequence<1, 0, 2>,  // BBlockTransferThreadClusterArrangeOrder
+        ck::Sequence<1, 0, 2>,  // BBlockTransferSrcAccessOrder
+        2,                      // BBlockTransferSrcVectorDim
+        8,                      // BBlockTransferSrcScalarPerVector
+        8,                      // BBlockTransferDstScalarPerVector_K1
+        false,                  // BBlockLdsAddExtraN
+        1,                      // CShuffleMXdlPerWavePerShuffle
+        1,                      // CShuffleNXdlPerWavePerShuffle
+        ck::Sequence<1,
+                     32,
+                     1,
+                     8>, // CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
+        8,               // CBlockTransferScalarPerVector_NWaveNPerXdl
+        ck::half_t,      // ComputeTypeA
+        ck::half_t,      // ComputeTypeB
+        1,               // MaxTransposeTransferSrcScalarPerVector
+        1>;              // MaxTransposeTransferDstScalarPerVector
+
+    std::string instance_str = ck_tile::reflect::instance_string<DeviceInstance>();
+
+    std::string expected_str = "DeviceGroupedConvBwdWeight_Xdl_CShuffle"
+                               "<2"             // NDimSpatial
+                               ",GNHWC"         // InLayout
+                               ",GKYXC"         // WeiLayout
+                               ",GNHWK"         // OutLayout
+                               ",fp16"          // InDataType
+                               ",fp16"          // WeiDataType
+                               ",fp16"          // OutDataType
+                               ",fp32"          // AccDataType
+                               ",PassThrough"   // InElementwiseOperation
+                               ",PassThrough"   // WeiElementwiseOperation
+                               ",PassThrough"   // OutElementwiseOperation
+                               ",Default"       // ConvBackwardWeightSpecialization
+                               ",256"           // BlockSize
+                               ",128"           // MPerBlock
+                               ",128"           // NPerBlock
+                               ",4"             // K0PerBlock
+                               ",8"             // K1
+                               ",32"            // MPerXDL
+                               ",32"            // NPerXDL
+                               ",2"             // MXdlPerWave
+                               ",2"             // NXdlPerWave
+                               ",Seq(4,64,1)"   // ABlockTransferThreadClusterLengths_K0_M_K1
+                               ",Seq(1,0,2)"    // ABlockTransferThreadClusterArrangeOrder
+                               ",Seq(1,0,2)"    // ABlockTransferSrcAccessOrder
+                               ",2"             // ABlockTransferSrcVectorDim
+                               ",8"             // ABlockTransferSrcScalarPerVector
+                               ",8"             // ABlockTransferDstScalarPerVector_K1
+                               ",false"         // ABlockLdsAddExtraM
+                               ",Seq(4,64,1)"   // BBlockTransferThreadClusterLengths_K0_N_K1
+                               ",Seq(1,0,2)"    // BBlockTransferThreadClusterArrangeOrder
+                               ",Seq(1,0,2)"    // BBlockTransferSrcAccessOrder
+                               ",2"             // BBlockTransferSrcVectorDim
+                               ",8"             // BBlockTransferSrcScalarPerVector
+                               ",8"             // BBlockTransferDstScalarPerVector_K1
+                               ",false"         // BBlockLdsAddExtraN
+                               ",1"             // CShuffleMXdlPerWavePerShuffle
+                               ",1"             // CShuffleNXdlPerWavePerShuffle
+                               ",Seq(1,32,1,8)" // CBlockTransferClusterLengths
+                               ",8"             // CBlockTransferScalarPerVector_NWaveNPerXdl
+                               ",fp16"          // ComputeTypeA
+                               ",fp16"          // ComputeTypeB
+                               ",1"             // MaxTransposeTransferSrcScalarPerVector
+                               ",1>";           // MaxTransposeTransferDstScalarPerVector
+
+    EXPECT_EQ(instance_str, expected_str);
+}
+
+} // anonymous namespace
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp
index 53199f1d77..73a731a69c 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp>
 
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp
index c68fdb4b24..75cb58018e 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_bias_bnorm_clamp.hpp>
 
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp
index a9e5e39e8f..b301bab966 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_bias_clamp.hpp>
 
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp
index cc9cbd3dda..8b77b11e0c 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_bilinear.hpp>
 #include "ck/utility/data_type.hpp"
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp
index a8d1d1763f..e678fa1258 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_clamp.hpp>
 #include "ck/utility/data_type.hpp"
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp
index 0da300e3d8..4882c6fde2 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_convscale.hpp>
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_convscale_relu.hpp>
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp
index e918785fd7..7437385e2a 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_dynamic_op.hpp>
 #include "ck/utility/data_type.hpp"
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp
index 428d1c81f3..8144d7bedd 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scale.hpp>
 
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp
index 774c30c05e..beebc3f853 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scaleadd_ab.hpp>
 
diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp
index ba8726a643..79feb298cd 100644
--- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp
+++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scaleadd_scaleadd_relu.hpp>
 
diff --git a/experimental/builder/test/test_conv_builder.cpp b/experimental/builder/test/test_conv_builder.cpp
index 4ec189daa4..27b30946b2 100644
--- a/experimental/builder/test/test_conv_builder.cpp
+++ b/experimental/builder/test/test_conv_builder.cpp
@@ -1,3 +1,6 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
 #include <gtest/gtest.h>
 
 class ConvBuilderTest : public ::testing::Test
diff --git a/experimental/builder/test/test_fwd_instance_traits.cpp b/experimental/builder/test/test_fwd_instance_traits.cpp
index 2e3b6ac264..b57b20eb7d 100644
--- a/experimental/builder/test/test_fwd_instance_traits.cpp
+++ b/experimental/builder/test/test_fwd_instance_traits.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <gmock/gmock.h>
diff --git a/experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp b/experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp
new file mode 100644
index 0000000000..68b43c6a99
--- /dev/null
+++ b/experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp
@@ -0,0 +1,86 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <gtest/gtest.h>
+#include <ck_tile/builder/reflect/instance_traits.hpp>
+#include <ck/tensor_operation/gpu/device/device_base.hpp>
+#include <ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_xdl_instance.hpp>
+
+// Test GetInstanceString through base class pointer for backward weight XDL variant
+TEST(GetInstanceString, ReturnsStringForBwdWeightGrpConvXdlInstance)
+{
+    // Use the template helper to get a working instance configuration
+    using InstanceTuple = ck::tensor_operation::device::instance::
+        device_grouped_conv_bwd_weight_xdl_c_shuffle_f16_instances<
+            2,                                             // NDimSpatial
+            ck::tensor_operation::device::instance::GNHWC, // InLayout
+            ck::tensor_operation::device::instance::GKYXC, // WeiLayout
+            ck::tensor_operation::device::instance::GNHWK, // OutLayout
+            ck::tensor_operation::device::instance::
+                ConvBwdWeightDefault>; // ConvBwdWeightSpecialization
+
+    // Get the first instance from the tuple
+    using DeviceInstance = typename std::tuple_element<0, InstanceTuple>::type;
+
+    // Define the base class type using the most general operator base
+    using BaseClass = ck::tensor_operation::device::BaseOperator;
+
+    // Create an instance of the derived class
+    DeviceInstance device_instance;
+
+    // Get a pointer to the base class
+    BaseClass* base_ptr = &device_instance;
+
+    // Call GetInstanceString through the base class pointer
+    std::string instance_str = base_ptr->GetInstanceString();
+
+    // Expected complete instance string based on the first instance from
+    // device_grouped_conv_bwd_weight_xdl_c_shuffle_f16_instances
+    // This corresponds to the configuration with BlockSize=64, MPerBlock=64, NPerBlock=64, etc.
+    std::string expected_str = "DeviceGroupedConvBwdWeight_Xdl_CShuffle"
+                               "<2"             // NDimSpatial
+                               ",GNHWC"         // InLayout
+                               ",GKYXC"         // WeiLayout
+                               ",GNHWK"         // OutLayout
+                               ",fp16"          // InDataType
+                               ",fp16"          // WeiDataType
+                               ",fp16"          // OutDataType
+                               ",fp32"          // AccDataType
+                               ",PassThrough"   // InElementwiseOperation
+                               ",PassThrough"   // WeiElementwiseOperation
+                               ",PassThrough"   // OutElementwiseOperation
+                               ",Default"       // ConvBackwardWeightSpecialization
+                               ",64"            // BlockSize
+                               ",64"            // MPerBlock
+                               ",64"            // NPerBlock
+                               ",4"             // K0PerBlock
+                               ",8"             // K1
+                               ",32"            // MPerXDL
+                               ",32"            // NPerXDL
+                               ",2"             // MXdlPerWave
+                               ",2"             // NXdlPerWave
+                               ",Seq(1,4,8,2)"  // ABlockTransferThreadClusterLengths_K0_M_K1
+                               ",Seq(0,3,1,2)"  // ABlockTransferThreadClusterArrangeOrder
+                               ",Seq(0,2,1,3)"  // ABlockTransferSrcAccessOrder
+                               ",2"             // ABlockTransferSrcVectorDim
+                               ",2"             // ABlockTransferSrcScalarPerVector
+                               ",4"             // ABlockTransferDstScalarPerVector_K1
+                               ",true"          // ABlockLdsAddExtraM
+                               ",Seq(1,4,8,2)"  // BBlockTransferThreadClusterLengths_K0_N_K1
+                               ",Seq(0,3,1,2)"  // BBlockTransferThreadClusterArrangeOrder
+                               ",Seq(0,2,1,3)"  // BBlockTransferSrcAccessOrder
+                               ",2"             // BBlockTransferSrcVectorDim
+                               ",2"             // BBlockTransferSrcScalarPerVector
+                               ",4"             // BBlockTransferDstScalarPerVector_K1
+                               ",true"          // BBlockLdsAddExtraN
+                               ",1"             // CShuffleMXdlPerWavePerShuffle
+                               ",1"             // CShuffleNXdlPerWavePerShuffle
+                               ",Seq(1,16,1,4)" // CBlockTransferClusterLengths
+                               ",2"             // CBlockTransferScalarPerVector_NWaveNPerXdl
+                               ",fp16"          // ComputeTypeA
+                               ",fp16"          // ComputeTypeB
+                               ",1"             // MaxTransposeTransferSrcScalarPerVector
+                               ",1>";           // MaxTransposeTransferDstScalarPerVector
+
+    EXPECT_EQ(instance_str, expected_str);
+}
diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp
index 7c43c428ce..ca683abedc 100644
--- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp
+++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <ck_tile/builder/reflect/instance_traits.hpp>
diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp
index 54e026b308..abf55322b6 100644
--- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp
+++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <ck_tile/builder/reflect/instance_traits.hpp>
diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp
index fd75fd9e6b..d6900d041f 100644
--- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp
+++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <ck_tile/builder/reflect/instance_traits.hpp>
diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp
index 45ab5db322..67e55173de 100644
--- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp
+++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <ck_tile/builder/reflect/instance_traits.hpp>
diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp
index e1505659f5..ccce080f9d 100644
--- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp
+++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <ck_tile/builder/reflect/instance_traits.hpp>
diff --git a/experimental/builder/test/test_inline_diff.cpp b/experimental/builder/test/test_inline_diff.cpp
index 41692fb40e..7a1b0ab8c3 100644
--- a/experimental/builder/test/test_inline_diff.cpp
+++ b/experimental/builder/test/test_inline_diff.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 
diff --git a/experimental/builder/test/test_instance_traits_util.cpp b/experimental/builder/test/test_instance_traits_util.cpp
index 04d3cb978b..5623028699 100644
--- a/experimental/builder/test/test_instance_traits_util.cpp
+++ b/experimental/builder/test/test_instance_traits_util.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <gtest/gtest.h>
 #include <gmock/gmock.h>
diff --git a/experimental/builder/test/test_testing_utils.cpp b/experimental/builder/test/test_testing_utils.cpp
index 24a1c9bc81..89d8c06e59 100644
--- a/experimental/builder/test/test_testing_utils.cpp
+++ b/experimental/builder/test/test_testing_utils.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp>
 
diff --git a/experimental/builder/test/testing_utils.cpp b/experimental/builder/test/testing_utils.cpp
index 34793b601e..02663317f1 100644
--- a/experimental/builder/test/testing_utils.cpp
+++ b/experimental/builder/test/testing_utils.cpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include "testing_utils.hpp"
 #include <gtest/gtest.h>
diff --git a/experimental/builder/test/testing_utils.hpp b/experimental/builder/test/testing_utils.hpp
index 3ff2eb32de..ae5811bd4b 100644
--- a/experimental/builder/test/testing_utils.hpp
+++ b/experimental/builder/test/testing_utils.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <ck/library/tensor_operation_instance/device_operation_instance_factory.hpp>
 #include <gtest/gtest.h>
diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_common.hpp
deleted file mode 100644
index d18a008015..0000000000
--- a/experimental/builder/test/utils/ckb_conv_test_common.hpp
+++ /dev/null
@@ -1,238 +0,0 @@
-// SPDX-License-Identifier: MIT
-// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
-
-#pragma once
-
-#include <gtest/gtest.h>
-#include "impl/conv_algorithm_types.hpp"
-#include "impl/conv_signature_types.hpp"
-#include "ck_tile/builder/conv_builder.hpp"
-
-namespace ck_tile::builder::test_utils {
-
-using namespace ck_tile::builder;
-using namespace test;
-
-// Common test implementation
-template <ConvSignature FwdConvSignature,
-          ThreadBlock FwdThreadBlock,
-          BlockGemmPipelineVersion FwdPipelineVersion,
-          ConvFwdSpecialization FwdConvSpecialization>
-constexpr void run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3()
-{
-    constexpr GridwiseXdlGemm FwdGemmParams{.ak1            = 8,
-                                            .bk1            = 8,
-                                            .m_per_xdl      = 32,
-                                            .n_per_xdl      = 32,
-                                            .m_xdl_per_wave = 4,
-                                            .n_xdl_per_wave = 4};
-
-    constexpr BlockTransferABC FwdBlockTransfer{.block_transfer_a = {.k0 = 4, .m_n = 64, .k1 = 1},
-                                                .block_transfer_b = {.k0 = 4, .m_n = 64, .k1 = 1},
-                                                .thread_cluster_dims_c = {.m_block        = 1,
-                                                                          .m_wave_per_xdl = 32,
-                                                                          .n_block        = 1,
-                                                                          .n_wave_per_xdl = 8},
-                                                .lds_transfer_a        = {.src_vector_dim            = 2,
-                                                                          .src_scalar_per_vector     = 2,
-                                                                          .lds_dst_scalar_per_vector = 8,
-                                                                          .is_direct_load = false,
-                                                                          .lds_padding    = false},
-                                                .lds_transfer_b        = {.src_vector_dim            = 2,
-                                                                          .src_scalar_per_vector     = 8,
-                                                                          .lds_dst_scalar_per_vector = 8,
-                                                                          .is_direct_load = false,
-                                                                          .lds_padding    = false},
-                                                .epilogue_c = {.m_per_wave_per_shuffle = 1,
-                                                               .n_per_wave_per_shuffle = 1,
-                                                               .scalar_per_vector      = 8},
-                                                .block_transfer_access_order_a = {1, 0, 2},
-                                                .block_transfer_access_order_b = {1, 0, 2},
-                                                .src_access_order_a            = {1, 0, 2},
-                                                .src_access_order_b            = {1, 0, 2}};
-
-    constexpr BlockGemm BlockGemmDesc = {.pipeline_version = FwdPipelineVersion,
-                                         .scheduler        = BlockGemmPipelineScheduler::INTRAWAVE};
-
-    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{
-        .thread_block        = FwdThreadBlock,
-        .gridwise_gemm       = FwdGemmParams,
-        .block_transfer      = FwdBlockTransfer,
-        .fwd_specialization  = FwdConvSpecialization,
-        .gemm_specialization = GemmSpecialization::MNKPadding,
-        .block_gemm          = BlockGemmDesc};
-
-    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
-
-    auto instance = typename Builder::Instance{};
-
-    const auto kernel_string = instance.GetTypeString();
-    std::cout << "Generated kernel: " << kernel_string << std::endl;
-    EXPECT_GT(kernel_string.size(), 0);
-
-    EXPECT_TRUE(kernel_string.starts_with("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3"));
-
-    // Verify pipeline version is correct
-    if(FwdPipelineVersion == BlockGemmPipelineVersion::V1)
-        EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v1") != std::string::npos);
-    else if(FwdPipelineVersion == BlockGemmPipelineVersion::V3)
-        EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v3") != std::string::npos);
-    else if(FwdPipelineVersion == BlockGemmPipelineVersion::V4)
-        EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v4") != std::string::npos);
-    else if(FwdPipelineVersion == BlockGemmPipelineVersion::V5)
-        EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v5") != std::string::npos);
-
-    // Verify specialization is correct
-    if(FwdConvSpecialization == ConvFwdSpecialization::DEFAULT)
-        EXPECT_TRUE(kernel_string.find("Default") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_PAD0)
-        EXPECT_TRUE(kernel_string.find("Filter1x1Pad0") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0)
-        EXPECT_TRUE(kernel_string.find("Filter1x1Stride1Pad0") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_3x3)
-        EXPECT_TRUE(kernel_string.find("Filter3x3") != std::string::npos);
-
-    const auto invoker_ptr = instance.MakeInvokerPointer();
-    EXPECT_NE(invoker_ptr, nullptr);
-}
-
-template <ConvSignature FwdConvSignature,
-          ThreadBlock FwdThreadBlock,
-          ConvFwdSpecialization FwdConvSpecialization>
-constexpr void run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle()
-{
-    constexpr GridwiseXdlGemm FwdGemmParams{.ak1            = 8,
-                                            .bk1            = 8,
-                                            .m_per_xdl      = 32,
-                                            .n_per_xdl      = 32,
-                                            .m_xdl_per_wave = 2,
-                                            .n_xdl_per_wave = 1};
-
-    constexpr BlockTransferABC FwdBlockTransfer{.block_transfer_a = {.k0 = 4, .m_n = 16, .k1 = 1},
-                                                .block_transfer_b = {.k0 = 4, .m_n = 16, .k1 = 1},
-                                                .thread_cluster_dims_c = {.m_block        = 1,
-                                                                          .m_wave_per_xdl = 16,
-                                                                          .n_block        = 1,
-                                                                          .n_wave_per_xdl = 4},
-                                                .lds_transfer_a        = {.src_vector_dim            = 2,
-                                                                          .src_scalar_per_vector     = 8,
-                                                                          .lds_dst_scalar_per_vector = 8,
-                                                                          .is_direct_load = false,
-                                                                          .lds_padding    = true},
-                                                .lds_transfer_b        = {.src_vector_dim            = 2,
-                                                                          .src_scalar_per_vector     = 8,
-                                                                          .lds_dst_scalar_per_vector = 8,
-                                                                          .is_direct_load = false,
-                                                                          .lds_padding    = true},
-                                                .epilogue_c = {.m_per_wave_per_shuffle = 1,
-                                                               .n_per_wave_per_shuffle = 1,
-                                                               .scalar_per_vector      = 8},
-                                                .block_transfer_access_order_a = {1, 0, 2},
-                                                .block_transfer_access_order_b = {1, 0, 2},
-                                                .src_access_order_a            = {1, 0, 2},
-                                                .src_access_order_b            = {1, 0, 2}};
-
-    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle FwdConvAlgorithm{
-        .thread_block               = FwdThreadBlock,
-        .gridwise_gemm              = FwdGemmParams,
-        .block_transfer             = FwdBlockTransfer,
-        .fwd_specialization         = FwdConvSpecialization,
-        .gemm_specialization        = GemmSpecialization::MNKPadding,
-        .num_gemm_k_prefetch_stages = 1,
-        .num_groups_to_merge        = 2,
-        .loop_scheduler             = LoopScheduler::DEFAULT};
-
-    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
-
-    auto instance = typename Builder::Instance{};
-
-    const auto kernel_string = instance.GetTypeString();
-    std::cout << "Generated kernel: " << kernel_string << std::endl;
-    EXPECT_GT(kernel_string.size(), 0);
-
-    EXPECT_TRUE(kernel_string.starts_with("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle"));
-
-    // Verify specialization is correct
-    if(FwdConvSpecialization == ConvFwdSpecialization::DEFAULT)
-        EXPECT_TRUE(kernel_string.find("Default") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_PAD0)
-        EXPECT_TRUE(kernel_string.find("Filter1x1Pad0") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0)
-        EXPECT_TRUE(kernel_string.find("Filter1x1Stride1Pad0") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_3x3)
-        EXPECT_TRUE(kernel_string.find("Filter3x3") != std::string::npos);
-
-    const auto invoker_ptr = instance.MakeInvokerPointer();
-    EXPECT_NE(invoker_ptr, nullptr);
-}
-
-template <ConvSignature FwdConvSignature,
-          ThreadBlock FwdThreadBlock,
-          ConvFwdSpecialization FwdConvSpecialization>
-constexpr void run_test_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle()
-{
-    constexpr GridwiseWmmaGemm FwdGemmParams{.k1               = 8,
-                                             .m_per_wmma       = 32,
-                                             .n_per_wmma       = 32,
-                                             .m_wmma_per_wave  = 2,
-                                             .n_wmma_per_wave  = 1,
-                                             .pipeline_version = GridwiseGemmPipelineVersion::V1};
-
-    constexpr BlockTransferABC FwdBlockTransfer{.block_transfer_a = {.k0 = 4, .m_n = 32, .k1 = 1},
-                                                .block_transfer_b = {.k0 = 4, .m_n = 32, .k1 = 1},
-                                                .thread_cluster_dims_c = {.m_block        = 1,
-                                                                          .m_wave_per_xdl = 32,
-                                                                          .n_block        = 1,
-                                                                          .n_wave_per_xdl = 4},
-                                                .lds_transfer_a        = {.src_vector_dim            = 2,
-                                                                          .src_scalar_per_vector     = 16,
-                                                                          .lds_dst_scalar_per_vector = 16,
-                                                                          .is_direct_load = false,
-                                                                          .lds_padding    = true},
-                                                .lds_transfer_b        = {.src_vector_dim            = 2,
-                                                                          .src_scalar_per_vector     = 16,
-                                                                          .lds_dst_scalar_per_vector = 16,
-                                                                          .is_direct_load = false,
-                                                                          .lds_padding    = true},
-                                                .epilogue_c = {.m_per_wave_per_shuffle = 1,
-                                                               .n_per_wave_per_shuffle = 1,
-                                                               .scalar_per_vector      = 8},
-                                                .block_transfer_access_order_a = {1, 0, 2},
-                                                .block_transfer_access_order_b = {1, 0, 2},
-                                                .src_access_order_a            = {1, 0, 2},
-                                                .src_access_order_b            = {1, 0, 2}};
-
-    constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle FwdConvAlgorithm{
-        .thread_block               = FwdThreadBlock,
-        .gridwise_gemm              = FwdGemmParams,
-        .block_transfer             = FwdBlockTransfer,
-        .fwd_specialization         = FwdConvSpecialization,
-        .gemm_specialization        = GemmSpecialization::MNKPadding,
-        .num_gemm_k_prefetch_stages = 1,
-        .loop_scheduler             = LoopScheduler::DEFAULT};
-
-    using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
-
-    auto instance = typename Builder::Instance{};
-
-    const auto kernel_string = instance.GetTypeString();
-    std::cout << "Generated kernel: " << kernel_string << std::endl;
-    EXPECT_GT(kernel_string.size(), 0);
-
-    EXPECT_TRUE(kernel_string.starts_with("DeviceGroupedConvFwdMultipleD_Wmma_CShuffle"));
-
-    // Verify specialization is correct
-    if(FwdConvSpecialization == ConvFwdSpecialization::DEFAULT)
-        EXPECT_TRUE(kernel_string.find("Default") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_PAD0)
-        EXPECT_TRUE(kernel_string.find("Filter1x1Pad0") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0)
-        EXPECT_TRUE(kernel_string.find("Filter1x1Stride1Pad0") != std::string::npos);
-    else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_3x3)
-        EXPECT_TRUE(kernel_string.find("Filter3x3") != std::string::npos);
-
-    const auto invoker_ptr = instance.MakeInvokerPointer();
-    EXPECT_NE(invoker_ptr, nullptr);
-}
-
-} // namespace ck_tile::builder::test_utils
diff --git a/experimental/builder/test/utils/ckb_conv_test_configs.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp
new file mode 100644
index 0000000000..16c667b64b
--- /dev/null
+++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp
@@ -0,0 +1,136 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include "impl/conv_algorithm_types.hpp"
+#include "impl/conv_signature_types.hpp"
+#include "ck_tile/builder/conv_builder.hpp"
+
+namespace ck_tile::builder::test_utils {
+
+using namespace ck_tile::builder;
+using namespace test;
+
+constexpr BlockTransferABC FwdBlockTransfer_4x64_1{
+    .block_transfer_a              = {.k0 = 4, .m_n = 64, .k1 = 1},
+    .block_transfer_b              = {.k0 = 4, .m_n = 64, .k1 = 1},
+    .thread_cluster_dims_c         = {.m_block        = 1,
+                                      .m_wave_per_xdl = 32,
+                                      .n_block        = 1,
+                                      .n_wave_per_xdl = 8},
+    .lds_transfer_a                = {.src_vector_dim            = 2,
+                                      .src_scalar_per_vector     = 2,
+                                      .lds_dst_scalar_per_vector = 8,
+                                      .is_direct_load            = false,
+                                      .lds_padding               = false},
+    .lds_transfer_b                = {.src_vector_dim            = 2,
+                                      .src_scalar_per_vector     = 8,
+                                      .lds_dst_scalar_per_vector = 8,
+                                      .is_direct_load            = false,
+                                      .lds_padding               = false},
+    .epilogue_c                    = {.m_per_wave_per_shuffle = 1,
+                                      .n_per_wave_per_shuffle = 1,
+                                      .scalar_per_vector      = 8},
+    .block_transfer_access_order_a = {1, 0, 2},
+    .block_transfer_access_order_b = {1, 0, 2},
+    .src_access_order_a            = {1, 0, 2},
+    .src_access_order_b            = {1, 0, 2}};
+
+constexpr BlockTransferABC FwdBlockTransfer_4x16x1{
+    .block_transfer_a              = {.k0 = 4, .m_n = 16, .k1 = 1},
+    .block_transfer_b              = {.k0 = 4, .m_n = 16, .k1 = 1},
+    .thread_cluster_dims_c         = {.m_block        = 1,
+                                      .m_wave_per_xdl = 16,
+                                      .n_block        = 1,
+                                      .n_wave_per_xdl = 4},
+    .lds_transfer_a                = {.src_vector_dim            = 2,
+                                      .src_scalar_per_vector     = 8,
+                                      .lds_dst_scalar_per_vector = 8,
+                                      .is_direct_load            = false,
+                                      .lds_padding               = true},
+    .lds_transfer_b                = {.src_vector_dim            = 2,
+                                      .src_scalar_per_vector     = 8,
+                                      .lds_dst_scalar_per_vector = 8,
+                                      .is_direct_load            = false,
+                                      .lds_padding               = true},
+    .epilogue_c                    = {.m_per_wave_per_shuffle = 1,
+                                      .n_per_wave_per_shuffle = 1,
+                                      .scalar_per_vector      = 8},
+    .block_transfer_access_order_a = {1, 0, 2},
+    .block_transfer_access_order_b = {1, 0, 2},
+    .src_access_order_a            = {1, 0, 2},
+    .src_access_order_b            = {1, 0, 2}};
+
+constexpr BlockTransferABC FwdBlockTransfer_4x32x1{
+    .block_transfer_a              = {.k0 = 4, .m_n = 32, .k1 = 1},
+    .block_transfer_b              = {.k0 = 4, .m_n = 32, .k1 = 1},
+    .thread_cluster_dims_c         = {.m_block        = 1,
+                                      .m_wave_per_xdl = 32,
+                                      .n_block        = 1,
+                                      .n_wave_per_xdl = 4},
+    .lds_transfer_a                = {.src_vector_dim            = 2,
+                                      .src_scalar_per_vector     = 16,
+                                      .lds_dst_scalar_per_vector = 16,
+                                      .is_direct_load            = false,
+                                      .lds_padding               = true},
+    .lds_transfer_b                = {.src_vector_dim            = 2,
+                                      .src_scalar_per_vector     = 16,
+                                      .lds_dst_scalar_per_vector = 16,
+                                      .is_direct_load            = false,
+                                      .lds_padding               = true},
+    .epilogue_c                    = {.m_per_wave_per_shuffle = 1,
+                                      .n_per_wave_per_shuffle = 1,
+                                      .scalar_per_vector      = 8},
+    .block_transfer_access_order_a = {1, 0, 2},
+    .block_transfer_access_order_b = {1, 0, 2},
+    .src_access_order_a            = {1, 0, 2},
+    .src_access_order_b            = {1, 0, 2}};
+
+constexpr GridwiseXdlGemm FwdGemmParams_Xdl_4x4_per_wave{
+    .ak1 = 8, .bk1 = 8, .m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4};
+
+constexpr GridwiseXdlGemm FwdGemmParams_Xdl_2x1_per_wave{
+    .ak1 = 8, .bk1 = 8, .m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 2, .n_xdl_per_wave = 1};
+
+constexpr GridwiseWmmaGemm FwdGemmParams_Wmma_2x1_per_wave{.k1              = 8,
+                                                           .m_per_wmma      = 32,
+                                                           .n_per_wmma      = 32,
+                                                           .m_wmma_per_wave = 2,
+                                                           .n_wmma_per_wave = 1,
+                                                           .pipeline_version =
+                                                               GridwiseGemmPipelineVersion::V1};
+
+constexpr ThreadBlock FwdThreadBlock_256x256x32{.block_size = 256,
+                                                .tile_size  = {.m = 256, .n = 256, .k = 32}};
+
+constexpr ThreadBlock FwdThreadBlock_128x128x32{.block_size = 256,
+                                                .tile_size  = {.m = 128, .n = 128, .k = 32}};
+
+constexpr ThreadBlock FwdThreadBlock_64x32x32{.block_size = 64,
+                                              .tile_size  = {.m = 64, .n = 32, .k = 32}};
+
+constexpr ThreadBlock FwdThreadBlock_64x64x64{.block_size = 128,
+                                              .tile_size  = {.m = 64, .n = 64, .k = 64}};
+
+constexpr BlockGemm BlockGemmDesc_v1_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V1,
+                                                  .scheduler =
+                                                      BlockGemmPipelineScheduler::INTRAWAVE};
+
+constexpr BlockGemm BlockGemmDesc_v2_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V2,
+                                                  .scheduler =
+                                                      BlockGemmPipelineScheduler::INTRAWAVE};
+
+constexpr BlockGemm BlockGemmDesc_v3_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V3,
+                                                  .scheduler =
+                                                      BlockGemmPipelineScheduler::INTRAWAVE};
+
+constexpr BlockGemm BlockGemmDesc_v4_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V4,
+                                                  .scheduler =
+                                                      BlockGemmPipelineScheduler::INTRAWAVE};
+
+constexpr BlockGemm BlockGemmDesc_v5_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V5,
+                                                  .scheduler =
+                                                      BlockGemmPipelineScheduler::INTRAWAVE};
+
+} // namespace ck_tile::builder::test_utils
diff --git a/experimental/builder/test/utils/ckb_conv_test_utils.hpp b/experimental/builder/test/utils/ckb_conv_test_utils.hpp
new file mode 100644
index 0000000000..6f9442d067
--- /dev/null
+++ b/experimental/builder/test/utils/ckb_conv_test_utils.hpp
@@ -0,0 +1,31 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <gtest/gtest.h>
+#include <gmock/gmock.h>
+
+namespace ck_tile::builder::test_utils {
+using namespace test;
+
+// Common test implementation
+template <typename Builder>
+constexpr void run_test(const std::vector<std::string>& kernel_instance_components)
+{
+    auto instance = typename Builder::Instance{};
+
+    const auto kernel_string = instance.GetTypeString();
+    std::cout << "Generated kernel: " << kernel_string << std::endl;
+    EXPECT_GT(kernel_string.size(), 0);
+
+    const auto invoker_ptr = instance.MakeInvokerPointer();
+    EXPECT_NE(invoker_ptr, nullptr);
+
+    for(const auto& component : kernel_instance_components)
+    {
+        EXPECT_THAT(kernel_string, ::testing::HasSubstr(component));
+    }
+}
+
+} // namespace ck_tile::builder::test_utils
diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
index c7ee3e9ecf..650c6f11d3 100644
--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
+++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
@@ -1,5 +1,5 @@
+// Copyright (C) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -24,6 +24,10 @@
 #include "ck/host_utility/device_prop.hpp"
 #include "ck/host_utility/kernel_launch.hpp"
 
+#ifdef CK_EXPERIMENTAL_BUILDER
+#include "ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp"
+#endif
+
 namespace ck {
 namespace tensor_operation {
 namespace device {
@@ -1225,6 +1229,19 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle
         return str.str();
     }
 
+#ifdef CK_EXPERIMENTAL_BUILDER
+    std::string GetInstanceString() const override
+    {
+        static_assert(ck_tile::reflect::HasInstanceTraits<DeviceOp>,
+                      "Specialization of instance_traits not found. Please check that a "
+                      "specialization exists in file "
+                      "ck_tile/builder/reflect/"
+                      "instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp "
+                      "for the given template parameters.");
+        return ck_tile::reflect::instance_string<DeviceOp>();
+    }
+#endif
+
     size_t GetWorkSpaceSize(const BaseArgument* p_arg) const override
     {
         auto arg = dynamic_cast<const Argument*>(p_arg);
diff --git a/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp
index 6f24bf3df8..1bb0b63792 100644
--- a/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp
+++ b/include/ck/tensor_operation/gpu/element/element_wise_operation.hpp
@@ -73,7 +73,7 @@ struct AddReluAdd
     __host__ __device__ constexpr void operator()<half_t, float, half_t, half_t>(
         half_t& y, const float& x0, const half_t& x1, const half_t& x2) const
     {
-        float y_float;
+        float y_float = 0.0;
         (*this)(y_float, x0, x1, x2);
         y = y_float;
     }
diff --git a/include/ck_tile/host/kernel_launch.hpp b/include/ck_tile/host/kernel_launch.hpp
index 9ac0b5ba0e..be38e92b1a 100644
--- a/include/ck_tile/host/kernel_launch.hpp
+++ b/include/ck_tile/host/kernel_launch.hpp
@@ -110,6 +110,10 @@ CK_TILE_HOST double timing_loop_impl(TimerType timer,
 {
     for(int i = 0; i < s.cold_niters_; i++)
     {
+        if constexpr(!std::is_same_v<PreprocessFunc, std::nullptr_t>)
+        {
+            preprocess();
+        }
         callables_func();
     }
     // Only profile preprocess if it's provided
diff --git a/include/ck_tile/ops/gemm/kernel/streamk_gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/streamk_gemm_kernel.hpp
index 915aebd1e6..d8850749f1 100644
--- a/include/ck_tile/ops/gemm/kernel/streamk_gemm_kernel.hpp
+++ b/include/ck_tile/ops/gemm/kernel/streamk_gemm_kernel.hpp
@@ -84,9 +84,10 @@ struct StreamKKernel
     using CLayout = typename GemmPipeline::CLayout;
 
     /// @brief  Specify the data type configurations for A, B, and C
-    using ADataType = typename GemmPipeline::ADataType;
-    using BDataType = typename GemmPipeline::BDataType;
-    using CDataType = typename EpiloguePipeline::ODataType;
+    using ADataType   = typename GemmPipeline::ADataType;
+    using BDataType   = typename GemmPipeline::BDataType;
+    using CDataType   = typename EpiloguePipeline::ODataType;
+    using AccDataType = typename EpiloguePipeline::AccDataType;
 
     template <typename T>
     static constexpr bool is_tuple_v = is_detected<is_tuple, T>::value;
@@ -243,14 +244,6 @@ struct StreamKKernel
 
     CK_TILE_HOST static bool IsSupportedArgument(const StreamKKernelArgs& kargs)
     {
-        if(kargs.reduction_strategy == StreamKReductionStrategy::Reduction)
-        {
-            if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
-            {
-                CK_TILE_ERROR("CK Tile Stream-K only supports the atomic reduction strategy.");
-            }
-            return false;
-        }
         return UniversalGemmKernel::IsSupportedArgument(kargs);
     }
 
@@ -258,7 +251,7 @@ struct StreamKKernel
     /// @return The buffer size needed.
     CK_TILE_HOST static uint32_t GetWorkSpaceSize(const StreamKKernelArgs& kargs)
     {
-        return kargs.tile_partitioner.GetWorkSpaceSize(sizeof(CDataType));
+        return kargs.tile_partitioner.get_workspace_size(sizeof(AccDataType));
     }
 
     /// @brief Sets the kargs' current workspace_ptr to the given workspace_ptr.
@@ -299,6 +292,118 @@ struct StreamKKernel
             {a_ptr}, {b_ptr}, {/*ds_ptr*/}, c_ptr, smem_ptr_0, kargs, num_loop, i_m, i_n, k_size);
     }
 
+    /// @brief Signals that the current thread block (CTA) has completed storing its partial
+    /// results.
+    /// @param kargs Kernel arguments, including the workspace pointer.
+    /// @param cta_idx The index of the current thread block (CTA).
+    /// @note This function utilizes a workgroup barrier to set a synchronization flag for the given
+    /// CTA index.
+    CK_TILE_DEVICE void SignalStorePartialDone(const StreamKKernelArgs& kargs,
+                                               index_t cta_idx) const
+    {
+        auto sk_flags_ptr = static_cast<uint32_t*>(kargs.workspace_ptr);
+        workgroup_barrier sk_flags(sk_flags_ptr);
+        sk_flags.wait_set(0, 1, cta_idx);
+    }
+
+    /// @brief Waits for the thread block (cta_idx) to complete storing its partial results.
+    /// @param kargs Kernel arguments, including the workspace pointer.
+    /// @param cta_idx The index of the thread block (CTA).
+    /// @note This function utilizes a workgroup barrier to wait for the synchronization flag to be
+    /// set by the given CTA index.
+    CK_TILE_DEVICE void WaitStorePartialDone(const StreamKKernelArgs& kargs, index_t cta_idx) const
+    {
+        auto sk_flags_ptr = static_cast<uint32_t*>(kargs.workspace_ptr);
+        workgroup_barrier sk_flags(sk_flags_ptr);
+        sk_flags.wait_eq(1, cta_idx);
+    }
+
+    /// @brief Adds the values of a block tile to an output block tile.
+    /// @param in_out_block_tile The output block tile to which values are added.
+    /// @param in_block_tile The input block tile whose values are added.
+    /// @note This function iterates over the distributed spans of the block tiles and updates the
+    /// output block tile with accumulated values.
+    template <typename OAccTile>
+    CK_TILE_DEVICE void AddBlockTile(OAccTile& in_out_block_tile,
+                                     const OAccTile& in_block_tile) const
+    {
+        using BlockType        = remove_cvref_t<decltype(in_out_block_tile)>;
+        constexpr auto o_spans = BlockType::get_distributed_spans();
+        sweep_tile_span(o_spans[number<0>{}], [&](auto idx0) {
+            sweep_tile_span(o_spans[number<1>{}], [&](auto idx1) {
+                constexpr auto idx     = make_tuple(idx0, idx1);
+                in_out_block_tile(idx) = in_out_block_tile[idx] + in_block_tile[idx];
+            });
+        });
+    }
+
+    /// @brief Loads a partial block tile from the workspace buffer.
+    /// @param kargs Kernel arguments, including the workspace pointer.
+    /// @param cta_idx The index of the thread block (CTA).
+    /// @param c_block_tile_dist The tile distribution for the block.
+    /// @return The loaded partial block tile.
+    /// @note This function calculates the buffer pointer and uses the tile distribution for loading
+    /// the partial block tile.
+    template <typename DataType, typename OAccTileDist>
+    CK_TILE_DEVICE auto LoadPartial(const StreamKKernelArgs& kargs,
+                                    index_t cta_idx,
+                                    const OAccTileDist& c_block_tile_dist) const
+    {
+        const auto c_block_tile_buffer_size =
+            TilePartitioner::MPerBlock * TilePartitioner::NPerBlock * sizeof(DataType);
+        void* partial_buffer_ptr = static_cast<char*>(kargs.workspace_ptr) +
+                                   kargs.tile_partitioner.get_flags_buffer_size() +
+                                   cta_idx * c_block_tile_buffer_size;
+
+        const auto& partial_tensor_view = make_naive_tensor_view<address_space_enum::global>(
+            static_cast<DataType*>(partial_buffer_ptr),
+            make_tuple(number<TilePartitioner::MPerBlock>{}, number<TilePartitioner::NPerBlock>{}),
+            make_tuple(TilePartitioner::NPerBlock, 1),
+            number<GemmPipeline::GetVectorSizeC()>{},
+            number<1>{});
+
+        auto partial_tile_window = make_tile_window(
+            partial_tensor_view,
+            make_tuple(number<TilePartitioner::MPerBlock>{}, number<TilePartitioner::NPerBlock>{}),
+            {0, 0},
+            c_block_tile_dist);
+
+        return load_tile(partial_tile_window);
+    }
+
+    /// @brief Stores a partial block tile to the workspace buffer.
+    /// @param kargs Kernel arguments, including the workspace pointer.
+    /// @param cta_idx The index of the thread block (CTA).
+    /// @param c_block_tile The block tile to be stored.
+    /// @note This function calculates the buffer pointer and uses the tile window for storing the
+    /// partial block tile.
+    template <typename OAccTile>
+    CK_TILE_DEVICE void StorePartial(const StreamKKernelArgs& kargs,
+                                     index_t cta_idx,
+                                     const OAccTile& c_block_tile) const
+    {
+        const auto c_block_tile_buffer_size = TilePartitioner::MPerBlock *
+                                              TilePartitioner::NPerBlock *
+                                              sizeof(typename OAccTile::DataType);
+        void* partial_buffer_ptr = static_cast<char*>(kargs.workspace_ptr) +
+                                   kargs.tile_partitioner.get_flags_buffer_size() +
+                                   cta_idx * c_block_tile_buffer_size;
+
+        const auto& partial_tensor_view = make_naive_tensor_view<address_space_enum::global>(
+            static_cast<typename OAccTile::DataType*>(partial_buffer_ptr),
+            make_tuple(number<TilePartitioner::MPerBlock>{}, number<TilePartitioner::NPerBlock>{}),
+            make_tuple(TilePartitioner::NPerBlock, 1),
+            number<GemmPipeline::GetVectorSizeC()>{},
+            number<1>{});
+
+        auto partial_tile_window = make_tile_window(
+            partial_tensor_view,
+            make_tuple(number<TilePartitioner::MPerBlock>{}, number<TilePartitioner::NPerBlock>{}),
+            {0, 0});
+
+        store_tile(partial_tile_window, c_block_tile);
+    }
+
     /// @brief Runs the main Stream-K algorithm.
     /// @param kargs Stream-K kernel arguments.
     /// @param cta_idx The current Stream-K workgroup's index.
@@ -347,7 +452,88 @@ struct StreamKKernel
             }
             else
             {
-                // TODO: Apply reduction logic.
+                const auto c_macro_tile_idx =
+                    kargs.tile_partitioner.get_output_tile_index(tile_idx);
+                index_t i_m =
+                    c_macro_tile_idx[UniversalGemmKernel::I0] * TilePartitioner::MPerBlock;
+                index_t i_n =
+                    c_macro_tile_idx[UniversalGemmKernel::I1] * TilePartitioner::NPerBlock;
+
+                const ADataType* a_ptr = static_cast<const ADataType*>(kargs.as_ptr[0]) + i_k_a;
+                const BDataType* b_ptr = static_cast<const BDataType*>(kargs.bs_ptr[0]) + i_k_b;
+                CDataType* c_ptr       = static_cast<CDataType*>(kargs.e_ptr);
+
+                // Create Gemm tensor views, pad views and tile windows
+                const auto& gemm_tensor_views_tuple =
+                    UniversalGemmKernel::template MakeGemmTensorViews<
+                        EpiloguePipeline::MemoryOperation>(
+                        {a_ptr}, {b_ptr}, {/*ds_ptr*/}, c_ptr, kargs, k_size);
+
+                const auto& gemm_pad_views =
+                    UniversalGemmKernel::MakeGemmPadViews(gemm_tensor_views_tuple);
+                auto gemm_tile_windows =
+                    UniversalGemmKernel::MakeGemmTileWindows(gemm_pad_views, i_m, i_n);
+
+                // Run GEMM cooperatively by whole workgroup.
+                const auto& as_block_window = gemm_tile_windows.at(UniversalGemmKernel::I0);
+                const auto& bs_block_window = gemm_tile_windows.at(UniversalGemmKernel::I1);
+                const auto& ds_block_window = gemm_tile_windows.at(UniversalGemmKernel::I2);
+
+                // Since num_loop can vary per WG and per iteration of the Stream-K while loop,
+                // we compute has_hot_loop and tail_num here. This is a similar pattern used by
+                // grouped GEMM. In this case, we call the GemmPipeline's operator() function
+                // that takes both has_hot_loop and tail_num.
+                const bool has_hot_loop   = GemmPipeline::BlockHasHotloop(num_loop_sk);
+                const TailNumber tail_num = GemmPipeline::GetBlockLoopTailNum(num_loop_sk);
+
+                const auto& c_block_tile = GemmPipeline{}(as_block_window[UniversalGemmKernel::I0],
+                                                          bs_block_window[UniversalGemmKernel::I0],
+                                                          num_loop_sk,
+                                                          has_hot_loop,
+                                                          tail_num,
+                                                          smem_ptr_0);
+
+                auto tile_started = iter_start == tile_iter_start;
+                auto tile_ended   = iter_end >= tile_iter_end;
+                if(!tile_started)
+                {
+                    StorePartial(kargs, cta_idx, c_block_tile);
+                    // Ensure device-wide visibility of partial results stored in global memory
+                    // before signaling completion. __threadfence() guarantees that all global
+                    // memory writes by this thread are visible to other threads on the device.
+                    __threadfence(); // send signal when the store is done
+                    SignalStorePartialDone(kargs, cta_idx);
+                }
+                else
+                {
+                    auto accum_block_tile = c_block_tile;
+                    if(!tile_ended)
+                    {
+                        const index_t iter_per_tile = kargs.tile_partitioner.get_iters_per_tile();
+                        const index_t iter_per_cta  = kargs.tile_partitioner.get_iters_per_sk_cta();
+                        const index_t extra_iters   = kargs.tile_partitioner.get_extra_iters();
+                        int accum_iters             = local_iter_end - local_iter_start;
+                        int next_cta                = cta_idx + 1;
+
+                        while(accum_iters < iter_per_tile)
+                        {
+                            WaitStorePartialDone(kargs, next_cta);
+
+                            using BlockType = remove_cvref_t<decltype(c_block_tile)>;
+                            AddBlockTile(
+                                accum_block_tile,
+                                LoadPartial<typename BlockType::DataType>(
+                                    kargs, next_cta, c_block_tile.get_tile_distribution()));
+
+                            accum_iters += iter_per_cta + (next_cta < extra_iters);
+                            ++next_cta;
+                        }
+                    }
+
+                    auto& c_block_window = gemm_tile_windows.at(UniversalGemmKernel::I3);
+                    EpiloguePipeline{}(
+                        c_block_window, accum_block_tile, ds_block_window, smem_ptr_0);
+                }
             }
 
             // Prepare for next Stream-K loop iteration.
diff --git a/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner.hpp b/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner.hpp
index e98c60e5f0..996ef5a7ef 100644
--- a/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner.hpp
+++ b/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner.hpp
@@ -31,21 +31,20 @@ struct StreamKTilePartitionerBase
 
     StreamKTilePartitionerBase(index_t m, index_t n, index_t k, index_t grid);
 
-    private:
     /**
      * @brief Calculates the total space needed for the partials buffer.
      *
      * @param acc_element_bytes  The number of bytes for the accumulator data type used in the GEMM.
      * @return index_t           The number of bytes needed for the partials buffer.
      */
-    CK_TILE_HOST index_t get_partials_buffer_size(index_t acc_element_bytes) const noexcept;
+    CK_TILE_HOST_DEVICE index_t get_partials_buffer_size(index_t acc_element_bytes) const noexcept;
 
     /**
      * @brief Calculates the total space needed for the flags buffer.
      *
      * @return index_t The number of bytes needed for the flags buffer.
      */
-    CK_TILE_HOST index_t get_flags_buffer_size() const noexcept;
+    CK_TILE_HOST_DEVICE index_t get_flags_buffer_size() const noexcept;
 
     public:
     /**
@@ -123,7 +122,7 @@ struct StreamKTilePartitionerBase
      * @param acc_element_bytes  The number of bytes for the accumulator data type used in the GEMM.
      * @return index_t           The number of bytes needed for the partials and flags buffers.
      */
-    CK_TILE_HOST index_t get_workspace_size(index_t acc_element_bytes) const noexcept;
+    CK_TILE_HOST_DEVICE index_t get_workspace_size(index_t acc_element_bytes) const noexcept;
 
     /**
      * @brief Returns the number of macro tiles in the C tensor.
diff --git a/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner_impl.hpp b/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner_impl.hpp
index 2eefe3420c..b3217624d1 100644
--- a/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner_impl.hpp
+++ b/include/ck_tile/ops/gemm/kernel/streamk_gemm_tile_partitioner_impl.hpp
@@ -45,7 +45,7 @@ StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::StreamKTi
 }
 
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
-CK_TILE_HOST index_t
+CK_TILE_HOST_DEVICE index_t
 StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_partials_buffer_size(
     index_t acc_element_bytes) const noexcept
 {
@@ -53,7 +53,7 @@ StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_parti
 }
 
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
-CK_TILE_HOST index_t
+CK_TILE_HOST_DEVICE index_t
 StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_flags_buffer_size()
     const noexcept
 {
@@ -116,7 +116,7 @@ StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_outpu
 }
 
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
-CK_TILE_HOST index_t
+CK_TILE_HOST_DEVICE index_t
 StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_workspace_size(
     index_t acc_element_bytes) const noexcept
 {
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,
diff --git a/script/launch_tests.sh b/script/launch_tests.sh
index 17a99e62a3..52151b71f6 100755
--- a/script/launch_tests.sh
+++ b/script/launch_tests.sh
@@ -40,20 +40,88 @@ python3 "$SCRIPT_DIR/dependency-parser/main.py" select "$JSON_FILE" origin/devel
 # Path to tests_to_run.json in the same directory
 TEST_FILE="tests_to_run.json"
 
-command=$(python3 -c "
+# Configuration: Adjust these defaults as needed
+# Number of tests per ctest command (can be overridden with CTEST_CHUNK_SIZE env var)
+DEFAULT_CHUNK_SIZE=10
+# Whether to stop on first failure (can be overridden with CTEST_FAIL_FAST env var)
+DEFAULT_FAIL_FAST=false
+
+# Split tests into chunks and run multiple ctest commands
+# Export variables so Python subprocess can access them
+export CHUNK_SIZE=${CTEST_CHUNK_SIZE:-$DEFAULT_CHUNK_SIZE}
+export FAIL_FAST=${CTEST_FAIL_FAST:-$DEFAULT_FAIL_FAST}
+
+python3 -c "
 import json
 import os
+import sys
+import subprocess
+
+CHUNK_SIZE = int(os.environ.get('CHUNK_SIZE', '10'))
+FAIL_FAST = os.environ.get('FAIL_FAST', 'false').lower() == 'true'
+
 with open('$TEST_FILE', 'r') as f:
     data = json.load(f)
     tests = data.get('tests_to_run', [])
-    if tests:
-        # Extract just the filename after the last '/'
-        clean_tests = [os.path.basename(test) for test in tests]
-        print('ctest --output-on-failure -R \"' + '|'.join(clean_tests) + '\"')
-    else:
+    
+    if not tests:
         print('# No tests to run')
-")
+        sys.exit(0)
+    
+    # Extract just the filename after the last '/'
+    clean_tests = [os.path.basename(test) for test in tests]
+    
+    total_tests = len(clean_tests)
+    total_chunks = (total_tests + CHUNK_SIZE - 1) // CHUNK_SIZE
+    
+    print(f'# Total tests to run: {total_tests}')
+    print(f'# Running in {total_chunks} chunk(s) of up to {CHUNK_SIZE} tests each')
+    print(f'# Fail-fast mode: {FAIL_FAST}')
+    print()
+    
+    failed_chunks = []
+    
+    # Split into chunks
+    for i in range(0, total_tests, CHUNK_SIZE):
+        chunk = clean_tests[i:i+CHUNK_SIZE]
+        chunk_num = (i // CHUNK_SIZE) + 1
+        
+        print(f'Running test chunk {chunk_num}/{total_chunks} ({len(chunk)} tests)...')
+        sys.stdout.flush()
+        
+        # Run ctest command, don't raise exception on failure
+        cmd = ['ctest', '--output-on-failure', '-R', '|'.join(chunk)]
+        try:
+            result = subprocess.run(cmd, cwd='$BUILD_DIR', check=False)
+            
+            if result.returncode != 0:
+                failed_chunks.append(chunk_num)
+                print(f'WARNING: Chunk {chunk_num} had test failures (exit code: {result.returncode})')
+                
+                # If fail-fast is enabled, exit immediately
+                if FAIL_FAST:
+                    print(f'FAIL-FAST: Stopping at chunk {chunk_num} due to failures')
+                    sys.exit(1)
+        except Exception as e:
+            print(f'ERROR: Failed to run chunk {chunk_num}: {e}')
+            failed_chunks.append(chunk_num)
+            if FAIL_FAST:
+                sys.exit(1)
+        
+        print()
+        sys.stdout.flush()
+    
+    # Print summary
+    print('=' * 60)
+    if failed_chunks:
+        print(f'SUMMARY: {len(failed_chunks)} of {total_chunks} chunk(s) had failures: {failed_chunks}')
+        print('=' * 60)
+        sys.exit(1)
+    else:
+        print(f'SUMMARY: All {total_chunks} chunk(s) passed successfully!')
+        print('=' * 60)
+        sys.exit(0)
+" 
+PYTHON_EXIT=$?
 
-echo "$command"
-
-eval "$command"
+exit $PYTHON_EXIT
diff --git a/script/sccache_wrapper.sh b/script/sccache_wrapper.sh
index b0ec08de45..30fd17e520 100755
--- a/script/sccache_wrapper.sh
+++ b/script/sccache_wrapper.sh
@@ -52,5 +52,22 @@ if [ "${ENFORCE_REDIS}" == "true" ]; then
 fi
 setup_rocm_compilers_hash_file
 $SCCACHE_BIN --version
+echo "=== Starting sccache server at $(date) ==="
 $SCCACHE_BIN --start-server
 
+# Log initial sccache statistics
+echo "=== Initial sccache statistics ==="
+$SCCACHE_BIN --show-stats || echo "Could not get initial stats"
+
+# Test Redis connectivity and performance
+echo "=== Testing Redis connectivity ==="
+start_time=$(date +%s%N)
+redis-cli -u ${SCCACHE_REDIS} ping || echo "Redis ping failed"
+end_time=$(date +%s%N)
+latency=$(( (end_time - start_time) / 1000000 ))
+echo "Redis ping latency: ${latency}ms"
+
+# Check Redis memory status
+echo "=== Redis memory status ==="
+redis-cli -u ${SCCACHE_REDIS} info memory | grep -E "(used_memory|maxmemory|evicted_keys)" || echo "Could not get Redis memory info"
+