From fd61987d73ef887b2180c228b5e7f3e50c3bec4c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <188998872+vpietila-amd@users.noreply.github.com> Date: Thu, 30 Oct 2025 14:19:07 +0200 Subject: [PATCH] [CK_TILE] Improve grouped conv kernel name generation (#3028) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Improve the grouped conv kernel name generation in CK Tile. * Fix building CShuffle epilogue tests. --------- Co-authored-by: Bartłomiej Kocot [ROCm/composable_kernel commit: 9ee9f4d2a3d967d99b6fb3d9dce75bfa2084ba18] --- include/ck_tile/ops/common/utils.hpp | 12 ++++++++++++ .../ops/epilogue/cshuffle_epilogue.hpp | 15 +++++++++++++++ ...ouped_convolution_backward_data_kernel.hpp | 7 ++++++- ...ped_convolution_backward_weight_kernel.hpp | 19 +++++++++++++++---- .../grouped_convolution_forward_kernel.hpp | 7 ++++++- 5 files changed, 54 insertions(+), 6 deletions(-) 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 }