diff --git a/include/ck_tile/ops/common/utils.hpp b/include/ck_tile/ops/common/utils.hpp index b422a0a896..f60a7e1441 100644 --- a/include/ck_tile/ops/common/utils.hpp +++ b/include/ck_tile/ops/common/utils.hpp @@ -19,6 +19,12 @@ template <> struct typeToStr { static constexpr const char * name = "fp8" template <> struct typeToStr { static constexpr const char * name = "bf8"; }; template <> struct typeToStr { static constexpr const char * name = "int8"; }; template <> struct typeToStr { static constexpr const char * name = "pk_int4"; }; + +template struct memOpToStr; +template <> struct memOpToStr { static constexpr const char * name = "set"; }; +template <> struct memOpToStr { static constexpr const char * name = "atomic_add"; }; +template <> struct memOpToStr { static constexpr const char * name = "atomic_max"; }; +template <> struct memOpToStr { static constexpr const char * name = "add"; }; // clang-format on template @@ -32,4 +38,10 @@ std::string gemm_prec_str() return base_str; } +template +std::string mem_op_string() +{ + return std::string(memOpToStr::name); +} + } // namespace ck_tile diff --git a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp index 50ac1328e1..8a84f7e9bf 100644 --- a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp +++ b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp @@ -3,7 +3,9 @@ #pragma once +#include "ck_tile/host/concat.hpp" #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" @@ -123,6 +125,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()); + // clang-format on + } + /** * @brief Get the vector store size for C tensor. * diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp index 1cff9b5733..7b8cdb3792 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp @@ -556,7 +556,12 @@ struct GroupedConvolutionBackwardDataKernel [[nodiscard]] CK_TILE_HOST static const std::string GetName() { // clang-format off - return concat('_', "grouped_convolution_backward_data", gemm_prec_str, GemmPipeline::GetName()); + return concat('_', "grouped_convolution_backward_data", + gemm_prec_str(), + "gemm", + GemmPipeline::GetName(), + "epilogue", + EpiloguePipeline::GetName()); // clang-format on } diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp index b4e0485702..2eb4f2dfd1 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp @@ -412,10 +412,21 @@ struct GroupedConvolutionBackwardWeightKernel { constexpr auto NumGroupsToMerge = GroupedConvTraitsType_::NumGroupsToMerge; // clang-format off - if (NumGroupsToMerge > 1) - return concat('_', "grouped_convolution_backward_weight", gemm_prec_str, GemmPipeline::GetName(), "merge", NumGroupsToMerge); - else - return concat('_', "grouped_convolution_backward_weight", gemm_prec_str, GemmPipeline::GetName()); + if (NumGroupsToMerge > 1) { + return concat('_', "grouped_convolution_backward_weight", + gemm_prec_str(), + "gemm", + GemmPipeline::GetName(), + "epilogue", + EpiloguePipeline::GetName()); + } else { + return concat('_', "grouped_convolution_backward_weight", + gemm_prec_str(), + "gemm", + GemmPipeline::GetName(), + "epilogue", + EpiloguePipeline::GetName(), "merge", NumGroupsToMerge); + } // clang-format on } 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 ce81fe24ed..110ec2cb54 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 @@ -452,7 +452,12 @@ struct GroupedConvolutionForwardKernel [[nodiscard]] CK_TILE_HOST static const std::string GetName() { // clang-format off - return concat('_', "grouped_convolution_forward", gemm_prec_str, GemmPipeline::GetName()); + return concat('_', "grouped_convolution_forward", + gemm_prec_str(), + "gemm", + GemmPipeline::GetName(), + "epilogue", + EpiloguePipeline::GetName()); // clang-format on }