From 75d20e08f115e6bcea4a70e0d5fd393ef94f07d3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Wed, 14 Jan 2026 09:29:04 -0500 Subject: [PATCH] clang-format --- .../builder/conv_algorithm_concepts.hpp | 16 +- .../builder/factory/conv_algorithms.hpp | 81 +++-- ...onv_bwd_weight_multi_d_wmma_v3_factory.hpp | 2 +- .../conv_bwd_weight_multi_d_xdl_factory.hpp | 4 +- ...v_bwd_weight_two_stage_wmma_v3_factory.hpp | 2 +- .../conv_bwd_weight_two_stage_xdl_factory.hpp | 4 +- .../factory/conv_bwd_weight_wmma_factory.hpp | 2 +- .../conv_bwd_weight_wmma_v3_factory.hpp | 2 +- .../factory/conv_bwd_weight_xdl_factory.hpp | 5 +- .../conv_bwd_weight_xdl_v3_factory.hpp | 4 +- .../builder/factory/conv_dispatcher.hpp | 67 ++-- .../builder/factory/conv_fwd_v3_factory.hpp | 4 +- .../helpers/ck/conv_block_transfer.hpp | 76 ++-- .../factory/helpers/ck/conv_tuning_params.hpp | 9 +- .../builder/include/ck_tile/builder/types.hpp | 10 +- ...test_ckb_conv_bwd_weight_wmma_cshuffle.cpp | 13 +- .../conv/ck/test_ckb_conv_fwd_2d_bf16.cpp | 15 +- .../test/impl/conv_algorithm_types.hpp | 28 +- .../builder/test/test_conv_description.cpp | 53 +-- .../test/utils/ckb_conv_test_configs.hpp | 330 ++++++++++-------- .../test/utils/conv_algorithm_type_utils.hpp | 83 +++-- 21 files changed, 432 insertions(+), 378 deletions(-) diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp index adec564306..ce60bf4a94 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -64,8 +64,9 @@ concept InputTileThreadClusterDescriptor4D = requires(T t) { }; template -concept InputTileThreadClusterDescriptor = (ThreadClusterRank == 3 && InputTileThreadClusterDescriptor3D) || - (ThreadClusterRank == 4 && InputTileThreadClusterDescriptor4D); +concept InputTileThreadClusterDescriptor = + (ThreadClusterRank == 3 && InputTileThreadClusterDescriptor3D) || + (ThreadClusterRank == 4 && InputTileThreadClusterDescriptor4D); // Concept for thread cluster dimensions for GEMM output tensor. template @@ -285,7 +286,6 @@ template concept TransposeTransferWellDefinedIfProvided = !HasTransposeTransfer || SpecifiesTransposeTransfer; - /******************************************** */ /* Algorithm specialization concepts */ /******************************************** */ @@ -328,14 +328,12 @@ concept SpecifiesGenericInstance = !requires { }; template -concept SpecifiesXdl = requires { - requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::XDL; -}; +concept SpecifiesXdl = + requires { requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::XDL; }; template -concept SpecifiesWmma = requires { - requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::WMMA; -}; +concept SpecifiesWmma = + requires { requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::WMMA; }; template concept SpecifiesValidWarpGemm = SpecifiesXdl || SpecifiesWmma; 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 5c169d96a7..ad72609f84 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -14,31 +14,32 @@ concept SpecifiesTileTransferParameters = // Base algorithm concepts template -concept ConvAlgorithm = - ConvAlgorithmDescriptor && SpecifiesThreadBlock && - SpecifiesTileTransferParameters && - SpecifiesWarpGemm && SpecifiesGemmPipeline;; +concept ConvAlgorithm = ConvAlgorithmDescriptor && SpecifiesThreadBlock && + SpecifiesTileTransferParameters && + SpecifiesWarpGemm && SpecifiesGemmPipeline; +; template -concept FwdAlgorithm = ConvAlgorithm && SpecifiesFwdConvSpecialization; +concept FwdAlgorithm = ConvAlgorithm && SpecifiesFwdConvSpecialization; template concept FwdAlgorithmV3 = FwdAlgorithm && SpecifiesPipelineV3; template -concept BwdAlgorithm = ConvAlgorithm && SpecifiesBwdWeightConvSpecialization; +concept BwdAlgorithm = + ConvAlgorithm && SpecifiesBwdWeightConvSpecialization; template concept BwdAlgorithmV3 = BwdAlgorithm && SpecifiesPipelineV3; template -concept DlAlgorithm = - ConvAlgorithmDescriptor && SpecifiesThreadBlock && - SpecifiesDlThreadConfig && SpecifiesDlThreadCluster && SpecifiesDlEpilogue; +concept DlAlgorithm = + ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesDlThreadConfig && + SpecifiesDlThreadCluster && SpecifiesDlEpilogue; template -concept FwdDlAlgorithmBase = DlAlgorithm && SpecifiesFwdConvSpecialization && - SpecifiesDlFwdBlockTransfer && SpecifiesGemmSpecialization; +concept FwdDlAlgorithmBase = DlAlgorithm && SpecifiesFwdConvSpecialization && + SpecifiesDlFwdBlockTransfer && SpecifiesGemmSpecialization; template concept FwdXdlAlgorithmBase = FwdAlgorithm && SpecifiesXdl; @@ -57,20 +58,23 @@ concept BwdWmmaV3AlgorithmBase = BwdAlgorithmV3 && SpecifiesWmm // Reference algorithm concept template -concept ReferenceAlgorithm = ConvAlgorithmDescriptor && SpecifiesReferenceAlgorithm; +concept ReferenceAlgorithm = + ConvAlgorithmDescriptor && SpecifiesReferenceAlgorithm; // Tile-based algorithm concept template -concept TileAlgorithm = ConvAlgorithmDescriptor && SpecifiesTileThreadBlock && - SpecifiesTileTransfer && SpecifiesTileConvSpecialization && - SpecifiesTileBlockGemm && SpecifiesTileOptimizations; +concept TileAlgorithm = + ConvAlgorithmDescriptor && SpecifiesTileThreadBlock && + SpecifiesTileTransfer && SpecifiesTileConvSpecialization && + SpecifiesTileBlockGemm && SpecifiesTileOptimizations; // FWD XDL algorithm concepts template concept FwdXdlAlgorithm = FwdXdlAlgorithmBase && SpecifiesGenericInstance; template -concept LargeTensorAlgorithm = FwdXdlAlgorithmBase && SpecifiesLargeTensorSupport; +concept LargeTensorAlgorithm = + FwdXdlAlgorithmBase && SpecifiesLargeTensorSupport; template concept FwdXdlV3Algorithm = FwdAlgorithmV3 && SpecifiesXdl; @@ -82,63 +86,68 @@ concept FwdWmmaAlgorithm = FwdAlgorithm && SpecifiesWmma // FWD DL algorithms template concept FwdDlAlgorithm = FwdDlAlgorithmBase; - + // BWD weight XDL algorithm concepts template concept BwdXdlAlgorithm = - BwdXdlAlgorithmBase && SpecifiesTransposeTransfer && + BwdXdlAlgorithmBase && SpecifiesTransposeTransfer && SpecifiesGenericInstance; template -concept BwdMultiDXdlAlgorithm = BwdXdlAlgorithmBase && SpecifiesMultipleDSupport; +concept BwdMultiDXdlAlgorithm = + BwdXdlAlgorithmBase && SpecifiesMultipleDSupport; template concept BwdXdlV3Algorithm = BwdXdlV3AlgorithmBase; template -concept BwdTwoStageXdlAlgorithm = BwdXdlV3AlgorithmBase && SpecifiesTransposeTransfer && - SpecifiesTwoStageSupport; +concept BwdTwoStageXdlAlgorithm = + BwdXdlV3AlgorithmBase && SpecifiesTransposeTransfer && + SpecifiesTwoStageSupport; // BWD weight WMMA algorithm concepts template concept BwdWmmaAlgorithm = - BwdWmmaAlgorithmBase && SpecifiesNumPrefetchStages && + BwdWmmaAlgorithmBase && SpecifiesNumPrefetchStages && SpecifiesGemmPipeline && SpecifiesGenericInstance; template -concept BwdMultiDWmmaV3Algorithm = BwdWmmaV3AlgorithmBase && SpecifiesMultipleDSupport; +concept BwdMultiDWmmaV3Algorithm = + BwdWmmaV3AlgorithmBase && SpecifiesMultipleDSupport; template concept BwdWmmaV3Algorithm = - BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer && + BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer && SpecifiesGenericInstance; template -concept BwdTwoStageWmmaV3Algorithm = BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer && - SpecifiesTwoStageSupport; +concept BwdTwoStageWmmaV3Algorithm = + BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer && + SpecifiesTwoStageSupport; // BWD weight DL algorithms template -concept BwdDlAlgorithm = DlAlgorithm && SpecifiesBwdWeightConvSpecialization && - SpecifiesDlBwdBlockTransfer; +concept BwdDlAlgorithm = + DlAlgorithm && SpecifiesBwdWeightConvSpecialization && + SpecifiesDlBwdBlockTransfer; // Concepts for valid XDL/WMMA algorithms template -concept SpecifiesValidFwdXdlAlgorithm = -FwdXdlAlgorithm || FwdXdlV3Algorithm || LargeTensorAlgorithm; +concept SpecifiesValidFwdXdlAlgorithm = + FwdXdlAlgorithm || FwdXdlV3Algorithm || LargeTensorAlgorithm; template concept SpecifiesValidFwdWmmaAlgorithm = FwdWmmaAlgorithm; template -concept SpecifiesValidBwdXdlAlgorithm = - BwdXdlAlgorithm || BwdXdlV3Algorithm || - BwdTwoStageXdlAlgorithm || BwdMultiDXdlAlgorithm; +concept SpecifiesValidBwdXdlAlgorithm = + BwdXdlAlgorithm || BwdXdlV3Algorithm || BwdTwoStageXdlAlgorithm || + BwdMultiDXdlAlgorithm; template -concept SpecifiesValidBwdWmmaAlgorithm = - BwdWmmaAlgorithm || BwdWmmaV3Algorithm || - BwdTwoStageWmmaV3Algorithm || BwdMultiDWmmaV3Algorithm; +concept SpecifiesValidBwdWmmaAlgorithm = + BwdWmmaAlgorithm || BwdWmmaV3Algorithm || BwdTwoStageWmmaV3Algorithm || + BwdMultiDWmmaV3Algorithm; template concept FwdWarpGemmOrDL = SpecifiesValidWarpGemm || FwdDlAlgorithm; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp index 8333a784ba..9f86028c33 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp @@ -34,7 +34,7 @@ struct ConvBwdWeightMultiDWmmaV3Factory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); 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 9f8dae1de8..ec67a3c6b0 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 @@ -34,8 +34,8 @@ struct ConvBwdWeightMultiDXdlFactory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp index 9409f7e8bd..afdc5e1a17 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp @@ -34,7 +34,7 @@ struct ConvBwdWeightTwoStageWmmaV3Factory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); 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 986b9fd946..cac6de591a 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 @@ -34,8 +34,8 @@ struct ConvBwdWeightTwoStageXdlFactory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp index a1fc53f7c3..0fd089b6c0 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp @@ -34,7 +34,7 @@ struct ConvBwdWeightWmmaFactory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm; static constexpr auto GRIDWISE_GEMM_PIPELINE_VERSION = internal::SetGridwiseGemmPipelineVersion(); 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 index ba5302f82c..cb3c87905d 100644 --- 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 @@ -34,7 +34,7 @@ struct ConvBwdWeightWmmaV3Factory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); 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 1b0cc36747..b382b7d3d9 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 @@ -34,8 +34,8 @@ struct ConvBwdWeightXdlFactory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = @@ -57,7 +57,6 @@ struct ConvBwdWeightXdlFactory "A nd B block transfer vector load size need to be the same"); static constexpr size_t GMEM_VECTOR_LOAD_SIZE = A_BLOCK_TRANSFER.global_memory_vector_load_size; - // The forward convolution kernel class instance. using Instance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle< SPATIAL_DIM, 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 f37da289c7..9c9bbd21af 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 @@ -34,8 +34,8 @@ struct ConvBwdWeightXdlV3Factory static constexpr auto BWD_CONV_SPECIALIZATION = internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = 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 965a231cd1..cb51cb70dc 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp @@ -117,7 +117,7 @@ constexpr auto make_conv_instance() // Forward direction (supports most algorithm variants) else if constexpr(ConvDirectionIsForward) { - if constexpr (SpecifiesXdl) + if constexpr(SpecifiesXdl) { if constexpr(FwdXdlV3Algorithm) { @@ -129,30 +129,31 @@ constexpr auto make_conv_instance() } else if constexpr(LargeTensorAlgorithm) { - return typename ConvFwdLargeTensorFactory::Instance{}; + return + typename ConvFwdLargeTensorFactory::Instance{}; } else { static_assert( SpecifiesValidFwdXdlAlgorithm, - "No suitable forward convolution XDL kernel factory found for the provided ALGORITHM. " + "No suitable forward convolution XDL kernel factory found for the provided " + "ALGORITHM. " "The ALGORITHM must satisfy requirements for one of: XDL V3, generic XDL, " "DL (NHWC layout), or Large Tensor variant."); } } - else if constexpr (SpecifiesWmma) + else if constexpr(SpecifiesWmma) { if constexpr(FwdWmmaAlgorithm) { return typename ConvFwdWmmaFactory::Instance{}; } - else + else { - static_assert(FwdWmmaAlgorithm, - "Did not find matching WMMA factory."); + static_assert(FwdWmmaAlgorithm, "Did not find matching WMMA factory."); } } - else if constexpr (FwdDlAlgorithm) + else if constexpr(FwdDlAlgorithm) { return typename ConvFwdDlFactory::Instance{}; } @@ -173,22 +174,23 @@ constexpr auto make_conv_instance() // Backward weight direction (will expand with more algorithms in the future) else if constexpr(ConvDirectionIsBackwardWeight) { - if constexpr (SpecifiesXdl) + if constexpr(SpecifiesXdl) { // Start from more specialized and end with least specialized. if constexpr(BwdTwoStageXdlAlgorithm) { - return - typename ConvBwdWeightTwoStageXdlFactory::Instance{}; + return typename ConvBwdWeightTwoStageXdlFactory:: + Instance{}; } else if constexpr(BwdMultiDXdlAlgorithm) { - return - typename ConvBwdWeightMultiDXdlFactory::Instance{}; + return typename ConvBwdWeightMultiDXdlFactory:: + Instance{}; } else if constexpr(BwdXdlV3Algorithm) { - return typename ConvBwdWeightXdlV3Factory::Instance{}; + return + typename ConvBwdWeightXdlV3Factory::Instance{}; } else if constexpr(BwdXdlAlgorithm) { @@ -196,46 +198,51 @@ constexpr auto make_conv_instance() } else { - static_assert( - SpecifiesValidBwdXdlAlgorithm, - "No suitable backward weight convolution XDL kernel factory found for the provided ALGORITHM. " - "The ALGORITHM must satisfy requirements for one of: Two-Stage XDL, Multi-D XDL, DL, " - "generic XDL, or XDL V3 variant."); + static_assert(SpecifiesValidBwdXdlAlgorithm, + "No suitable backward weight convolution XDL kernel factory found " + "for the provided ALGORITHM. " + "The ALGORITHM must satisfy requirements for one of: Two-Stage XDL, " + "Multi-D XDL, DL, " + "generic XDL, or XDL V3 variant."); } } - else if constexpr (SpecifiesWmma) + else if constexpr(SpecifiesWmma) { // Start from more specialized and end with least specialized. if constexpr(BwdTwoStageWmmaV3Algorithm) { - return typename ConvBwdWeightTwoStageWmmaV3Factory::Instance{}; + return typename ConvBwdWeightTwoStageWmmaV3Factory:: + Instance{}; } else if constexpr(BwdMultiDWmmaV3Algorithm) { - return typename ConvBwdWeightMultiDWmmaV3Factory::Instance{}; + return typename ConvBwdWeightMultiDWmmaV3Factory:: + Instance{}; } else if constexpr(BwdWmmaV3Algorithm) { - return typename ConvBwdWeightWmmaV3Factory::Instance{}; + return + typename ConvBwdWeightWmmaV3Factory::Instance{}; } else if constexpr(BwdWmmaAlgorithm) { return typename ConvBwdWeightWmmaFactory::Instance{}; } - else + else { - static_assert( - SpecifiesValidBwdWmmaAlgorithm, - "No suitable backward weight convolution WMMA kernel factory found for the provided ALGORITHM. " - "The ALGORITHM must satisfy requirements for one of: Two-Stage WMMA V3, Multi-D WMMA V3, " - "WMMA V3, or generic WMMA variant."); + static_assert(SpecifiesValidBwdWmmaAlgorithm, + "No suitable backward weight convolution WMMA kernel factory found " + "for the provided ALGORITHM. " + "The ALGORITHM must satisfy requirements for one of: Two-Stage WMMA " + "V3, Multi-D WMMA V3, " + "WMMA V3, or generic WMMA variant."); } } else if constexpr(BwdDlAlgorithm) { return typename ConvBwdWeightDlFactory::Instance{}; } - else + else { static_assert(BwdWarpGemmOrDL, "Backward convolution: Algorithm must specify either DL, XDL or WMMA."); diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp index e6f71b0b25..ad6f91b0a4 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp @@ -41,8 +41,8 @@ struct ConvFwdXdlV3Factory static constexpr internal::ConvSpec SPECIALIZATION{.conv_spec = FWD_CONV_SPECIALIZATION, .gemm_spec = GEMM_SPECIALIZATION}; - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm; static constexpr auto A_BLOCK_TRANSFER = internal::SetFwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_block_transfer.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_block_transfer.hpp index c48404dd47..6259e239a9 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_block_transfer.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_block_transfer.hpp @@ -17,11 +17,11 @@ struct BlockTransfer ck::Array thread_cluster_order{}; ck::Array src_access_order{}; size_t global_memory_vector_load_size = 0; - size_t src_vector_dim = 0; - size_t src_scalar_per_vector = 0; - size_t lds_dst_scalar_per_vector = 0; - bool is_direct_load = false; - bool lds_padding = false; + size_t src_vector_dim = 0; + size_t src_scalar_per_vector = 0; + size_t lds_dst_scalar_per_vector = 0; + bool is_direct_load = false; + bool lds_padding = false; }; template @@ -33,15 +33,15 @@ constexpr BlockTransfer<> SetFwdConvBlockTransfer() auto& lds_cfg = TRANSFER.lds_transfer; return BlockTransfer<>{ - .thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1}, - .thread_cluster_order = {block_order.order[0], block_order.order[1], block_order.order[2]}, - .src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]}, + .thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1}, + .thread_cluster_order = {block_order.order[0], block_order.order[1], block_order.order[2]}, + .src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]}, .global_memory_vector_load_size = lds_cfg.global_memory_vector_load_size, - .src_vector_dim = lds_cfg.src_vector_dim, - .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, - .lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector, - .is_direct_load = lds_cfg.is_direct_load, - .lds_padding = lds_cfg.lds_padding, + .src_vector_dim = lds_cfg.src_vector_dim, + .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, + .lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector, + .is_direct_load = lds_cfg.is_direct_load, + .lds_padding = lds_cfg.lds_padding, }; } @@ -60,38 +60,38 @@ constexpr auto SetBwdConvBlockTransfer() if constexpr(array_length == 3) { return BlockTransfer<3>{ - .thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1}, - .thread_cluster_order = {block_order.order[0], - block_order.order[1], - block_order.order[2]}, - .src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]}, + .thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1}, + .thread_cluster_order = {block_order.order[0], + block_order.order[1], + block_order.order[2]}, + .src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]}, .global_memory_vector_load_size = lds_cfg.global_memory_vector_load_size, - .src_vector_dim = lds_cfg.src_vector_dim, - .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, - .lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector, - .lds_padding = lds_cfg.lds_padding, + .src_vector_dim = lds_cfg.src_vector_dim, + .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, + .lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector, + .lds_padding = lds_cfg.lds_padding, }; } else if constexpr(array_length == 4) { return BlockTransfer<4>{ - .thread_cluster_dims = {block_xfer.k_batch_size, - block_xfer.k0, - block_xfer.m_n, - block_xfer.k1}, - .thread_cluster_order = {block_order.order[0], - block_order.order[1], - block_order.order[2], - block_order.order[3]}, - .src_access_order = {src_order.order[0], - src_order.order[1], - src_order.order[2], - src_order.order[3]}, + .thread_cluster_dims = {block_xfer.k_batch_size, + block_xfer.k0, + block_xfer.m_n, + block_xfer.k1}, + .thread_cluster_order = {block_order.order[0], + block_order.order[1], + block_order.order[2], + block_order.order[3]}, + .src_access_order = {src_order.order[0], + src_order.order[1], + src_order.order[2], + src_order.order[3]}, .global_memory_vector_load_size = lds_cfg.global_memory_vector_load_size, - .src_vector_dim = lds_cfg.src_vector_dim, - .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, - .lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector, - .lds_padding = lds_cfg.lds_padding, + .src_vector_dim = lds_cfg.src_vector_dim, + .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, + .lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector, + .lds_padding = lds_cfg.lds_padding, }; } else diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp index 8db9607f34..f123edbfa9 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp @@ -65,11 +65,10 @@ consteval BlockGemmSpec SetBlockGemm() default: throw "Unknown PipelineVersion"; } - return BlockGemmSpec{ - .num_conv_groups_to_merge = BG.num_conv_groups_to_merge, - .num_gemm_k_prefetch_stages = BG.num_gemm_k_prefetch_stages, - .pipeline_version = version, - .scheduler = scheduler}; + return BlockGemmSpec{.num_conv_groups_to_merge = BG.num_conv_groups_to_merge, + .num_gemm_k_prefetch_stages = BG.num_gemm_k_prefetch_stages, + .pipeline_version = version, + .scheduler = scheduler}; } template diff --git a/experimental/builder/include/ck_tile/builder/types.hpp b/experimental/builder/include/ck_tile/builder/types.hpp index af9cc44115..4d09ad8cb4 100644 --- a/experimental/builder/include/ck_tile/builder/types.hpp +++ b/experimental/builder/include/ck_tile/builder/types.hpp @@ -232,12 +232,12 @@ enum class PipelineScheduler enum class ConvAlgorithmSpecialization { - NONE = 0, + NONE = 0, LARGE_TENSOR = 1 << 0, - REFERENCE = 1 << 1, // GPU reference implementation for validation, - TWO_STAGE = 1 << 2, - MULTIPLE_D = 1 << 3, - PIPELINE_V3 = 1 << 4 + REFERENCE = 1 << 1, // GPU reference implementation for validation, + TWO_STAGE = 1 << 2, + MULTIPLE_D = 1 << 3, + PIPELINE_V3 = 1 << 4 }; constexpr ConvAlgorithmSpecialization operator|(ConvAlgorithmSpecialization lhs, diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp index 562599da93..136c8fea40 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp @@ -19,12 +19,13 @@ constexpr auto SIGNATURE = ckt::ConvSignature{.spatial_dim = 3, .weight = {.config = {.layout = GKZYXC}}, .output = {.config = {.layout = NGKDHW}}}; -constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle{} - .with_thread_block(cku::ThreadBlock_64_32x32x32) - .with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave) - .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3) - .with_bwd_specialization(ckb::ConvSpecialization::DEFAULT) - .with_gemm_pipeline(ckb::PipelineVersion::V1, ckb::PipelineScheduler::DEFAULT); +constexpr auto ALGORITHM = + cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle{} + .with_thread_block(cku::ThreadBlock_64_32x32x32) + .with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave) + .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3) + .with_bwd_specialization(ckb::ConvSpecialization::DEFAULT) + .with_gemm_pipeline(ckb::PipelineVersion::V1, ckb::PipelineScheduler::DEFAULT); using Builder = ckb::ConvBuilder; using Instance = Builder::Instance; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp index ebc30272c9..9ed5758b4e 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp @@ -62,14 +62,13 @@ TEST(FwdConvInstances, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = NHWGK}}}; - constexpr auto FwdConvAlgorithm = - ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle_V3{} - .with_thread_block(ThreadBlock_256_256x256x32) - .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) - .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvSpecialization::FILTER_3x3, - GemmSpecialization::MNKPadding) - .with_gemm_pipeline(BlockGemmDesc_v5_intrawave); + constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle_V3{} + .with_thread_block(ThreadBlock_256_256x256x32) + .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) + .with_transfer(Transfer_4x64x1) + .with_fwd_specializations(ConvSpecialization::FILTER_3x3, + GemmSpecialization::MNKPadding) + .with_gemm_pipeline(BlockGemmDesc_v5_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp index 264f69ad81..cf4bc25c0f 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -31,10 +31,10 @@ static_assert(ckb::ThreadBlockDescriptor); struct WarpGemmParams { MatrixInstructionType matrix_instruction; - size_t gemm_m_per_instruction = 0; - size_t gemm_n_per_instruction = 0; - size_t gemm_m_iters_per_wave = 0; - size_t gemm_n_iters_per_wave = 0; + size_t gemm_m_per_instruction = 0; + size_t gemm_n_per_instruction = 0; + size_t gemm_m_iters_per_wave = 0; + size_t gemm_n_iters_per_wave = 0; }; static_assert(ckb::WarpGemmDescriptor); @@ -370,7 +370,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_num_conv_groups_to_merge(size_t num_groups_to_merge) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.gemm_pipeline.num_conv_groups_to_merge = num_groups_to_merge; return result; } @@ -378,7 +378,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_num_gemm_k_prefetch_stages(size_t num_prefetch_stages) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.gemm_pipeline.num_gemm_k_prefetch_stages = num_prefetch_stages; return result; } @@ -387,7 +387,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_gemm_pipeline(const PL& pl) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.gemm_pipeline = pl; return result; } @@ -395,7 +395,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_gemm_pipeline(const PipelineVersion plv) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.gemm_pipeline.pipeline_version = plv; return result; } @@ -403,7 +403,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_gemm_pipeline(const PipelineScheduler sch) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.gemm_pipeline.scheduler = sch; return result; } @@ -411,7 +411,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_gemm_pipeline(const PipelineVersion plv, const PipelineScheduler sch) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.gemm_pipeline.pipeline_version = plv; result.gemm_pipeline.scheduler = sch; return result; @@ -535,7 +535,7 @@ using ConvAlgorithm_Tile_GroupedConvolutionKernel = ConvAlgorithmTemplate; // Reference algorithm descriptor - for GPU reference validation -using ConvAlgorithm_Reference = ConvAlgorithmTemplate>; +using ConvAlgorithm_Reference = ConvAlgorithmTemplate>; // Bwd weight algorithm types using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle = @@ -543,7 +543,8 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle = WarpGemm_, InputOutputTileTransfer_<4>, ConvSpecializationBwdWeight_, - GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm concept. + GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm + // concept. TransposeParams_, AlgorithmSpecialization_<>>; @@ -593,7 +594,8 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle = ConvAlgorithmTemplate, - GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm concept. + GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm + // concept. ConvSpecializationBwdWeight_, AlgorithmSpecialization_>; diff --git a/experimental/builder/test/test_conv_description.cpp b/experimental/builder/test/test_conv_description.cpp index 760ada78ea..c8cfbe6b19 100644 --- a/experimental/builder/test/test_conv_description.cpp +++ b/experimental/builder/test/test_conv_description.cpp @@ -117,57 +117,60 @@ static_assert(!ckb::ConvSignatureDescriptor transfer{ .a = { - .thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2}, - .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 2, - .is_direct_load = false, - .lds_padding = false}, + .thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2}, + .lds_transfer = {.global_memory_vector_load_size = 8, + .src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 2, + .is_direct_load = false, + .lds_padding = false}, .thread_cluster_access_order = {.order = {0, 1, 2}}, .src_access_order = {.order = {0, 1, 2}}, }, .b = { - .thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2}, - .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 2, - .is_direct_load = false, - .lds_padding = false}, + .thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2}, + .lds_transfer = {.global_memory_vector_load_size = 8, + .src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 2, + .is_direct_load = false, + .lds_padding = false}, .thread_cluster_access_order = {.order = {0, 1, 2}}, .src_access_order = {.order = {0, 1, 2}}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 2}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 32, + .gemm_n_block_size = 1, + .gemm_n_per_block = 8}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 2}, }, }; ckb::ConvSpecialization fwd_specialization = ckb::ConvSpecialization::DEFAULT; ckb::GemmSpecialization gemm_specialization = ckb::GemmSpecialization::Default; ckb::test::GemmPipeline gemm_pipeline{.pipeline_version = ckb::PipelineVersion::V4, - .scheduler =ckb::PipelineScheduler::INTRAWAVE}; + .scheduler = ckb::PipelineScheduler::INTRAWAVE}; }; static_assert(ckb::ConvAlgorithmDescriptor); diff --git a/experimental/builder/test/utils/ckb_conv_test_configs.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp index a281ca28ad..4209d708cd 100644 --- a/experimental/builder/test/utils/ckb_conv_test_configs.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp @@ -55,13 +55,13 @@ constexpr DlTransfer<5> DlTransfer5D{.a = DlBlockTransfer_1x8x1x1x1, constexpr InputOutputTileTransfer<> Transfer_4x64x1{ .a = { - .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1}, + .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = false}, + .src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = false}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, @@ -69,21 +69,23 @@ constexpr InputOutputTileTransfer<> Transfer_4x64x1{ { .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = false}, + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = false}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 4}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 32, + .gemm_n_block_size = 1, + .gemm_n_per_block = 8}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 4}, }, }; @@ -92,11 +94,11 @@ constexpr InputOutputTileTransfer<4> BwdTransfer_4x64x1{ { .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1, .k_batch_size = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 4, - .is_direct_load = false, - .lds_padding = true}, + .src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 4, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {0, 3, 1, 2}, .src_access_order = {0, 2, 1, 3}, }, @@ -104,21 +106,23 @@ constexpr InputOutputTileTransfer<4> BwdTransfer_4x64x1{ { .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1, .k_batch_size = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 4, - .is_direct_load = false, - .lds_padding = true}, + .src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 4, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {0, 3, 1, 2}, .src_access_order = {0, 2, 1, 3}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 32, + .gemm_n_block_size = 1, + .gemm_n_per_block = 8}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, }, }; @@ -127,11 +131,11 @@ constexpr InputOutputTileTransfer<> BwdTransfer_4x8x1_4x16x1_v3{ { .thread_cluster = {.k0 = 4, .m_n = 8, .k1 = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 1, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 2, - .is_direct_load = false, - .lds_padding = false}, + .src_vector_dim = 1, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 2, + .is_direct_load = false, + .lds_padding = false}, .thread_cluster_access_order = {2, 0, 1}, .src_access_order = {1, 0, 2}, }, @@ -139,21 +143,23 @@ constexpr InputOutputTileTransfer<> BwdTransfer_4x8x1_4x16x1_v3{ { .thread_cluster = {.k0 = 4, .m_n = 16, .k1 = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 1, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 2, - .is_direct_load = false, - .lds_padding = false}, + .src_vector_dim = 1, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 2, + .is_direct_load = false, + .lds_padding = false}, .thread_cluster_access_order = {2, 0, 1}, .src_access_order = {1, 0, 2}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 8, .gemm_n_block_size = 1, .gemm_n_per_block = 8}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 2}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 8, + .gemm_n_block_size = 1, + .gemm_n_per_block = 8}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 2}, }, }; @@ -162,11 +168,11 @@ constexpr InputOutputTileTransfer<> Transfer_4x64x1_fp8{ { .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, @@ -174,21 +180,23 @@ constexpr InputOutputTileTransfer<> Transfer_4x64x1_fp8{ { .thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1}, .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 32, + .gemm_n_block_size = 1, + .gemm_n_per_block = 8}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, }, }; @@ -196,34 +204,36 @@ constexpr InputOutputTileTransfer<> Transfer_4x16x1{ .a = { .thread_cluster = {.k0 = 4, .m_n = 16, .k1 = 1}, - .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, + .lds_transfer = {.global_memory_vector_load_size = 8, + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .b = { .thread_cluster = {.k0 = 4, .m_n = 16, .k1 = 1}, - .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, + .lds_transfer = {.global_memory_vector_load_size = 8, + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 16, .gemm_n_block_size = 1, .gemm_n_per_block = 4}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 16, + .gemm_n_block_size = 1, + .gemm_n_per_block = 4}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, }, }; @@ -232,34 +242,36 @@ constexpr InputOutputTileTransfer<> Transfer_4x32x1{ .a = { .thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1}, - .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 16, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, + .lds_transfer = {.global_memory_vector_load_size = 8, + .src_vector_dim = 2, + .src_scalar_per_vector = 16, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .b = { .thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1}, - .lds_transfer = {.global_memory_vector_load_size = 8, - .src_vector_dim = 2, - .src_scalar_per_vector = 16, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, + .lds_transfer = {.global_memory_vector_load_size = 8, + .src_vector_dim = 2, + .src_scalar_per_vector = 16, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 4}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 32, + .gemm_n_block_size = 1, + .gemm_n_per_block = 4}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, }, }; @@ -267,68 +279,94 @@ constexpr InputOutputTileTransfer<> Transfer_4x32x1_vector_load_16_generic{ .a = { .thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1}, - .lds_transfer = {.global_memory_vector_load_size = 16, - .src_vector_dim = 2, - .src_scalar_per_vector = 1, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, + .lds_transfer = {.global_memory_vector_load_size = 16, + .src_vector_dim = 2, + .src_scalar_per_vector = 1, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .b = { .thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1}, - .lds_transfer = {.global_memory_vector_load_size = 16, - .src_vector_dim = 2, - .src_scalar_per_vector = 1, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, + .lds_transfer = {.global_memory_vector_load_size = 16, + .src_vector_dim = 2, + .src_scalar_per_vector = 1, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, .thread_cluster_access_order = {1, 0, 2}, .src_access_order = {1, 0, 2}, }, .c = { - .thread_cluster = - {.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 4}, - .epilogue = {.m_xdl_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 1}, + .thread_cluster = {.gemm_m_block_size = 1, + .gemm_m_per_block = 32, + .gemm_n_block_size = 1, + .gemm_n_per_block = 4}, + .epilogue = {.m_xdl_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 1}, }, }; -constexpr WarpGemmParams BwdGemmParams_Xdl_4x4_per_wave{ - .matrix_instruction = MatrixInstructionType::XDL, - .gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 4, .gemm_n_iters_per_wave = 4}; +constexpr WarpGemmParams BwdGemmParams_Xdl_4x4_per_wave{.matrix_instruction = + MatrixInstructionType::XDL, + .gemm_m_per_instruction = 32, + .gemm_n_per_instruction = 32, + .gemm_m_iters_per_wave = 4, + .gemm_n_iters_per_wave = 4}; -constexpr WarpGemmParams BwdGemmParams_Xdl_1x1_per_wave{ - .matrix_instruction = MatrixInstructionType::XDL, - .gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 1, .gemm_n_iters_per_wave = 1}; +constexpr WarpGemmParams BwdGemmParams_Xdl_1x1_per_wave{.matrix_instruction = + MatrixInstructionType::XDL, + .gemm_m_per_instruction = 32, + .gemm_n_per_instruction = 32, + .gemm_m_iters_per_wave = 1, + .gemm_n_iters_per_wave = 1}; -constexpr WarpGemmParams FwdGemmParams_Xdl_4x4_per_wave{ - .matrix_instruction = MatrixInstructionType::XDL, - .gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 4, .gemm_n_iters_per_wave = 4}; +constexpr WarpGemmParams FwdGemmParams_Xdl_4x4_per_wave{.matrix_instruction = + MatrixInstructionType::XDL, + .gemm_m_per_instruction = 32, + .gemm_n_per_instruction = 32, + .gemm_m_iters_per_wave = 4, + .gemm_n_iters_per_wave = 4}; -constexpr WarpGemmParams FwdGemmParams_Xdl_4x2_per_wave{ - .matrix_instruction = MatrixInstructionType::XDL, - .gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 4, .gemm_n_iters_per_wave = 2}; +constexpr WarpGemmParams FwdGemmParams_Xdl_4x2_per_wave{.matrix_instruction = + MatrixInstructionType::XDL, + .gemm_m_per_instruction = 32, + .gemm_n_per_instruction = 32, + .gemm_m_iters_per_wave = 4, + .gemm_n_iters_per_wave = 2}; -constexpr WarpGemmParams FwdGemmParams_Xdl_2x2_per_wave{ - .matrix_instruction = MatrixInstructionType::XDL, - .gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 2}; +constexpr WarpGemmParams FwdGemmParams_Xdl_2x2_per_wave{.matrix_instruction = + MatrixInstructionType::XDL, + .gemm_m_per_instruction = 32, + .gemm_n_per_instruction = 32, + .gemm_m_iters_per_wave = 2, + .gemm_n_iters_per_wave = 2}; -constexpr WarpGemmParams FwdGemmParams_Xdl_2x1_per_wave{ - .matrix_instruction = MatrixInstructionType::XDL, - .gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 1}; +constexpr WarpGemmParams FwdGemmParams_Xdl_2x1_per_wave{.matrix_instruction = + MatrixInstructionType::XDL, + .gemm_m_per_instruction = 32, + .gemm_n_per_instruction = 32, + .gemm_m_iters_per_wave = 2, + .gemm_n_iters_per_wave = 1}; -constexpr WarpGemmParams GemmParams_Wmma_16x16_2x1_per_wave{ - .matrix_instruction = MatrixInstructionType::WMMA, - .gemm_m_per_instruction = 16, .gemm_n_per_instruction = 16, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 1}; +constexpr WarpGemmParams GemmParams_Wmma_16x16_2x1_per_wave{.matrix_instruction = + MatrixInstructionType::WMMA, + .gemm_m_per_instruction = 16, + .gemm_n_per_instruction = 16, + .gemm_m_iters_per_wave = 2, + .gemm_n_iters_per_wave = 1}; -constexpr WarpGemmParams GemmParams_Wmma_16x16_2x2_per_wave{ - .matrix_instruction = MatrixInstructionType::WMMA, - .gemm_m_per_instruction = 16, .gemm_n_per_instruction = 16, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 2}; +constexpr WarpGemmParams GemmParams_Wmma_16x16_2x2_per_wave{.matrix_instruction = + MatrixInstructionType::WMMA, + .gemm_m_per_instruction = 16, + .gemm_n_per_instruction = 16, + .gemm_m_iters_per_wave = 2, + .gemm_n_iters_per_wave = 2}; constexpr ThreadBlock ThreadBlock_256_256x256x32{.block_size = 256, .tile_size = {.m = 256, .n = 256, .k = 32}}; @@ -357,19 +395,19 @@ constexpr ThreadBlock ThreadBlock_128_128x128x32{.block_size = 128, constexpr ThreadBlock ThreadBlock_128_64x64x64{.block_size = 128, .tile_size = {.m = 64, .n = 64, .k = 64}}; -constexpr GemmPipeline BlockGemmDesc_v1_intrawave = { - .pipeline_version = PipelineVersion::V1, .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr GemmPipeline BlockGemmDesc_v1_intrawave = {.pipeline_version = PipelineVersion::V1, + .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr GemmPipeline BlockGemmDesc_v2_intrawave = { - .pipeline_version = PipelineVersion::V2, .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr GemmPipeline BlockGemmDesc_v2_intrawave = {.pipeline_version = PipelineVersion::V2, + .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr GemmPipeline BlockGemmDesc_v3_intrawave = { - .pipeline_version = PipelineVersion::V3, .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr GemmPipeline BlockGemmDesc_v3_intrawave = {.pipeline_version = PipelineVersion::V3, + .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr GemmPipeline BlockGemmDesc_v4_intrawave = { - .pipeline_version = PipelineVersion::V4, .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr GemmPipeline BlockGemmDesc_v4_intrawave = {.pipeline_version = PipelineVersion::V4, + .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr GemmPipeline BlockGemmDesc_v5_intrawave = { - .pipeline_version = PipelineVersion::V5, .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr GemmPipeline BlockGemmDesc_v5_intrawave = {.pipeline_version = PipelineVersion::V5, + .scheduler = PipelineScheduler::INTRAWAVE}; } // 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 index db5aa93a7c..2f66a633fa 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -98,7 +98,8 @@ template <> inline std::string to_string(GemmPipeline t) { std::ostringstream oss; - oss << t.num_gemm_k_prefetch_stages << "," << t.num_conv_groups_to_merge << "," << to_string(t.scheduler) << "," << to_string(t.pipeline_version); + oss << t.num_gemm_k_prefetch_stages << "," << t.num_conv_groups_to_merge << "," + << to_string(t.scheduler) << "," << to_string(t.pipeline_version); return oss.str(); } @@ -123,17 +124,17 @@ inline std::string to_string(InputThreadCluster t) template <> inline std::string to_string(OutputThreadCluster t) { - return array_to_seq( - std::array{t.gemm_m_block_size, t.gemm_m_per_block, t.gemm_n_block_size, t.gemm_n_per_block}); + return array_to_seq(std::array{ + t.gemm_m_block_size, t.gemm_m_per_block, t.gemm_n_block_size, t.gemm_n_per_block}); } template <> inline std::string to_string(LdsTransfer t) { std::ostringstream oss; - oss << t.global_memory_vector_load_size << "," << t.src_vector_dim << "," << t.src_scalar_per_vector << "," << t.lds_dst_scalar_per_vector - << "," << (t.lds_padding ? "true" : "false") << "," - << (t.is_direct_load ? "true" : "false"); + oss << t.global_memory_vector_load_size << "," << t.src_vector_dim << "," + << t.src_scalar_per_vector << "," << t.lds_dst_scalar_per_vector << "," + << (t.lds_padding ? "true" : "false") << "," << (t.is_direct_load ? "true" : "false"); return oss.str(); } @@ -259,7 +260,6 @@ inline std::string to_string(WarpGemm_ t) return to_string(t.warp_gemm); } - template inline std::string to_string(InputOutputTileTransfer_ t) { @@ -319,21 +319,21 @@ inline std::string to_string(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); } - else + else { oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," << t.transfer.b.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); - } + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); + } return oss.str(); } @@ -342,11 +342,11 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," << t.transfer.b.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -367,11 +367,11 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," << t.transfer.b.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -380,10 +380,10 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -392,10 +392,10 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -404,10 +404,10 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -416,10 +416,10 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -430,8 +430,8 @@ inline std::string to_string(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } @@ -441,14 +441,13 @@ inline std::string to_string(t)) - << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size - << "," << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + oss << to_string(static_cast(t)) << "," + << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); } - template <> inline std::string to_string( ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl t) @@ -466,10 +465,10 @@ inline std::string to_string(t)) << "," + oss << to_string(static_cast(t)) << "," << t.transfer.a.lds_transfer.global_memory_vector_load_size << "," - << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); return oss.str(); }