From 4bb53408cb8527bd422e9e2e6aeec47eb6de75c3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Fri, 5 Dec 2025 14:55:29 +0000 Subject: [PATCH] Improve testing transfer parameters. --- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_1d_i8.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 8 +- ...est_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp | 5 +- .../conv/test_ckb_conv_fwd_2d_dl_fp16.cpp | 9 +- .../test/conv/test_ckb_conv_fwd_2d_fp16.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_2d_fp32.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_2d_fp8.cpp | 5 +- ...test_ckb_conv_fwd_2d_large_tensor_fp16.cpp | 9 +- .../test/conv/test_ckb_conv_fwd_3d_bf16.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 5 +- .../test/conv/test_ckb_conv_fwd_3d_fp32.cpp | 5 +- .../test/utils/ckb_conv_test_utils.hpp | 15 + .../test/utils/conv_algorithm_type_utils.hpp | 381 ++++++++++++++++++ 15 files changed, 457 insertions(+), 15 deletions(-) create mode 100644 experimental/builder/test/utils/conv_algorithm_type_utils.hpp 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 937d17a1ff..ac1c779232 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -33,8 +34,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v2_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,256,256,32", + expected_transfer_parameters, "NGCW,GKXC,EmptyTuple,NGKW", "PassThrough,PassThrough,Scale", "Filter1x1Stride1Pad0", 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 e8cd8fb136..1410e63245 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -29,11 +30,13 @@ TEST(FwdConvInstances, .with_prefetch_config(1, 2, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle", + expected_transfer_parameters, "NWGC,GKXC,EmptyTuple,NWGK", "PassThrough,PassThrough,PassThrough", "MNKPadding", - "64,64,32,32", "Default"}); } 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 014e221101..ade245513a 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -31,8 +32,10 @@ TEST(FwdConvInstances, .with_prefetch_config(1, 0, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleD_Wmma_CShuffle", - "128,64,64,64", + expected_transfer_parameters, "GNWC,GKXC,EmptyTuple,GNWK", "PassThrough,PassThrough,PassThrough", "Default"}); 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 b98e28c45a..8406ba80c1 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -29,8 +30,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v1_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,256,256,32", + expected_transfer_parameters, "Default", "NHWGC,GKYXC,EmptyTuple,NHWGK", "PassThrough,PassThrough,PassThrough", @@ -60,7 +63,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v5_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + expected_transfer_parameters, "Filter3x3", "NHWGC,GKYXC,EmptyTuple,NHWGK", "PassThrough,PassThrough,PassThrough", diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp index bc4a5e1047..32ebbfdc82 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -35,10 +36,12 @@ TEST(FwdConvInstances, .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle", + expected_transfer_parameters, "NHWGC,GKYXC,Tuple(NHWGK,G_K),NHWGK", "PassThrough,PassThrough,ScaleAddScaleAddRelu", - "64,64,32,32", "MNKPadding", "Default"}); } diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_dl_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_dl_fp16.cpp index 7af1448403..c24cf47518 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_dl_fp16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_dl_fp16.cpp @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -27,8 +28,10 @@ TEST(FwdConvInstances, Create_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK_Ins .with_dl_transfer(DlFwdTransfer); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK", - "256,128,128,16", + expected_transfer_parameters, "Default", "MNKPadding", "GNHWC,GKYXC,EmptyTuple,GNHWK", @@ -56,8 +59,10 @@ TEST(FwdConvInstances, .with_dl_transfer(DlFwdTransfer); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK", - "256,128,128,16", + expected_transfer_parameters, "Filter1x1Pad0", "MNKPadding", "GNHWC,GKYXC,EmptyTuple,GNHWK", 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 7b522403d3..0d51f04221 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -29,8 +30,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v3_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,256,256,32", + expected_transfer_parameters, "Filter1x1Pad0", "Intrawave", "v3", 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 615d098c7c..df27437fed 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -29,8 +30,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v4_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,128,128,32", + expected_transfer_parameters, "Filter1x1Stride1Pad0", "Intrawave", "v4", diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp8.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp8.cpp index 4dd9e2beef..63e3172a4e 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp8.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp8.cpp @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -29,8 +30,10 @@ TEST(FwdConvInstances, .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle", - "256,256,128,32", + expected_transfer_parameters, "Default", "NHWGC,GKYXC,EmptyTuple,NHWGK", "PassThrough,PassThrough,PassThrough", diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp index 8fe58dbe82..26059d4c5f 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -30,8 +31,10 @@ TEST(FwdConvInstances, .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT)}; using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor", - "256,256,128,32", + expected_transfer_parameters, "Default", "GNHWC,GKYXC,EmptyTuple,GNHWK", "PassThrough,PassThrough,PassThrough", @@ -61,8 +64,10 @@ TEST( .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT)}; using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor", - "128,128,128,32", + expected_transfer_parameters, "Filter1x1Pad0", "GNHWC,GKYXC,EmptyTuple,GNHWK", "PassThrough,PassThrough,PassThrough", 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 2df76ab3e0..7e15ccc8b8 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -30,8 +31,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v3_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,256,256,32", + expected_transfer_parameters, "Default", "Intrawave", "v3", 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 ad626d9a15..6057a9ac53 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -31,8 +32,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v4_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,128,128,32", + expected_transfer_parameters, "Filter1x1Pad0", "Intrawave", "v4", 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 85974ace5d..9705a7231c 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 @@ -3,6 +3,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" namespace { @@ -31,8 +32,10 @@ TEST(FwdConvInstances, .with_block_gemm(BlockGemmDesc_v1_intrawave); using Builder = ConvBuilder; + + const auto expected_transfer_parameters = to_string(FwdConvAlgorithm); run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256,256,256,32", + expected_transfer_parameters, "Filter1x1Pad0", "Intrawave", "v1", diff --git a/experimental/builder/test/utils/ckb_conv_test_utils.hpp b/experimental/builder/test/utils/ckb_conv_test_utils.hpp index 508c621c2e..ae2e9da769 100644 --- a/experimental/builder/test/utils/ckb_conv_test_utils.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_utils.hpp @@ -28,4 +28,19 @@ constexpr void run_test(const std::vector& kernel_instance_componen } } +template +constexpr void run_test(const std::string& expected_kernel_instance_string) +{ + auto instance = typename Builder::Instance{}; + + const auto kernel_string = instance.GetInstanceString(); + 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); + + EXPECT_THAT(kernel_string, ::testing::HasSubstr(expected_kernel_instance_string)); +} + } // namespace ck_tile::builder::test_utils diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp new file mode 100644 index 0000000000..ab076cc986 --- /dev/null +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -0,0 +1,381 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "../impl/conv_algorithm_types.hpp" +#include +#include + +namespace ck_tile::builder::test { + +namespace ckb = ck_tile::builder; + +// Helper function to convert arrays to Seq(...) format +template +std::string array_to_seq(const std::array& arr) +{ + std::ostringstream oss; + oss << "Seq("; + for(size_t i = 0; i < N; ++i) + { + if(i > 0) oss << ","; + oss << arr[i]; + } + oss << ")"; + return oss.str(); +} + +// Base template - will cause compilation error for unsupported types +template +std::string to_string(T) +{ + static_assert(sizeof(T) == 0, "Unsupported type"); + return ""; +} + +// Template specializations for enum types + +template<> +inline std::string to_string(PipelineVersion t) +{ + std::ostringstream oss; + oss << t; + return oss.str(); +} + +template<> +inline std::string to_string(PipelineScheduler t) +{ + std::ostringstream oss; + oss << t; + return oss.str(); +} + +template<> +inline std::string to_string(ConvFwdSpecialization t) +{ + std::ostringstream oss; + oss << t; + return oss.str(); +} + +template<> +inline std::string to_string(GemmSpecialization t) +{ + std::ostringstream oss; + oss << t; + return oss.str(); +} + +// Template specializations for struct types + +template<> +inline std::string to_string>(MNK t) +{ + return array_to_seq(std::array{t.m, t.n, t.k}); +} + +template<> +inline std::string to_string(ThreadBlock t) +{ + std::ostringstream oss; + oss << t.block_size << "," + << t.tile_size.m << "," + << t.tile_size.n << "," + << t.tile_size.k; + return oss.str(); +} + +template<> +inline std::string to_string(GridwiseXdlGemm t) +{ + std::ostringstream oss; + oss << t.ak1 << "," + << t.bk1 << "," + << t.m_per_xdl << "," + << t.n_per_xdl << "," + << t.m_xdl_per_wave << "," + << t.n_xdl_per_wave; + return oss.str(); +} + +template<> +inline std::string to_string(GridwiseWmmaGemm t) +{ + std::ostringstream oss; + oss << t.k1 << "," + << t.m_per_wmma << "," + << t.n_per_wmma << "," + << t.m_wmma_per_wave << "," + << t.n_wmma_per_wave << "," + << to_string(t.pipeline_version); + return oss.str(); +} + +template<> +inline std::string to_string(BlockGemm t) +{ + std::ostringstream oss; + oss << to_string(t.scheduler) << "," + << to_string(t.pipeline_version); + return oss.str(); +} + +template<> +inline std::string to_string(BlockTransfer t) +{ + return array_to_seq(std::array{t.k0, t.m_n, t.k1}); +} + +template<> +inline std::string to_string(ThreadCluster t) +{ + return array_to_seq(std::array{t.m_block, t.m_wave_per_xdl, t.n_block, t.n_wave_per_xdl}); +} + +template<> +inline std::string to_string(LdsTransfer t) +{ + std::ostringstream oss; + oss << t.src_vector_dim << "," + << t.src_scalar_per_vector << "," + << t.lds_dst_scalar_per_vector << "," + << (t.is_direct_load ? "1" : "0") << "," + << (t.lds_padding ? "1" : "0"); + return oss.str(); +} + +template<> +inline std::string to_string(AccessOrder t) +{ + return array_to_seq(t.order); +} + +template<> +inline std::string to_string(TransferAB t) +{ + std::ostringstream oss; + oss << to_string(t.block_transfer) << "," + << to_string(t.block_transfer_access_order) << "," + << to_string(t.src_access_order) << "," + << t.lds_transfer.src_vector_dim << "," + << t.lds_transfer.src_scalar_per_vector << "," + << t.lds_transfer.lds_dst_scalar_per_vector << "," + << (t.lds_transfer.lds_padding ? "1" : "0"); + return oss.str(); +} + +template<> +inline std::string to_string(TransferC t) +{ + std::ostringstream oss; + oss << t.epilogue.m_xdl_per_wave_per_shuffle << "," + << t.epilogue.n_per_wave_per_shuffle << "," + << to_string(t.thread_cluster_dims) << "," + << t.epilogue.scalar_per_vector; + return oss.str(); +} + +template<> +inline std::string to_string(TransferABC t) +{ + std::ostringstream oss; + oss << to_string(t.a) << "," + << to_string(t.b) << "," + << to_string(t.c); + return oss.str(); +} + +template<> +inline std::string to_string(DlThreadConfig t) +{ + std::ostringstream oss; + oss << t.k0_per_block << "," + << t.k1 << "," + << t.m1_per_thread << "," + << t.n1_per_thread << "," + << t.k_per_thread; + return oss.str(); +} + +template<> +inline std::string to_string(DlThreadCluster t) +{ + std::ostringstream oss; + oss << array_to_seq(t.m1_xs) << "," + << array_to_seq(t.n1_xs); + return oss.str(); +} + +template<> +inline std::string to_string(DlBlockTransfer t) +{ + std::ostringstream oss; + oss << array_to_seq(t.thread_slice_lengths) << "," + << array_to_seq(t.thread_cluster_lengths) << "," + << array_to_seq(t.thread_cluster_arrange_order) << "," + << array_to_seq(t.src_access_order) << "," + << array_to_seq(t.src_vector_tensor_lengths) << "," + << array_to_seq(t.src_vector_tensor_contiguous_dim_order) << "," + << array_to_seq(t.dst_vector_tensor_lengths); + return oss.str(); +} + +template<> +inline std::string to_string(DlEpilogue t) +{ + std::ostringstream oss; + oss << t.src_dst_vector_dim << "," + << t.dst_scalar_per_vector << "," + << array_to_seq(t.src_dst_access_order); + return oss.str(); +} + +template<> +inline std::string to_string(DlBlockTransferAB t) +{ + return to_string(t.block_transfer); +} + +template<> +inline std::string to_string(DlBlockTransferC t) +{ + return to_string(t.epilogue); +} + +template<> +inline std::string to_string(DlTransferABC t) +{ + std::ostringstream oss; + oss << to_string(t.a) << "," + << to_string(t.b) << "," + << to_string(t.c); + return oss.str(); +} + +// Template specializations for factory wrapper types + +template<> +inline std::string to_string(ThreadBlock_ t) +{ + return to_string(t.thread_block); +} + +template<> +inline std::string to_string(XdlGemm_ t) +{ + return to_string(t.gridwise_gemm); +} + +template<> +inline std::string to_string(WmmaGemm_ t) +{ + return to_string(t.gridwise_gemm); +} + +template<> +inline std::string to_string(Transfer_ t) +{ + return to_string(t.transfer); +} + +template<> +inline std::string to_string(ConvSpecialization_ t) +{ + std::ostringstream oss; + oss << to_string(t.fwd_specialization) << "," + << to_string(t.gemm_specialization); + return oss.str(); +} + +template<> +inline std::string to_string(Prefetch_ t) +{ + std::ostringstream oss; + oss << t.num_gemm_k_prefetch_stages << "," + << t.num_groups_to_merge << "," + << to_string(t.loop_scheduler); + return oss.str(); +} + +template<> +inline std::string to_string(BlockGemm_ t) +{ + return to_string(t.block_gemm); +} + +template<> +inline std::string to_string(DlThreadConfig_ t) +{ + return to_string(t.thread_config); +} + +template<> +inline std::string to_string(DlThreadCluster_ t) +{ + return to_string(t.thread_cluster); +} + +template<> +inline std::string to_string(DlTransfer_ t) +{ + return to_string(t.transfer); +} + +// Template specializations for algorithm types + +template<> +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle t) +{ + std::ostringstream oss; + oss << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast(t)); + return oss.str(); +} + +template<> +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 t) +{ + std::ostringstream oss; + oss << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast(t)); + return oss.str(); +} + +template<> +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle t) +{ + std::ostringstream oss; + oss << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast(t)); + return oss.str(); +} + +template<> +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK t) +{ + std::ostringstream oss; + oss << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast(t)); + return oss.str(); +} + +template<> +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor t) +{ + return to_string(t.base_algorithm); +} + +} // namespace ck_tile::builder::test