From bc3cba873bc523c24f64328dbe5a0d30798cf539 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Fri, 2 Jan 2026 06:08:29 -0500 Subject: [PATCH] Add factory and tests for DeviceGroupedConvBwdWeight_Wmma_CShuffleV3. --- .../builder/factory/conv_algorithms.hpp | 122 +++++++++++++----- .../factory/conv_bwd_weight_dl_factory.hpp | 3 +- .../conv_bwd_weight_multi_d_xdl_factory.hpp | 4 +- .../conv_bwd_weight_two_stage_xdl_factory.hpp | 4 +- .../conv_bwd_weight_wmma_v3_factory.hpp | 104 +++++++++++++++ .../factory/conv_bwd_weight_xdl_factory.hpp | 4 +- .../conv_bwd_weight_xdl_v3_factory.hpp | 4 +- .../builder/factory/conv_dispatcher.hpp | 4 +- experimental/builder/test/CMakeLists.txt | 1 + ...t_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp | 44 +++++++ .../test/utils/conv_algorithm_type_utils.hpp | 19 +++ 11 files changed, 268 insertions(+), 45 deletions(-) create mode 100644 experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp create mode 100644 experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp index 588e5eb698..1347b132e8 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -570,6 +570,12 @@ consteval int count_matches_bwd_two_stage_xdl() { return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9 + Alg::c10 + Alg::c11 + Alg::c12; } +template +consteval int count_matches_bwd_wmma_v3() { + using Alg = BwdWmmaV3Algorithm; + return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9 + Alg::c10 + Alg::c11; +} + template consteval int count_matches_bwd_dl() { using Alg = BwdDlAlgorithm; @@ -599,26 +605,26 @@ consteval void diagnose_fwd_algorithm_signature() constexpr int large_tensor_matches = count_matches_large_tensor(); constexpr int tile_matches = count_matches_tile(); - // Find maximum matches across all variants - constexpr int max_1 = xdl_v3_matches > xdl_matches ? xdl_v3_matches : xdl_matches; - constexpr int max_2 = wmma_matches > dl_matches ? wmma_matches : dl_matches; - constexpr int max_3 = max_1 > max_2 ? max_1 : max_2; - constexpr int max_4 = max_3 > large_tensor_matches ? max_3 : large_tensor_matches; - constexpr int max_matches = max_4 > tile_matches ? max_4 : tile_matches; - // Check whether we have XDL or WMMA algorithm if constexpr (SpecifiesGridwiseFwdXdlGemm) { + constexpr int max_1 = xdl_v3_matches > xdl_matches ? xdl_v3_matches : xdl_matches; + constexpr int max_2 = max_1 > dl_matches ? max_1 : dl_matches; + constexpr int max_matches = large_tensor_matches > max_2 ? large_tensor_matches : max_2; + if constexpr(max_matches == xdl_v3_matches) { using Alg = FwdXdlV3Algorithm; static_assert(Alg::is_valid(), Alg::message()); } else if constexpr(max_matches == xdl_matches) { using Alg = FwdXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); + } else if constexpr(max_matches == dl_matches) { + using Alg = FwdDlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); } else if constexpr (max_matches == large_tensor_matches) { using Alg = LargeTensorAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } + } } else if constexpr (SpecifiesGridwiseWmmaGemm) { @@ -627,6 +633,13 @@ consteval void diagnose_fwd_algorithm_signature() } else { + // Find maximum matches across all variants + constexpr int max_1 = xdl_v3_matches > xdl_matches ? xdl_v3_matches : xdl_matches; + constexpr int max_2 = wmma_matches > dl_matches ? wmma_matches : dl_matches; + constexpr int max_3 = max_1 > max_2 ? max_1 : max_2; + constexpr int max_4 = max_3 > large_tensor_matches ? max_3 : large_tensor_matches; + constexpr int max_matches = max_4 > tile_matches ? max_4 : tile_matches; + // If we cannot match with neither WMMA nor XDL, try all algorithms for diagnostics // and see whichi is the closest match. if constexpr(max_matches == xdl_v3_matches) { @@ -663,35 +676,78 @@ consteval void diagnose_bwd_weight_algorithm_signature() constexpr int two_stage_xdl_matches = count_matches_bwd_two_stage_xdl(); constexpr int dl_matches = count_matches_bwd_dl(); constexpr int multi_d_xdl_matches = count_matches_bwd_multi_d_xdl(); + constexpr int wmma_v3_matches = count_matches_bwd_wmma_v3(); - constexpr int max1 = xdl_v3_matches > xdl_matches ? xdl_v3_matches : xdl_matches; - constexpr int max2 = max1 > two_stage_xdl_matches ? max1 : two_stage_xdl_matches; - constexpr int max3 = max2 > dl_matches ? max2 : dl_matches; - constexpr int max_matches = max3 > multi_d_xdl_matches ? max3 : multi_d_xdl_matches; + // Check whether we have XDL or WMMA algorithm + if constexpr (SpecifiesGridwiseBwdXdlGemm) + { + constexpr int max1 = xdl_v3_matches > xdl_matches ? xdl_v3_matches : xdl_matches; + constexpr int max2 = max1 > two_stage_xdl_matches ? max1 : two_stage_xdl_matches; + constexpr int max3 = max2 > dl_matches ? max2 : dl_matches; + constexpr int max_matches = max3 > multi_d_xdl_matches ? max3 : multi_d_xdl_matches; - if constexpr (max_matches == xdl_matches) { - using Alg = BwdXdlAlgorithm; - static_assert(Alg::is_valid(), Alg::message()); - } - else if constexpr (max_matches == xdl_v3_matches) { - using Alg = BwdXdlV3Algorithm; - static_assert(Alg::is_valid(), Alg::message()); + if constexpr (max_matches == xdl_matches) { + using Alg = BwdXdlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == xdl_v3_matches) { + using Alg = BwdXdlV3Algorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == two_stage_xdl_matches) { + using Alg = BwdTwoStageXdlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == dl_matches) { + using Alg = BwdDlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == multi_d_xdl_matches) { + using Alg = BwdMultiDXdlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } } - else if constexpr (max_matches == two_stage_xdl_matches) { - using Alg = BwdTwoStageXdlAlgorithm; - static_assert(Alg::is_valid(), Alg::message()); + else if constexpr (SpecifiesGridwiseWmmaGemm) + { + constexpr int max_matches = wmma_v3_matches; + if constexpr (max_matches == wmma_v3_matches) { + using Alg = BwdWmmaV3Algorithm; + static_assert(Alg::is_valid(), Alg::message()); + } } - else if constexpr (max_matches == dl_matches) { - using Alg = BwdDlAlgorithm; - static_assert(Alg::is_valid(), Alg::message()); - } - else if constexpr (max_matches == multi_d_xdl_matches) { - using Alg = BwdMultiDXdlAlgorithm; - static_assert(Alg::is_valid(), Alg::message()); - } - else { - // This should never happen - static_assert(false, "Internal Error: No matching algorithm variant found for diagnostics."); + else + { + // If we cannot match with neither WMMA nor XDL, try all algorithms for diagnostics + // and see which is the closest match. + constexpr int max1 = xdl_v3_matches > xdl_matches ? xdl_v3_matches : xdl_matches; + constexpr int max2 = max1 > two_stage_xdl_matches ? max1 : two_stage_xdl_matches; + constexpr int max3 = max2 > dl_matches ? max2 : dl_matches; + constexpr int max_matches = max3 > multi_d_xdl_matches ? max3 : multi_d_xdl_matches; + + if constexpr (max_matches == xdl_matches) { + using Alg = BwdXdlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == xdl_v3_matches) { + using Alg = BwdXdlV3Algorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == two_stage_xdl_matches) { + using Alg = BwdTwoStageXdlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == dl_matches) { + using Alg = BwdDlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else if constexpr (max_matches == multi_d_xdl_matches) { + using Alg = BwdMultiDXdlAlgorithm; + static_assert(Alg::is_valid(), Alg::message()); + } + else { + // This should never happen + static_assert(false, "Internal Error: No matching algorithm variant found for diagnostics."); + } } } diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp index 0ddeae2eed..203280c063 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp @@ -10,13 +10,12 @@ #include "ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp" #include "ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp" #include "ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp" -//#include "ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp" #include "ck_tile/builder/factory/helpers/ck/conv_thread_block.hpp" namespace ck_tile::builder::factory { // Factory for DeviceGroupedConvBwdWeight_Dl instance -// of a grouped forward convolution kernel. +// of a grouped bwd weight convolution kernel. template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp index f4787ddc4a..2c11f3b436 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp @@ -17,8 +17,8 @@ namespace ck_tile::builder::factory { -// Factory for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle instance -// of a grouped forward convolution kernel. +// Factory for DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle instance +// of a grouped bwd weight convolution kernel. template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp index b9852127e8..fae8ad7b87 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp @@ -17,8 +17,8 @@ namespace ck_tile::builder::factory { -// Factory for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle instance -// of a grouped forward convolution kernel. +// Factory for DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle instance +// of a grouped bwd weight convolution kernel. template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp new file mode 100644 index 0000000000..ab75f5b072 --- /dev/null +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp @@ -0,0 +1,104 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_wmma_cshuffle_v3.hpp" +#include "ck_tile/builder/conv_signature_concepts.hpp" +#include "ck_tile/builder/conv_algorithm_concepts.hpp" +#include "ck_tile/builder/conv_algorithm_limits.hpp" +#include "ck_tile/builder/builder_utils.hpp" +#include "ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp" +#include "ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp" +#include "ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp" +#include "ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp" +#include "ck_tile/builder/factory/helpers/ck/conv_block_transfer.hpp" +#include "ck_tile/builder/factory/helpers/ck/conv_thread_block.hpp" + +namespace ck_tile::builder::factory { + +// Factory for DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 instance +// of a grouped bwd weight convolution kernel. +template + requires ConvDirectionIsBackwardWeight +struct ConvBwdWeightWmmaV3Factory +{ + static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; + using Layouts = internal::ConvTensorLayouts; + using Types = internal::BwdWeightConvTensorDataTypes; + using Ops = internal::ElementwiseOps; + using AlgorithmType = decltype(ALGORITHM); + + static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); + + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + static constexpr auto A_BLOCK_TRANSFER = + internal::SetBwdConvBlockTransfer(); + static constexpr auto B_BLOCK_TRANSFER = + internal::SetBwdConvBlockTransfer(); + static constexpr auto C_BLOCK_TRANSFER = internal::SetCBlockTransfer(); + static constexpr auto BLOCK_GEMM = internal::SetBlockGemm(); + + // Check limits for the algorithm parameters. + // TODO: Add more limits checks as needed. + static_assert(InputVectorTransferLimits, "Invalid A block transfer config"); + static_assert(InputVectorTransferLimits, "Invalid B block transfer config"); + static_assert(OutputVectorTransferLimits, "Invalid C block transfer config"); + static_assert(AccessOrderLimits3D, "Invalid A thread cluster access order"); + static_assert(AccessOrderLimits3D, "Invalid B thread cluster access order"); + static_assert(AccessOrderLimits3D, "Invalid A source access order"); + static_assert(AccessOrderLimits3D, "Invalid B source access order"); + + // The forward convolution kernel class instance. + using Instance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Wmma_CShuffleV3< + SPATIAL_DIM, + typename Layouts::InLayout, + typename Layouts::WeiLayout, + typename Layouts::OutLayout, + typename Types::InDataType, + typename Types::WeiDataType, + typename Types::OutDataType, + typename Types::AccDataType, + typename Ops::InElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::OutElementwiseOp, + BWD_CONV_SPECIALIZATION, + BLOCK.block_size, + BLOCK.per_block.m, + BLOCK.per_block.n, + BLOCK.per_block.k, + GRIDWISE_GEMM.k1, + GRIDWISE_GEMM.m_per_wmma, + GRIDWISE_GEMM.n_per_wmma, + GRIDWISE_GEMM.m_wmma_per_wave, + GRIDWISE_GEMM.n_wmma_per_wave, + to_sequence_v, + to_sequence_v, + to_sequence_v, + A_BLOCK_TRANSFER.src_vector_dim, + A_BLOCK_TRANSFER.src_scalar_per_vector, + A_BLOCK_TRANSFER.lds_dst_scalar_per_vector, + A_BLOCK_TRANSFER.lds_padding, + to_sequence_v, + to_sequence_v, + to_sequence_v, + B_BLOCK_TRANSFER.src_vector_dim, + B_BLOCK_TRANSFER.src_scalar_per_vector, + B_BLOCK_TRANSFER.lds_dst_scalar_per_vector, + B_BLOCK_TRANSFER.lds_padding, + C_BLOCK_TRANSFER.m_xdl_per_wave_per_shuffle, + C_BLOCK_TRANSFER.n_xdl_per_wave_per_shuffle, + to_sequence_v, + C_BLOCK_TRANSFER.scalar_per_vector, + BLOCK_GEMM.scheduler, + BLOCK_GEMM.pipeline_version, + typename Types::InComputeType, + typename Types::WeiComputeType, + ALGORITHM.max_transpose_transfer_src_scalar_per_vector, + ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; +}; + +} // namespace ck_tile::builder::factory diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp index 8790121ed9..8ffbb495ec 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp @@ -17,8 +17,8 @@ namespace ck_tile::builder::factory { -// Factory for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle instance -// of a grouped forward convolution kernel. +// Factory for DeviceGroupedConvBwdWeight_Xdl_CShuffle instance +// of a grouped bwd weight convolution kernel. template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp index 14121be940..20aac55f31 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp @@ -17,8 +17,8 @@ namespace ck_tile::builder::factory { -// Factory for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle instance -// of a grouped forward convolution kernel. +// Factory for DeviceGroupedConvBwdWeight_Xdl_CShuffleV3 instance +// of a grouped bwd weight convolution kernel. template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp index e5789ec05f..6aecd779db 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp @@ -73,6 +73,7 @@ #include "ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp" #include "ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp" #include "ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp" +#include "ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp" namespace ck_tile::builder::factory { @@ -173,8 +174,7 @@ constexpr auto make_conv_instance() } else if constexpr (BwdWmmaV3Algorithm::is_valid()) { - static_assert(false, - "Backward weight convolution: WMMA V3 algorithm not yet implemented."); + return typename ConvBwdWeightWmmaV3Factory::Instance{}; } else { diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index 90c6d20c83..f72a72f5fb 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -170,6 +170,7 @@ add_ck_builder_test(test_ckb_build_bwd_weight_instances conv/ck/test_ckb_conv_bwd_weight_multi_d_xdl_cshuffle.cpp conv/ck/test_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp conv/ck/test_ckb_conv_bwd_weight_dl.cpp + conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp conv/ck_tile/test_ckb_conv_bwd_weight_2d_fp16_v3.cpp ) target_link_libraries(test_ckb_build_bwd_weight_instances PRIVATE utility) diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp new file mode 100644 index 0000000000..268bdee5be --- /dev/null +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp @@ -0,0 +1,44 @@ +// 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" +#include "utils/conv_algorithm_type_utils.hpp" +#include "ck_tile/host/device_prop.hpp" + +namespace ckb = ck_tile::builder; +namespace ckt = ck_tile::builder::test; +namespace cku = ck_tile::builder::test_utils; +using enum ck_tile::builder::TensorLayout; + +constexpr auto SIGNATURE = + ckt::ConvSignature{.spatial_dim = 1, + .direction = ckb::ConvDirection::BACKWARD_WEIGHT, + .data_type = ckb::DataType::BF16, + .accumulation_data_type = ckb::DataType::FP32, + .input = {.config = {.layout = NGCW}}, + .weight = {.config = {.layout = GKXC}}, + .output = {.config = {.layout = NGKW}}}; + +constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3{} + .with_thread_block(cku::ThreadBlock_64_32x32x32) + .with_gemm_config(cku::GemmParams_Wmma_2x1_per_wave) + .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3) + .with_bwd_specialization(ckb::ConvSpecialization::FILTER_1X1_STRIDE1_PAD0) + .with_block_gemm(cku::BlockGemmDesc_v2_intrawave) + .with_transpose_params(4,4); + +using Builder = ckb::ConvBuilder; +using Instance = Builder::Instance; + +TEST(BwdWeight_1DBf16_Wmma_CShuffle_V3, Create) +{ + const auto expected_transfer_parameters = to_string(ALGORITHM); + cku::run_test({"DeviceGroupedConvBwdWeight_Wmma_CShuffleV3", + expected_transfer_parameters, + "Filter1x1Stride1Pad0", + "NGCW,GKXC,NGKW", + "PassThrough,PassThrough,PassThrough", + "Intrawave", + "v2"}); +} diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp index ba3173cc9a..bfd16814da 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -237,6 +237,15 @@ inline std::string to_string(DlEpilogue t) return oss.str(); } +template <> +inline std::string to_string(TransposeParams_ t) +{ + std::ostringstream oss; + oss << t.max_transpose_transfer_src_scalar_per_vector << "," + << t.max_transpose_transfer_dst_scalar_per_vector; + return oss.str(); +} + template <> inline std::string to_string>(DlTransfer<4> t) { @@ -414,6 +423,16 @@ inline std::string to_string +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_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_DeviceGroupedConvBwdWeight_TwoStage_Xdl_CShuffle t)