Improve the grouped conv kernel name generation in CK Tile.

include/ck_tile/ops/common/utils.hpp

@@ -19,6 +19,12 @@ template <> struct typeToStr<fp8_t> { static constexpr const char * name = "fp8"
 template <> struct typeToStr<bf8_t> { static constexpr const char * name = "bf8"; };
 template <> struct typeToStr<int8_t> { static constexpr const char * name = "int8"; };
 template <> struct typeToStr<pk_int4_t> { static constexpr const char * name = "pk_int4"; };
+
+template <memory_operation_enum MemOp> struct memOpToStr;
+template <> struct memOpToStr<memory_operation_enum::set> { static constexpr const char * name = "set"; };
+template <> struct memOpToStr<memory_operation_enum::atomic_add> { static constexpr const char * name = "atomic_add"; };
+template <> struct memOpToStr<memory_operation_enum::atomic_max> { static constexpr const char * name = "atomic_max"; };
+template <> struct memOpToStr<memory_operation_enum::add> { static constexpr const char * name = "add"; };
 // clang-format on
 
 template <typename ADataType_, typename BDataType_>
@@ -32,4 +38,10 @@ std::string gemm_prec_str()
     return base_str;
 }
 
+template <memory_operation_enum MemOp_>
+std::string mem_op_string()
+{
+    return std::string(memOpToStr<MemOp_>::name);
+}
+
 } // namespace ck_tile

include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp

@@ -4,6 +4,7 @@
 #pragma once
 
 #include "ck_tile/core.hpp"
+#include "ck_tile/ops/common/utils.hpp"
 #include "ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp"
 #include "ck_tile/ops/common/tensor_layout.hpp"
 #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp"
@@ -116,6 +117,19 @@ struct CShuffleEpilogue
 
     static_assert(NumDTensor == DsLayout::size(),
                   "The size of DsDataType and DsLayout should be the same");
+
+    [[nodiscard]] CK_TILE_HOST static const std::string GetName()
+    {
+        // clang-format off
+        return concat('_', "CShuffleEpilogue", 
+                      concat('x', MWave, NWave),
+                      concat('x', MPerXdl, NPerXdl, KPerXdl),
+                      VectorSizeC,
+                      isCTransposed ? "CTransposed" : "CNotTransposed",
+                      mem_op_string<MemoryOperation>());
+        // clang-format on
+    }
+
     /**
      * @brief Get the vector store size for C tensor.
      *

include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp

@@ -529,7 +529,12 @@ struct GroupedConvolutionBackwardDataKernel
     [[nodiscard]] CK_TILE_HOST static const std::string GetName()
     {
         // clang-format off
-        return concat('_', "grouped_convolution_backward_data", gemm_prec_str<InDataType, WeiDataType>, GemmPipeline::GetName());
+        return concat('_', "grouped_convolution_backward_data", 
+            gemm_prec_str<InDataType, WeiDataType>(), 
+            "gemm",
+            GemmPipeline::GetName(),
+            "epilogue",
+            EpiloguePipeline::GetName());
         // clang-format on
     }
 

include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp

@@ -391,7 +391,12 @@ struct GroupedConvolutionBackwardWeightKernel
     [[nodiscard]] CK_TILE_HOST static const std::string GetName()
     {
         // clang-format off
-        return concat('_', "grouped_convolution_backward_weight", gemm_prec_str<InDataType, WeiDataType>, GemmPipeline::GetName());
+        return concat('_', "grouped_convolution_backward_weight", 
+            gemm_prec_str<InDataType, WeiDataType>(),
+            "gemm",
+            GemmPipeline::GetName(),
+            "epilogue",
+            EpiloguePipeline::GetName());
         // clang-format on
     }
 

include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp

@@ -442,7 +442,12 @@ struct GroupedConvolutionForwardKernel
     [[nodiscard]] CK_TILE_HOST static const std::string GetName()
     {
         // clang-format off
-        return concat('_', "grouped_convolution_forward", gemm_prec_str<InDataType, WeiDataType>, GemmPipeline::GetName());
+        return concat('_', "grouped_convolution_forward", 
+            gemm_prec_str<InDataType, WeiDataType>(), 
+            "gemm",
+            GemmPipeline::GetName(),
+            "epilogue",
+            EpiloguePipeline::GetName());
         // clang-format on
     }