From 881bf916fe62097fc159745c38ab407b5ac40a4f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Mon, 5 Jan 2026 04:45:51 -0500 Subject: [PATCH] clang-format --- .../builder/conv_algorithm_concepts.hpp | 17 +- .../builder/conv_algorithm_diagnostics.hpp | 1790 +++++++++++------ .../builder/factory/conv_algorithms.hpp | 646 +++--- .../factory/conv_bwd_weight_dl_factory.hpp | 11 +- ...onv_bwd_weight_multi_d_wmma_v3_factory.hpp | 120 +- .../conv_bwd_weight_multi_d_xdl_factory.hpp | 17 +- ...v_bwd_weight_two_stage_wmma_v3_factory.hpp | 122 +- .../conv_bwd_weight_two_stage_xdl_factory.hpp | 29 +- .../factory/conv_bwd_weight_wmma_factory.hpp | 27 +- .../conv_bwd_weight_wmma_v3_factory.hpp | 27 +- .../factory/conv_bwd_weight_xdl_factory.hpp | 17 +- .../conv_bwd_weight_xdl_v3_factory.hpp | 29 +- .../builder/factory/conv_dispatcher.hpp | 30 +- .../builder/factory/conv_fwd_dl_factory.hpp | 8 +- .../factory/conv_fwd_large_tensor_factory.hpp | 16 +- .../builder/factory/conv_fwd_v3_factory.hpp | 8 +- .../builder/factory/conv_fwd_wmma_factory.hpp | 8 +- .../builder/factory/conv_fwd_xdl_factory.hpp | 8 +- .../helpers/ck/conv_block_transfer.hpp | 45 +- .../helpers/ck/conv_elementwise_op.hpp | 6 +- .../factory/helpers/ck/conv_tensor_layout.hpp | 14 +- .../factory/helpers/ck/conv_tensor_type.hpp | 4 +- .../factory/helpers/ck/conv_tuning_params.hpp | 8 +- .../conv/ck/test_ckb_conv_bwd_weight_dl.cpp | 10 +- ..._bwd_weight_two_stage_wmma_cshuffle_v3.cpp | 17 +- ...conv_bwd_weight_two_stage_xdl_cshuffle.cpp | 2 +- ...test_ckb_conv_bwd_weight_wmma_cshuffle.cpp | 15 +- ...t_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp | 30 +- ...st_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp | 28 +- .../conv/ck/test_ckb_conv_fwd_1d_bf16.cpp | 2 +- .../test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp | 2 +- .../conv/ck/test_ckb_conv_fwd_2d_bf16.cpp | 3 +- .../conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp | 2 +- .../conv/ck/test_ckb_conv_fwd_2d_fp16.cpp | 2 +- .../conv/ck/test_ckb_conv_fwd_2d_fp32.cpp | 2 +- ...test_ckb_conv_fwd_2d_large_tensor_fp16.cpp | 27 +- .../conv/ck/test_ckb_conv_fwd_3d_fp16.cpp | 2 +- .../conv/ck/test_ckb_conv_fwd_3d_fp32.cpp | 2 +- .../test_ckb_conv_bwd_data_2d_fp16_v3.cpp | 15 +- .../test_ckb_conv_bwd_weight_2d_fp16_v3.cpp | 15 +- .../test/impl/conv_algorithm_types.hpp | 139 +- .../test/test_concept_diagnostics_sync.cpp | 87 +- .../builder/test/test_conv_description.cpp | 16 +- .../test/utils/ckb_conv_test_configs.hpp | 108 +- .../test/utils/conv_algorithm_type_utils.hpp | 15 +- ...bwd_weight_multiple_d_wmma_cshuffle_v3.hpp | 84 +- ...onv_bwd_weight_multiple_d_xdl_cshuffle.hpp | 27 +- 47 files changed, 2195 insertions(+), 1464 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 5a5dfd17b8..2b0f63296b 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -16,7 +16,7 @@ namespace ck_tile::builder { /********************************************************************/ // Common concept for size-related fields -template +template concept SizeType = std::unsigned_integral>; // Concept for thread block dimensions for a GEMM problem. @@ -170,7 +170,7 @@ concept SpecifiesTileThreadBlock = requires { // Concept to check if a struct specifies gridwise XDL GEMM info. template -concept GridwiseFwdXdlGemmDescriptor = requires (T t){ +concept GridwiseFwdXdlGemmDescriptor = requires(T t) { { t.ak1 } -> SizeType; { t.bk1 } -> SizeType; { t.xdl_params } -> GridwiseXdlGemmDescriptor; @@ -178,26 +178,26 @@ concept GridwiseFwdXdlGemmDescriptor = requires (T t){ // Concept to check if a struct specifies gridwise XDL GEMM info. template -concept GridwiseBwdXdlGemmDescriptor = requires (T t){ +concept GridwiseBwdXdlGemmDescriptor = requires(T t) { { t.k1 } -> SizeType; { t.xdl_params } -> GridwiseXdlGemmDescriptor; }; // Concept to check if a struct specifies gridwise XDL GEMM info. template -concept SpecifiesGridwiseFwdXdlGemm = requires (T t) { +concept SpecifiesGridwiseFwdXdlGemm = requires(T t) { { t.gridwise_gemm } -> GridwiseFwdXdlGemmDescriptor; }; // Concept to check if a struct specifies gridwise XDL GEMM info. template -concept SpecifiesGridwiseBwdXdlGemm = requires (T t) { +concept SpecifiesGridwiseBwdXdlGemm = requires(T t) { { t.gridwise_gemm } -> GridwiseBwdXdlGemmDescriptor; }; // Concept to check if a struct specifies gridwise WMMA GEMM info. template -concept SpecifiesGridwiseWmmaGemm = requires (T t){ +concept SpecifiesGridwiseWmmaGemm = requires(T t) { { t.gridwise_gemm } -> GridwiseWmmaGemmDescriptor; }; @@ -209,7 +209,7 @@ concept SpecifiesBlockTransfer = requires(T t) { { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor; }; -// Concept to check if a struct specifies convolution input and output block transfer info +// Concept to check if a struct specifies convolution input and output block transfer info // for 4D thread slices. template concept SpecifiesBlockTransfer4D = requires(T t) { @@ -255,8 +255,7 @@ concept SpecifiesBlockGemm = requires { }; template -concept SpecifiesGridwiseGemmPipeline = requires -{ +concept SpecifiesGridwiseGemmPipeline = requires { { T::pipeline_version } -> std::convertible_to; }; diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp index 46e205d68b..6db35a9ba0 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp @@ -10,373 +10,520 @@ namespace ck_tile::builder::diagnostics { #define CHECK_MARK(cond) (cond ? "[✓]" : "[✗]") // Macro to check a concept and generate both the boolean and the string representation -#define CHECK_CONCEPT(Type, Concept) \ - static constexpr bool c_##Concept = Concept; \ +#define CHECK_CONCEPT(Type, Concept) \ + static constexpr bool c_##Concept = Concept; \ static constexpr const char* s_##Concept = #Concept; // Helper to create diagnostic message line -#define DIAGNOSTIC_LINE(Concept) \ +#define DIAGNOSTIC_LINE(Concept) \ " " + std::string(s_##Concept) + ": " + std::string(CHECK_MARK(c_##Concept)) + "\n" + \ - (c_##Concept ? std::string("") : detailed_diagnostic_##Concept()) + (c_##Concept ? std::string("") : detailed_diagnostic_##Concept()) namespace detail { // Helper to get type information -template -consteval auto get_type_info() -> const char* { +template +consteval auto get_type_info() -> const char* +{ // Returns a descriptive string about the type - if constexpr (std::is_same_v) { + if constexpr(std::is_same_v) + { return " (type: size_t)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: int)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: bool)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: PipelineVersion)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: PipelineScheduler)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: ConvSpecialization)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: GemmSpecialization)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: TileConvSpecialization)"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return " (type: ConvAlgorithmSpecialization)"; - } else if constexpr (std::is_same_v>) { + } + else if constexpr(std::is_same_v>) + { return " (type: std::array)"; - } else if constexpr (std::is_same_v>) { + } + else if constexpr(std::is_same_v>) + { return " (type: std::array)"; - } else if constexpr (std::is_same_v>) { + } + else if constexpr(std::is_same_v>) + { return " (type: std::array)"; - } else if constexpr (std::is_same_v>) { + } + else if constexpr(std::is_same_v>) + { return " (type: std::array)"; - } else { + } + else + { return " (type: found but unknown)"; } } // ThreadBlockDescriptor diagnostics template -consteval auto diagnose_thread_block_descriptor() -> std::string { - if constexpr (!requires { T::thread_block; }) { +consteval auto diagnose_thread_block_descriptor() -> std::string +{ + if constexpr(!requires { T::thread_block; }) + { return " → T::thread_block member: [✗] (missing member)\n"; - } else { + } + else + { using TB = decltype(T::thread_block); std::string msg; - - if constexpr (requires(TB t) { t.block_size; }) { - using BlockSizeType = decltype(std::declval().block_size); + + if constexpr(requires(TB t) { t.block_size; }) + { + using BlockSizeType = decltype(std::declval().block_size); constexpr bool convertible = SizeType; - msg += " → thread_block.block_size: " + std::string(CHECK_MARK(convertible)) + + msg += " → thread_block.block_size: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → thread_block.block_size: [✗] (missing member)\n"; } - - if constexpr (requires(TB t) { t.tile_size.m; }) { - using TileMType = decltype(std::declval().tile_size.m); + + if constexpr(requires(TB t) { t.tile_size.m; }) + { + using TileMType = decltype(std::declval().tile_size.m); constexpr bool convertible = SizeType; - msg += " → thread_block.tile_size.m: " + std::string(CHECK_MARK(convertible)) + + msg += " → thread_block.tile_size.m: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → thread_block.tile_size.m: [✗] (missing member)\n"; } - - if constexpr (requires(TB t) { t.tile_size.n; }) { - using TileNType = decltype(std::declval().tile_size.n); + + if constexpr(requires(TB t) { t.tile_size.n; }) + { + using TileNType = decltype(std::declval().tile_size.n); constexpr bool convertible = SizeType; - msg += " → thread_block.tile_size.n: " + std::string(CHECK_MARK(convertible)) + + msg += " → thread_block.tile_size.n: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → thread_block.tile_size.n: [✗] (missing member)\n"; } - - if constexpr (requires(TB t) { t.tile_size.k; }) { - using TileKType = decltype(std::declval().tile_size.k); + + if constexpr(requires(TB t) { t.tile_size.k; }) + { + using TileKType = decltype(std::declval().tile_size.k); constexpr bool convertible = SizeType; - msg += " → thread_block.tile_size.k: " + std::string(CHECK_MARK(convertible)) + + msg += " → thread_block.tile_size.k: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → thread_block.tile_size.k: [✗] (missing member)\n"; } - + return msg; } } // GridwiseXdlGemmDescriptor diagnostics template -consteval auto diagnose_xdl_params() -> std::string { +consteval auto diagnose_xdl_params() -> std::string +{ std::string msg; - - if constexpr (requires(XdlParams t) { t.m_per_xdl; }) { - using MPerXdlType = decltype(std::declval().m_per_xdl); + + if constexpr(requires(XdlParams t) { t.m_per_xdl; }) + { + using MPerXdlType = decltype(std::declval().m_per_xdl); constexpr bool convertible = SizeType; - msg += " → xdl_params.m_per_xdl: " + std::string(CHECK_MARK(convertible)) + + msg += " → xdl_params.m_per_xdl: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → xdl_params.m_per_xdl: [✗] (missing member)\n"; } - - if constexpr (requires(XdlParams t) { t.n_per_xdl; }) { - using NPerXdlType = decltype(std::declval().n_per_xdl); + + if constexpr(requires(XdlParams t) { t.n_per_xdl; }) + { + using NPerXdlType = decltype(std::declval().n_per_xdl); constexpr bool convertible = SizeType; - msg += " → xdl_params.n_per_xdl: " + std::string(CHECK_MARK(convertible)) + + msg += " → xdl_params.n_per_xdl: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → xdl_params.n_per_xdl: [✗] (missing member)\n"; } - - if constexpr (requires(XdlParams t) { t.m_xdl_per_wave; }) { - using MXdlPerWaveType = decltype(std::declval().m_xdl_per_wave); + + if constexpr(requires(XdlParams t) { t.m_xdl_per_wave; }) + { + using MXdlPerWaveType = decltype(std::declval().m_xdl_per_wave); constexpr bool convertible = SizeType; - msg += " → xdl_params.m_xdl_per_wave: " + std::string(CHECK_MARK(convertible)) + + msg += " → xdl_params.m_xdl_per_wave: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → xdl_params.m_xdl_per_wave: [✗] (missing member)\n"; } - - if constexpr (requires(XdlParams t) { t.n_xdl_per_wave; }) { - using NXdlPerWaveType = decltype(std::declval().n_xdl_per_wave); + + if constexpr(requires(XdlParams t) { t.n_xdl_per_wave; }) + { + using NXdlPerWaveType = decltype(std::declval().n_xdl_per_wave); constexpr bool convertible = SizeType; - msg += " → xdl_params.n_xdl_per_wave: " + std::string(CHECK_MARK(convertible)) + + msg += " → xdl_params.n_xdl_per_wave: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += " → xdl_params.n_xdl_per_wave: [✗] (missing member)\n"; } - + return msg; } // BlockTransferDescriptor diagnostics template -consteval auto diagnose_block_transfer(const char* prefix) -> std::string { +consteval auto diagnose_block_transfer(const char* prefix) -> std::string +{ std::string msg; - - if constexpr (requires(BT t) { t.k0; }) { - using K0Type = decltype(std::declval().k0); + + if constexpr(requires(BT t) { t.k0; }) + { + using K0Type = decltype(std::declval().k0); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".k0: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".k0: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".k0: [✗] (missing member)\n"; } - - if constexpr (requires(BT t) { t.m_n; }) { - using MNType = decltype(std::declval().m_n); + + if constexpr(requires(BT t) { t.m_n; }) + { + using MNType = decltype(std::declval().m_n); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".m_n: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".m_n: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".m_n: [✗] (missing member)\n"; } - - if constexpr (requires(BT t) { t.k1; }) { - using K1Type = decltype(std::declval().k1); + + if constexpr(requires(BT t) { t.k1; }) + { + using K1Type = decltype(std::declval().k1); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".k1: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".k1: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".k1: [✗] (missing member)\n"; } - + return msg; } // BlockTransferDescriptor4D diagnostics (requires k_batch_size) template -consteval auto diagnose_block_transfer_4d(const char* prefix) -> std::string { +consteval auto diagnose_block_transfer_4d(const char* prefix) -> std::string +{ std::string msg; - - if constexpr (requires(BT t) { t.k0; }) { - using K0Type = decltype(std::declval().k0); + + if constexpr(requires(BT t) { t.k0; }) + { + using K0Type = decltype(std::declval().k0); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".k0: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".k0: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".k0: [✗] (missing member)\n"; } - - if constexpr (requires(BT t) { t.m_n; }) { - using MNType = decltype(std::declval().m_n); + + if constexpr(requires(BT t) { t.m_n; }) + { + using MNType = decltype(std::declval().m_n); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".m_n: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".m_n: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".m_n: [✗] (missing member)\n"; } - - if constexpr (requires(BT t) { t.k1; }) { - using K1Type = decltype(std::declval().k1); + + if constexpr(requires(BT t) { t.k1; }) + { + using K1Type = decltype(std::declval().k1); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".k1: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".k1: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".k1: [✗] (missing member)\n"; } - + // k_batch_size is required for Bwd descriptor - if constexpr (requires(BT t) { t.k_batch_size; }) { - using KBatchType = decltype(std::declval().k_batch_size); + if constexpr(requires(BT t) { t.k_batch_size; }) + { + using KBatchType = decltype(std::declval().k_batch_size); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".k_batch_size: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".k_batch_size: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".k_batch_size: [✗] (missing member)\n"; } - + return msg; } // LdsTransferDescriptor diagnostics template -consteval auto diagnose_lds_transfer(const char* prefix) -> std::string { +consteval auto diagnose_lds_transfer(const char* prefix) -> std::string +{ std::string msg; - - if constexpr (requires(LT t) { t.src_vector_dim; }) { - using SrcVectorDimType = decltype(std::declval().src_vector_dim); + + if constexpr(requires(LT t) { t.src_vector_dim; }) + { + using SrcVectorDimType = decltype(std::declval().src_vector_dim); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".src_vector_dim: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".src_vector_dim: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".src_vector_dim: [✗] (missing member)\n"; } - - if constexpr (requires(LT t) { t.src_scalar_per_vector; }) { - using SrcScalarType = decltype(std::declval().src_scalar_per_vector); + + if constexpr(requires(LT t) { t.src_scalar_per_vector; }) + { + using SrcScalarType = decltype(std::declval().src_scalar_per_vector); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".src_scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".src_scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { - msg += std::string(" → ") + prefix + ".src_scalar_per_vector: [✗] (missing member)\n"; } - - if constexpr (requires(LT t) { t.lds_dst_scalar_per_vector; }) { - using LdsDstScalarType = decltype(std::declval().lds_dst_scalar_per_vector); + else + { + msg += + std::string(" → ") + prefix + ".src_scalar_per_vector: [✗] (missing member)\n"; + } + + if constexpr(requires(LT t) { t.lds_dst_scalar_per_vector; }) + { + using LdsDstScalarType = decltype(std::declval().lds_dst_scalar_per_vector); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".lds_dst_scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".lds_dst_scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { - msg += std::string(" → ") + prefix + ".lds_dst_scalar_per_vector: [✗] (missing member)\n"; } - - if constexpr (requires(LT t) { t.is_direct_load; }) { - using IsDirectLoadType = decltype(std::declval().is_direct_load); + else + { + msg += std::string(" → ") + prefix + + ".lds_dst_scalar_per_vector: [✗] (missing member)\n"; + } + + if constexpr(requires(LT t) { t.is_direct_load; }) + { + using IsDirectLoadType = decltype(std::declval().is_direct_load); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".is_direct_load: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".is_direct_load: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".is_direct_load: [✗] (missing member)\n"; } - - if constexpr (requires(LT t) { t.lds_padding; }) { - using LdsPaddingType = decltype(std::declval().lds_padding); + + if constexpr(requires(LT t) { t.lds_padding; }) + { + using LdsPaddingType = decltype(std::declval().lds_padding); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".lds_padding: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".lds_padding: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".lds_padding: [✗] (missing member)\n"; } - + return msg; } // ThreadClusterDescriptor diagnostics template -consteval auto diagnose_thread_cluster(const char* prefix) -> std::string { +consteval auto diagnose_thread_cluster(const char* prefix) -> std::string +{ std::string msg; - - if constexpr (requires(TC t) { t.m_block; }) { - using MBlockType = decltype(std::declval().m_block); + + if constexpr(requires(TC t) { t.m_block; }) + { + using MBlockType = decltype(std::declval().m_block); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".m_block: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".m_block: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".m_block: [✗] (missing member)\n"; } - - if constexpr (requires(TC t) { t.m_wave_per_xdl; }) { - using MWaveType = decltype(std::declval().m_wave_per_xdl); + + if constexpr(requires(TC t) { t.m_wave_per_xdl; }) + { + using MWaveType = decltype(std::declval().m_wave_per_xdl); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".m_wave_per_xdl: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".m_wave_per_xdl: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".m_wave_per_xdl: [✗] (missing member)\n"; } - - if constexpr (requires(TC t) { t.n_block; }) { - using NBlockType = decltype(std::declval().n_block); + + if constexpr(requires(TC t) { t.n_block; }) + { + using NBlockType = decltype(std::declval().n_block); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".n_block: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".n_block: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".n_block: [✗] (missing member)\n"; } - - if constexpr (requires(TC t) { t.n_wave_per_xdl; }) { - using NWaveType = decltype(std::declval().n_wave_per_xdl); + + if constexpr(requires(TC t) { t.n_wave_per_xdl; }) + { + using NWaveType = decltype(std::declval().n_wave_per_xdl); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".n_wave_per_xdl: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".n_wave_per_xdl: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".n_wave_per_xdl: [✗] (missing member)\n"; } - + return msg; } // AccessOrderDescriptor diagnostics template -consteval auto diagnose_access_order(const char* prefix) -> std::string { +consteval auto diagnose_access_order(const char* prefix) -> std::string +{ std::string msg; - - if constexpr (requires(AO t) { t.order; }) { - using OrderType = decltype(std::declval().order); + + if constexpr(requires(AO t) { t.order; }) + { + using OrderType = decltype(std::declval().order); constexpr bool convertible_3 = std::convertible_to>; constexpr bool convertible_4 = std::convertible_to>; - constexpr bool convertible = convertible_3 || convertible_4; - msg += std::string(" → ") + prefix + ".order: " + std::string(CHECK_MARK(convertible)) + + constexpr bool convertible = convertible_3 || convertible_4; + msg += std::string(" → ") + prefix + + ".order: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".order: [✗] (missing member)\n"; } - + return msg; } // EpilogueDescriptor diagnostics template -consteval auto diagnose_epilogue(const char* prefix) -> std::string { +consteval auto diagnose_epilogue(const char* prefix) -> std::string +{ std::string msg; - - if constexpr (requires(E t) { t.m_xdl_per_wave_per_shuffle; }) { - using MXdlType = decltype(std::declval().m_xdl_per_wave_per_shuffle); + + if constexpr(requires(E t) { t.m_xdl_per_wave_per_shuffle; }) + { + using MXdlType = decltype(std::declval().m_xdl_per_wave_per_shuffle); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".m_xdl_per_wave_per_shuffle: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".m_xdl_per_wave_per_shuffle: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { - msg += std::string(" → ") + prefix + ".m_xdl_per_wave_per_shuffle: [✗] (missing member)\n"; } - - if constexpr (requires(E t) { t.n_per_wave_per_shuffle; }) { - using NPerWaveType = decltype(std::declval().n_per_wave_per_shuffle); + else + { + msg += std::string(" → ") + prefix + + ".m_xdl_per_wave_per_shuffle: [✗] (missing member)\n"; + } + + if constexpr(requires(E t) { t.n_per_wave_per_shuffle; }) + { + using NPerWaveType = decltype(std::declval().n_per_wave_per_shuffle); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".n_per_wave_per_shuffle: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".n_per_wave_per_shuffle: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { - msg += std::string(" → ") + prefix + ".n_per_wave_per_shuffle: [✗] (missing member)\n"; } - - if constexpr (requires(E t) { t.scalar_per_vector; }) { - using ScalarType = decltype(std::declval().scalar_per_vector); + else + { + msg += std::string(" → ") + prefix + + ".n_per_wave_per_shuffle: [✗] (missing member)\n"; + } + + if constexpr(requires(E t) { t.scalar_per_vector; }) + { + using ScalarType = decltype(std::declval().scalar_per_vector); constexpr bool convertible = std::convertible_to; - msg += std::string(" → ") + prefix + ".scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + + msg += std::string(" → ") + prefix + + ".scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(get_type_info())) + "\n"; - } else { + } + else + { msg += std::string(" → ") + prefix + ".scalar_per_vector: [✗] (missing member)\n"; } - + return msg; } @@ -384,398 +531,544 @@ consteval auto diagnose_epilogue(const char* prefix) -> std::string { // Detailed diagnostic functions for high-level concepts template -consteval auto detailed_diagnostic_ConvAlgorithmDescriptor() -> std::string { +consteval auto detailed_diagnostic_ConvAlgorithmDescriptor() -> std::string +{ return ""; // Base concept, no sub-requirements to check } template -consteval auto detailed_diagnostic_SpecifiesThreadBlock() -> std::string { - if constexpr (!requires { { T::thread_block } -> ThreadBlockDescriptor; }) { +consteval auto detailed_diagnostic_SpecifiesThreadBlock() -> std::string +{ + if constexpr(!requires { + { T::thread_block } -> ThreadBlockDescriptor; + }) + { return " → T::thread_block member: [✗] (missing or wrong type)\n"; - } else { - return " → T::thread_block member: [✓]\n" + + } + else + { + return " → T::thread_block member: [✓]\n" + detail::diagnose_thread_block_descriptor(); } } template -consteval auto detailed_diagnostic_SpecifiesGridwiseFwdXdlGemm() -> std::string { +consteval auto detailed_diagnostic_SpecifiesGridwiseFwdXdlGemm() -> std::string +{ std::string msg; - - if constexpr (!requires(T t) { { t.gridwise_gemm } -> GridwiseFwdXdlGemmDescriptor; }) { + + if constexpr(!requires(T t) { + { t.gridwise_gemm } -> GridwiseFwdXdlGemmDescriptor; + }) + { return " → T::gridwise_gemm member: [✗] (missing or wrong type)\n"; } - + msg += " → T::gridwise_gemm member: [✓]\n"; using GG = decltype(T::gridwise_gemm); - - if constexpr (requires(GG t) { t.ak1; }) { - using AK1Type = decltype(std::declval().ak1); + + if constexpr(requires(GG t) { t.ak1; }) + { + using AK1Type = decltype(std::declval().ak1); constexpr bool convertible = std::convertible_to; - msg += " → gridwise_gemm.ak1: " + std::string(CHECK_MARK(convertible)) + + msg += " → gridwise_gemm.ak1: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → gridwise_gemm.ak1: [✗] (missing member)\n"; } - - if constexpr (requires(GG t) { t.bk1; }) { - using BK1Type = decltype(std::declval().bk1); + + if constexpr(requires(GG t) { t.bk1; }) + { + using BK1Type = decltype(std::declval().bk1); constexpr bool convertible = std::convertible_to; - msg += " → gridwise_gemm.bk1: " + std::string(CHECK_MARK(convertible)) + + msg += " → gridwise_gemm.bk1: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → gridwise_gemm.bk1: [✗] (missing member)\n"; } - - if constexpr (requires(GG t) { t.xdl_params; }) { + + if constexpr(requires(GG t) { t.xdl_params; }) + { msg += " → gridwise_gemm.xdl_params member: [✓]\n"; msg += detail::diagnose_xdl_params().xdl_params)>(); - } else { + } + else + { msg += " → gridwise_gemm.xdl_params: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesGridwiseBwdXdlGemm() -> std::string { +consteval auto detailed_diagnostic_SpecifiesGridwiseBwdXdlGemm() -> std::string +{ std::string msg; - - if constexpr (!requires(T t) { { t.gridwise_gemm } -> GridwiseBwdXdlGemmDescriptor; }) { + + if constexpr(!requires(T t) { + { t.gridwise_gemm } -> GridwiseBwdXdlGemmDescriptor; + }) + { return " → T::gridwise_gemm member: [✗] (missing or wrong type)\n"; } - + msg += " → T::gridwise_gemm member: [✓]\n"; using GG = decltype(T::gridwise_gemm); - - if constexpr (requires(GG t) { t.k1; }) { - using K1Type = decltype(std::declval().k1); + + if constexpr(requires(GG t) { t.k1; }) + { + using K1Type = decltype(std::declval().k1); constexpr bool convertible = std::convertible_to; - msg += " → gridwise_gemm.k1: " + std::string(CHECK_MARK(convertible)) + + msg += " → gridwise_gemm.k1: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → gridwise_gemm.k1: [✗] (missing member)\n"; } - - if constexpr (requires(GG t) { t.xdl_params; }) { + + if constexpr(requires(GG t) { t.xdl_params; }) + { msg += " → gridwise_gemm.xdl_params member: [✓]\n"; msg += detail::diagnose_xdl_params().xdl_params)>(); - } else { + } + else + { msg += " → gridwise_gemm.xdl_params: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesBlockTransfer() -> std::string { +consteval auto detailed_diagnostic_SpecifiesBlockTransfer() -> std::string +{ std::string msg; - + constexpr bool has_transfer = requires { T::transfer; }; msg += " → T::transfer member: " + std::string(CHECK_MARK(has_transfer)) + "\n"; - - if constexpr (!has_transfer) { + + if constexpr(!has_transfer) + { return msg; } - - constexpr bool has_a = requires { { T::transfer.a.block_transfer } -> BlockTransferDescriptor; }; + + constexpr bool has_a = requires { + { T::transfer.a.block_transfer } -> BlockTransferDescriptor; + }; msg += " → T::transfer.a: " + std::string(CHECK_MARK(has_a)) + "\n"; - if constexpr (!has_a) { + if constexpr(!has_a) + { msg += " → T::transfer.a.block_transfer: [✗] (missing or wrong type)\n"; } - - constexpr bool has_b = requires { { T::transfer.b.block_transfer } -> BlockTransferDescriptor; }; + + constexpr bool has_b = requires { + { T::transfer.b.block_transfer } -> BlockTransferDescriptor; + }; msg += " → T::transfer.b: " + std::string(CHECK_MARK(has_b)) + "\n"; - if constexpr (!has_b) { + if constexpr(!has_b) + { msg += " → T::transfer.b.block_transfer: [✗] (missing or wrong type)\n"; } - - constexpr bool has_c = requires { { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor; }; + + constexpr bool has_c = requires { + { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor; + }; msg += " → T::transfer.c: " + std::string(CHECK_MARK(has_c)) + "\n"; - if constexpr (!has_c) { + if constexpr(!has_c) + { msg += " → T::transfer.c.thread_cluster_dims: [✗] (missing or wrong type)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesBlockTransfer4D() -> std::string { +consteval auto detailed_diagnostic_SpecifiesBlockTransfer4D() -> std::string +{ std::string msg; - + constexpr bool has_transfer = requires { T::transfer; }; msg += " → T::transfer member: " + std::string(CHECK_MARK(has_transfer)) + "\n"; - - if constexpr (!has_transfer) { + + if constexpr(!has_transfer) + { return msg; } - - constexpr bool has_a = requires { { T::transfer.a.block_transfer } -> BlockTransferDescriptor4D; }; + + constexpr bool has_a = requires { + { T::transfer.a.block_transfer } -> BlockTransferDescriptor4D; + }; msg += " → T::transfer.a: " + std::string(CHECK_MARK(has_a)) + "\n"; - if constexpr (!has_a) { + if constexpr(!has_a) + { msg += " → T::transfer.a.block_transfer: [✗] (missing or wrong type)\n"; - } else { - msg += detail::diagnose_block_transfer_4d("transfer.a.block_transfer"); } - - constexpr bool has_b = requires { { T::transfer.b.block_transfer } -> BlockTransferDescriptor4D; }; + else + { + msg += detail::diagnose_block_transfer_4d( + "transfer.a.block_transfer"); + } + + constexpr bool has_b = requires { + { T::transfer.b.block_transfer } -> BlockTransferDescriptor4D; + }; msg += " → T::transfer.b: " + std::string(CHECK_MARK(has_b)) + "\n"; - if constexpr (!has_b) { + if constexpr(!has_b) + { msg += " → T::transfer.b.block_transfer: [✗] (missing or wrong type)\n"; - } else { - msg += detail::diagnose_block_transfer_4d("transfer.b.block_transfer"); } - - constexpr bool has_c = requires { { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor; }; + else + { + msg += detail::diagnose_block_transfer_4d( + "transfer.b.block_transfer"); + } + + constexpr bool has_c = requires { + { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor; + }; msg += " → T::transfer.c: " + std::string(CHECK_MARK(has_c)) + "\n"; - if constexpr (!has_c) { + if constexpr(!has_c) + { msg += " → T::transfer.c.thread_cluster_dims: [✗] (missing or wrong type)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesThreadClusterAccessOrder() -> std::string { +consteval auto detailed_diagnostic_SpecifiesThreadClusterAccessOrder() -> std::string +{ std::string msg; - + constexpr bool has_transfer = requires { T::transfer; }; - if constexpr (!has_transfer) { + if constexpr(!has_transfer) + { return " → T::transfer member: [✗] (missing member)\n"; } - + constexpr bool has_a = requires { T::transfer.a; }; constexpr bool has_b = requires { T::transfer.b; }; - - if constexpr (has_a && requires { T::transfer.a.block_transfer_access_order; }) { - msg += detail::diagnose_access_order("transfer.a.block_transfer_access_order"); - } else if constexpr (has_a) { + + if constexpr(has_a && requires { T::transfer.a.block_transfer_access_order; }) + { + msg += + detail::diagnose_access_order( + "transfer.a.block_transfer_access_order"); + } + else if constexpr(has_a) + { msg += " → T::transfer.a.block_transfer_access_order: [✗] (missing member)\n"; } - - if constexpr (has_b && requires { T::transfer.b.block_transfer_access_order; }) { - msg += detail::diagnose_access_order("transfer.b.block_transfer_access_order"); - } else if constexpr (has_b) { + + if constexpr(has_b && requires { T::transfer.b.block_transfer_access_order; }) + { + msg += + detail::diagnose_access_order( + "transfer.b.block_transfer_access_order"); + } + else if constexpr(has_b) + { msg += " → T::transfer.b.block_transfer_access_order: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesSourceAccessOrder() -> std::string { +consteval auto detailed_diagnostic_SpecifiesSourceAccessOrder() -> std::string +{ std::string msg; - + constexpr bool has_transfer = requires { T::transfer; }; - if constexpr (!has_transfer) { + if constexpr(!has_transfer) + { return " → T::transfer member: [✗] (missing member)\n"; } - + constexpr bool has_a = requires { T::transfer.a; }; constexpr bool has_b = requires { T::transfer.b; }; - - if constexpr (has_a && requires { T::transfer.a.src_access_order; }) { - msg += detail::diagnose_access_order("transfer.a.src_access_order"); - } else if constexpr (has_a) { + + if constexpr(has_a && requires { T::transfer.a.src_access_order; }) + { + msg += detail::diagnose_access_order( + "transfer.a.src_access_order"); + } + else if constexpr(has_a) + { msg += " → T::transfer.a.src_access_order: [✗] (missing member)\n"; } - - if constexpr (has_b && requires { T::transfer.b.src_access_order; }) { - msg += detail::diagnose_access_order("transfer.b.src_access_order"); - } else if constexpr (has_b) { + + if constexpr(has_b && requires { T::transfer.b.src_access_order; }) + { + msg += detail::diagnose_access_order( + "transfer.b.src_access_order"); + } + else if constexpr(has_b) + { msg += " → T::transfer.b.src_access_order: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesBlockGemm() -> std::string { +consteval auto detailed_diagnostic_SpecifiesBlockGemm() -> std::string +{ std::string msg; - - if constexpr (!requires { { T::block_gemm_pipeline } -> BlockGemmPipelineDescriptor; }) { + + if constexpr(!requires { + { T::block_gemm_pipeline } -> BlockGemmPipelineDescriptor; + }) + { return " → T::block_gemm_pipeline: [✗] (missing or wrong type)\n"; } - + msg += " → T::block_gemm_pipeline member: [✓]\n"; - - if constexpr (requires { T::block_gemm_pipeline.pipeline_version; }) { - using PipelineType = decltype(T::block_gemm_pipeline.pipeline_version); + + if constexpr(requires { T::block_gemm_pipeline.pipeline_version; }) + { + using PipelineType = decltype(T::block_gemm_pipeline.pipeline_version); constexpr bool convertible = std::convertible_to; - msg += " → block_gemm_pipeline.pipeline_version: " + std::string(CHECK_MARK(convertible)) + + msg += " → block_gemm_pipeline.pipeline_version: " + + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → block_gemm_pipeline.pipeline_version: [✗] (missing member)\n"; } - - if constexpr (requires { T::block_gemm_pipeline.scheduler; }) { - using SchedulerType = decltype(T::block_gemm_pipeline.scheduler); + + if constexpr(requires { T::block_gemm_pipeline.scheduler; }) + { + using SchedulerType = decltype(T::block_gemm_pipeline.scheduler); constexpr bool convertible = std::convertible_to; - msg += " → block_gemm_pipeline.scheduler: " + std::string(CHECK_MARK(convertible)) + + msg += " → block_gemm_pipeline.scheduler: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → block_gemm_pipeline.scheduler: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesFwdConvSpecialization() -> std::string { - if constexpr (requires { T::fwd_specialization; }) { - using FwdSpecType = decltype(T::fwd_specialization); +consteval auto detailed_diagnostic_SpecifiesFwdConvSpecialization() -> std::string +{ + if constexpr(requires { T::fwd_specialization; }) + { + using FwdSpecType = decltype(T::fwd_specialization); constexpr bool convertible = std::convertible_to; - return " → T::fwd_specialization: " + std::string(CHECK_MARK(convertible)) + + return " → T::fwd_specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::fwd_specialization: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesBwdWeightConvSpecialization() -> std::string { - if constexpr (requires { T::bwd_weight_specialization; }) { - using BwdSpecType = decltype(T::bwd_weight_specialization); +consteval auto detailed_diagnostic_SpecifiesBwdWeightConvSpecialization() -> std::string +{ + if constexpr(requires { T::bwd_weight_specialization; }) + { + using BwdSpecType = decltype(T::bwd_weight_specialization); constexpr bool convertible = std::convertible_to; - return " → T::bwd_weight_specialization: " + std::string(CHECK_MARK(convertible)) + + return " → T::bwd_weight_specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::bwd_weight_specialization: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesGemmSpecialization() -> std::string { - if constexpr (requires { T::gemm_specialization; }) { - using GemmSpecType = decltype(T::gemm_specialization); +consteval auto detailed_diagnostic_SpecifiesGemmSpecialization() -> std::string +{ + if constexpr(requires { T::gemm_specialization; }) + { + using GemmSpecType = decltype(T::gemm_specialization); constexpr bool convertible = std::convertible_to; - return " → T::gemm_specialization: " + std::string(CHECK_MARK(convertible)) + + return " → T::gemm_specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::gemm_specialization: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesNumPrefetchStages() -> std::string { - if constexpr (requires { T::num_gemm_k_prefetch_stages; }) { - using NumPrefetchType = decltype(T::num_gemm_k_prefetch_stages); +consteval auto detailed_diagnostic_SpecifiesNumPrefetchStages() -> std::string +{ + if constexpr(requires { T::num_gemm_k_prefetch_stages; }) + { + using NumPrefetchType = decltype(T::num_gemm_k_prefetch_stages); constexpr bool convertible = std::convertible_to; - return " → T::num_gemm_k_prefetch_stages: " + std::string(CHECK_MARK(convertible)) + + return " → T::num_gemm_k_prefetch_stages: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::num_gemm_k_prefetch_stages: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesNumGroupsToMerge() -> std::string { - if constexpr (requires { T::num_conv_groups_to_merge; }) { - using NumGroupsType = decltype(T::num_conv_groups_to_merge); +consteval auto detailed_diagnostic_SpecifiesNumGroupsToMerge() -> std::string +{ + if constexpr(requires { T::num_conv_groups_to_merge; }) + { + using NumGroupsType = decltype(T::num_conv_groups_to_merge); constexpr bool convertible = std::convertible_to; - return " → T::num_conv_groups_to_merge: " + std::string(CHECK_MARK(convertible)) + + return " → T::num_conv_groups_to_merge: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::num_conv_groups_to_merge: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesLoopScheduler() -> std::string { - if constexpr (requires { T::loop_scheduler; }) { - using LoopSchedulerType = decltype(T::loop_scheduler); +consteval auto detailed_diagnostic_SpecifiesLoopScheduler() -> std::string +{ + if constexpr(requires { T::loop_scheduler; }) + { + using LoopSchedulerType = decltype(T::loop_scheduler); constexpr bool convertible = std::convertible_to; - return " → T::loop_scheduler: " + std::string(CHECK_MARK(convertible)) + + return " → T::loop_scheduler: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::loop_scheduler: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesLargeTensorSupport() -> std::string { +consteval auto detailed_diagnostic_SpecifiesLargeTensorSupport() -> std::string +{ std::string msg; - if constexpr (requires { T::specialization; }) { - using SpecType = decltype(T::specialization); + if constexpr(requires { T::specialization; }) + { + using SpecType = decltype(T::specialization); constexpr bool convertible = std::convertible_to; - msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + + msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - - if constexpr (convertible) { - constexpr bool is_large_tensor = (T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR); - msg += " → specialization == LARGE_TENSOR: " + std::string(CHECK_MARK(is_large_tensor)) + "\n"; + + if constexpr(convertible) + { + constexpr bool is_large_tensor = + (T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR); + msg += " → specialization == LARGE_TENSOR: " + + std::string(CHECK_MARK(is_large_tensor)) + "\n"; } - } else { + } + else + { msg += " → T::specialization: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesReferenceAlgorithm() -> std::string { +consteval auto detailed_diagnostic_SpecifiesReferenceAlgorithm() -> std::string +{ std::string msg; - if constexpr (requires { T::specialization; }) { - using SpecType = decltype(T::specialization); + if constexpr(requires { T::specialization; }) + { + using SpecType = decltype(T::specialization); constexpr bool convertible = std::convertible_to; - msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + + msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - - if constexpr (convertible) { - constexpr bool is_reference = (T::specialization == ConvAlgorithmSpecialization::REFERENCE); - msg += " → specialization == REFERENCE: " + std::string(CHECK_MARK(is_reference)) + "\n"; + + if constexpr(convertible) + { + constexpr bool is_reference = + (T::specialization == ConvAlgorithmSpecialization::REFERENCE); + msg += " → specialization == REFERENCE: " + std::string(CHECK_MARK(is_reference)) + + "\n"; } - } else { + } + else + { msg += " → T::specialization: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesTwoStageSupport() -> std::string { +consteval auto detailed_diagnostic_SpecifiesTwoStageSupport() -> std::string +{ std::string msg; - if constexpr (requires { T::specialization; }) { - using SpecType = decltype(T::specialization); + if constexpr(requires { T::specialization; }) + { + using SpecType = decltype(T::specialization); constexpr bool convertible = std::convertible_to; - msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + + msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - - if constexpr (convertible) { - constexpr bool is_two_stage = (T::specialization == ConvAlgorithmSpecialization::TWO_STAGE); - msg += " → specialization == TWO_STAGE: " + std::string(CHECK_MARK(is_two_stage)) + "\n"; + + if constexpr(convertible) + { + constexpr bool is_two_stage = + (T::specialization == ConvAlgorithmSpecialization::TWO_STAGE); + msg += " → specialization == TWO_STAGE: " + std::string(CHECK_MARK(is_two_stage)) + + "\n"; } - } else { + } + else + { msg += " → T::specialization: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesMultipleDSupport() -> std::string { +consteval auto detailed_diagnostic_SpecifiesMultipleDSupport() -> std::string +{ std::string msg; - if constexpr (requires { T::specialization; }) { - using SpecType = decltype(T::specialization); + if constexpr(requires { T::specialization; }) + { + using SpecType = decltype(T::specialization); constexpr bool convertible = std::convertible_to; - msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + + msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - - if constexpr (convertible) { - constexpr bool is_multiple_d = (T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D); - msg += " → specialization == MULTIPLE_D: " + std::string(CHECK_MARK(is_multiple_d)) + "\n"; + + if constexpr(convertible) + { + constexpr bool is_multiple_d = + (T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D); + msg += + " → specialization == MULTIPLE_D: " + std::string(CHECK_MARK(is_multiple_d)) + + "\n"; } - } else { + } + else + { msg += " → T::specialization: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesGenericInstance() -> std::string { +consteval auto detailed_diagnostic_SpecifiesGenericInstance() -> std::string +{ std::string msg; - if constexpr (requires { T::specialization; }) { + if constexpr(requires { T::specialization; }) + { msg += " → T::specialization: [✗] (member should NOT exist for generic instance)\n"; msg += " → This concept requires the absence of the specialization member\n"; } @@ -783,423 +1076,692 @@ consteval auto detailed_diagnostic_SpecifiesGenericInstance() -> std::string { } template -consteval auto detailed_diagnostic_SpecifiesTransposeTransfer() -> std::string { +consteval auto detailed_diagnostic_SpecifiesTransposeTransfer() -> std::string +{ std::string msg; - - if constexpr (requires { T::max_transpose_transfer_src_scalar_per_vector; }) { - using SrcType = decltype(T::max_transpose_transfer_src_scalar_per_vector); + + if constexpr(requires { T::max_transpose_transfer_src_scalar_per_vector; }) + { + using SrcType = decltype(T::max_transpose_transfer_src_scalar_per_vector); constexpr bool convertible = std::convertible_to; - msg += " → T::max_transpose_transfer_src_scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + + msg += " → T::max_transpose_transfer_src_scalar_per_vector: " + + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → T::max_transpose_transfer_src_scalar_per_vector: [✗] (missing member)\n"; } - - if constexpr (requires { T::max_transpose_transfer_dst_scalar_per_vector; }) { - using DstType = decltype(T::max_transpose_transfer_dst_scalar_per_vector); + + if constexpr(requires { T::max_transpose_transfer_dst_scalar_per_vector; }) + { + using DstType = decltype(T::max_transpose_transfer_dst_scalar_per_vector); constexpr bool convertible = std::convertible_to; - msg += " → T::max_transpose_transfer_dst_scalar_per_vector: " + std::string(CHECK_MARK(convertible)) + + msg += " → T::max_transpose_transfer_dst_scalar_per_vector: " + + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { msg += " → T::max_transpose_transfer_dst_scalar_per_vector: [✗] (missing member)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_TransposeTransferWellDefinedIfProvided() -> std::string { +consteval auto detailed_diagnostic_TransposeTransferWellDefinedIfProvided() -> std::string +{ std::string msg; - + constexpr bool has_src = requires { T::max_transpose_transfer_src_scalar_per_vector; }; constexpr bool has_dst = requires { T::max_transpose_transfer_dst_scalar_per_vector; }; constexpr bool has_transpose_transfer = has_src || has_dst; - - if constexpr (!has_transpose_transfer) { + + if constexpr(!has_transpose_transfer) + { msg += " → Transpose transfer fields not provided: [✓] (optional, not required)\n"; - } else { + } + else + { msg += " → Transpose transfer fields provided, checking if well-defined:\n"; - - if constexpr (has_src) { + + if constexpr(has_src) + { using SrcType = decltype(T::max_transpose_transfer_src_scalar_per_vector); constexpr bool src_convertible = std::convertible_to; - msg += " → T::max_transpose_transfer_src_scalar_per_vector: " + - std::string(CHECK_MARK(src_convertible)) + + msg += " → T::max_transpose_transfer_src_scalar_per_vector: " + + std::string(CHECK_MARK(src_convertible)) + (src_convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { - msg += " → T::max_transpose_transfer_src_scalar_per_vector: [✗] (missing, but dst is provided)\n"; } - - if constexpr (has_dst) { + else + { + msg += " → T::max_transpose_transfer_src_scalar_per_vector: [✗] (missing, but " + "dst is provided)\n"; + } + + if constexpr(has_dst) + { using DstType = decltype(T::max_transpose_transfer_dst_scalar_per_vector); constexpr bool dst_convertible = std::convertible_to; - msg += " → T::max_transpose_transfer_dst_scalar_per_vector: " + - std::string(CHECK_MARK(dst_convertible)) + + msg += " → T::max_transpose_transfer_dst_scalar_per_vector: " + + std::string(CHECK_MARK(dst_convertible)) + (dst_convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { - msg += " → T::max_transpose_transfer_dst_scalar_per_vector: [✗] (missing, but src is provided)\n"; + } + else + { + msg += " → T::max_transpose_transfer_dst_scalar_per_vector: [✗] (missing, but " + "src is provided)\n"; } } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesGemmBatchOptions() -> std::string { - if constexpr (requires { T::num_conv_groups_to_merge; }) { - using NumGroupsType = decltype(T::num_conv_groups_to_merge); +consteval auto detailed_diagnostic_SpecifiesGemmBatchOptions() -> std::string +{ + if constexpr(requires { T::num_conv_groups_to_merge; }) + { + using NumGroupsType = decltype(T::num_conv_groups_to_merge); constexpr bool convertible = std::convertible_to; - return " → T::num_conv_groups_to_merge: " + std::string(CHECK_MARK(convertible)) + + return " → T::num_conv_groups_to_merge: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::num_conv_groups_to_merge: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesGridwiseWmmaGemm() -> std::string { +consteval auto detailed_diagnostic_SpecifiesGridwiseWmmaGemm() -> std::string +{ std::string msg; - constexpr bool has_gridwise_gemm = requires(T t) { { t.gridwise_gemm } -> GridwiseWmmaGemmDescriptor; }; + constexpr bool has_gridwise_gemm = requires(T t) { + { t.gridwise_gemm } -> GridwiseWmmaGemmDescriptor; + }; msg += " → T::gridwise_gemm member: " + std::string(CHECK_MARK(has_gridwise_gemm)) + "\n"; - - if constexpr (!has_gridwise_gemm) { + + if constexpr(!has_gridwise_gemm) + { return msg; } - - using GG = decltype(T::gridwise_gemm); - constexpr bool has_k1 = requires(GG t) { { t.k1 } -> std::convertible_to; }; - constexpr bool has_m_per_wmma = requires(GG t) { { t.m_per_wmma } -> std::convertible_to; }; - constexpr bool has_n_per_wmma = requires(GG t) { { t.n_per_wmma } -> std::convertible_to; }; - constexpr bool has_m_wmma_per_wave = requires(GG t) { { t.m_wmma_per_wave } -> std::convertible_to; }; - constexpr bool has_n_wmma_per_wave = requires(GG t) { { t.n_wmma_per_wave } -> std::convertible_to; }; - - msg += " → gridwise_gemm.k1: " + std::string(CHECK_MARK(has_k1)) + (has_k1 ? "\n" : " (missing or wrong type)\n"); - msg += " → gridwise_gemm.m_per_wmma: " + std::string(CHECK_MARK(has_m_per_wmma)) + (has_m_per_wmma ? "\n" : " (missing or wrong type)\n"); - msg += " → gridwise_gemm.n_per_wmma: " + std::string(CHECK_MARK(has_n_per_wmma)) + (has_n_per_wmma ? "\n" : " (missing or wrong type)\n"); - msg += " → gridwise_gemm.m_wmma_per_wave: " + std::string(CHECK_MARK(has_m_wmma_per_wave)) + (has_m_wmma_per_wave ? "\n" : " (missing or wrong type)\n"); - msg += " → gridwise_gemm.n_wmma_per_wave: " + std::string(CHECK_MARK(has_n_wmma_per_wave)) + (has_n_wmma_per_wave ? "\n" : " (missing or wrong type)\n"); - + + using GG = decltype(T::gridwise_gemm); + constexpr bool has_k1 = requires(GG t) { + { t.k1 } -> std::convertible_to; + }; + constexpr bool has_m_per_wmma = requires(GG t) { + { t.m_per_wmma } -> std::convertible_to; + }; + constexpr bool has_n_per_wmma = requires(GG t) { + { t.n_per_wmma } -> std::convertible_to; + }; + constexpr bool has_m_wmma_per_wave = requires(GG t) { + { t.m_wmma_per_wave } -> std::convertible_to; + }; + constexpr bool has_n_wmma_per_wave = requires(GG t) { + { t.n_wmma_per_wave } -> std::convertible_to; + }; + + msg += " → gridwise_gemm.k1: " + std::string(CHECK_MARK(has_k1)) + + (has_k1 ? "\n" : " (missing or wrong type)\n"); + msg += " → gridwise_gemm.m_per_wmma: " + std::string(CHECK_MARK(has_m_per_wmma)) + + (has_m_per_wmma ? "\n" : " (missing or wrong type)\n"); + msg += " → gridwise_gemm.n_per_wmma: " + std::string(CHECK_MARK(has_n_per_wmma)) + + (has_n_per_wmma ? "\n" : " (missing or wrong type)\n"); + msg += + " → gridwise_gemm.m_wmma_per_wave: " + std::string(CHECK_MARK(has_m_wmma_per_wave)) + + (has_m_wmma_per_wave ? "\n" : " (missing or wrong type)\n"); + msg += + " → gridwise_gemm.n_wmma_per_wave: " + std::string(CHECK_MARK(has_n_wmma_per_wave)) + + (has_n_wmma_per_wave ? "\n" : " (missing or wrong type)\n"); + return msg; } // Tile-specific diagnostics template -consteval auto detailed_diagnostic_SpecifiesTileThreadBlock() -> std::string { - if constexpr (!requires { { T::thread_block } -> TileThreadBlockDescriptor; }) { +consteval auto detailed_diagnostic_SpecifiesTileThreadBlock() -> std::string +{ + if constexpr(!requires { + { T::thread_block } -> TileThreadBlockDescriptor; + }) + { return " → T::thread_block member: [✗] (missing or wrong type)\n"; - } else { - using TB = decltype(T::thread_block); + } + else + { + using TB = decltype(T::thread_block); std::string msg = " → T::thread_block member: [✓]\n"; - - constexpr bool has_tile_m = requires(TB t) { { t.tile_size.m } -> std::convertible_to; }; - constexpr bool has_tile_n = requires(TB t) { { t.tile_size.n } -> std::convertible_to; }; - constexpr bool has_tile_k = requires(TB t) { { t.tile_size.k } -> std::convertible_to; }; - - msg += " → thread_block.tile_size.m: " + std::string(CHECK_MARK(has_tile_m)) + (has_tile_m ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_block.tile_size.n: " + std::string(CHECK_MARK(has_tile_n)) + (has_tile_n ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_block.tile_size.k: " + std::string(CHECK_MARK(has_tile_k)) + (has_tile_k ? "\n" : " (missing or wrong type)\n"); - + + constexpr bool has_tile_m = requires(TB t) { + { t.tile_size.m } -> std::convertible_to; + }; + constexpr bool has_tile_n = requires(TB t) { + { t.tile_size.n } -> std::convertible_to; + }; + constexpr bool has_tile_k = requires(TB t) { + { t.tile_size.k } -> std::convertible_to; + }; + + msg += " → thread_block.tile_size.m: " + std::string(CHECK_MARK(has_tile_m)) + + (has_tile_m ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_block.tile_size.n: " + std::string(CHECK_MARK(has_tile_n)) + + (has_tile_n ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_block.tile_size.k: " + std::string(CHECK_MARK(has_tile_k)) + + (has_tile_k ? "\n" : " (missing or wrong type)\n"); + return msg; } } template -consteval auto detailed_diagnostic_SpecifiesTileTransfer() -> std::string { +consteval auto detailed_diagnostic_SpecifiesTileTransfer() -> std::string +{ std::string msg; constexpr bool has_transfer = requires { T::transfer; }; msg += " → T::transfer member: " + std::string(CHECK_MARK(has_transfer)) + "\n"; - - if constexpr (!has_transfer) { + + if constexpr(!has_transfer) + { return msg; } - - constexpr bool has_a_scalar = requires { { T::transfer.a_scalar_per_vector } -> std::convertible_to; }; - constexpr bool has_b_scalar = requires { { T::transfer.b_scalar_per_vector } -> std::convertible_to; }; - constexpr bool has_c_scalar = requires { { T::transfer.c_scalar_per_vector } -> std::convertible_to; }; - - msg += " → transfer.a_scalar_per_vector: " + std::string(CHECK_MARK(has_a_scalar)) + (has_a_scalar ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.b_scalar_per_vector: " + std::string(CHECK_MARK(has_b_scalar)) + (has_b_scalar ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.c_scalar_per_vector: " + std::string(CHECK_MARK(has_c_scalar)) + (has_c_scalar ? "\n" : " (missing or wrong type)\n"); - + + constexpr bool has_a_scalar = requires { + { T::transfer.a_scalar_per_vector } -> std::convertible_to; + }; + constexpr bool has_b_scalar = requires { + { T::transfer.b_scalar_per_vector } -> std::convertible_to; + }; + constexpr bool has_c_scalar = requires { + { T::transfer.c_scalar_per_vector } -> std::convertible_to; + }; + + msg += " → transfer.a_scalar_per_vector: " + std::string(CHECK_MARK(has_a_scalar)) + + (has_a_scalar ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.b_scalar_per_vector: " + std::string(CHECK_MARK(has_b_scalar)) + + (has_b_scalar ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.c_scalar_per_vector: " + std::string(CHECK_MARK(has_c_scalar)) + + (has_c_scalar ? "\n" : " (missing or wrong type)\n"); + return msg; } template -consteval auto detailed_diagnostic_SpecifiesTileBlockGemm() -> std::string { +consteval auto detailed_diagnostic_SpecifiesTileBlockGemm() -> std::string +{ std::string msg; - constexpr bool has_block_gemm = requires { { T::block_gemm } -> TileBlockGemmDescriptor; }; + constexpr bool has_block_gemm = requires { + { T::block_gemm } -> TileBlockGemmDescriptor; + }; msg += " → T::block_gemm member: " + std::string(CHECK_MARK(has_block_gemm)) + "\n"; - - if constexpr (!has_block_gemm) { + + if constexpr(!has_block_gemm) + { return msg; } - - using BG = decltype(T::block_gemm); - constexpr bool has_warps_m = requires(BG t) { { t.warps.m } -> std::convertible_to; }; - constexpr bool has_warps_n = requires(BG t) { { t.warps.n } -> std::convertible_to; }; - constexpr bool has_warps_k = requires(BG t) { { t.warps.k } -> std::convertible_to; }; - constexpr bool has_warp_tile_m = requires(BG t) { { t.warp_tile.m } -> std::convertible_to; }; - constexpr bool has_warp_tile_n = requires(BG t) { { t.warp_tile.n } -> std::convertible_to; }; - constexpr bool has_warp_tile_k = requires(BG t) { { t.warp_tile.k } -> std::convertible_to; }; - constexpr bool has_double_smem = requires(BG t) { { t.double_smem_buffer } -> std::convertible_to; }; - constexpr bool has_num_wave_groups = requires(BG t) { { t.num_wave_groups } -> std::convertible_to; }; - constexpr bool has_pipeline = requires(BG t) { { t.pipeline_version } -> std::convertible_to; }; - constexpr bool has_scheduler = requires(BG t) { { t.scheduler } -> std::convertible_to; }; - - msg += " → block_gemm.warps.m: " + std::string(CHECK_MARK(has_warps_m)) + (has_warps_m ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.warps.n: " + std::string(CHECK_MARK(has_warps_n)) + (has_warps_n ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.warps.k: " + std::string(CHECK_MARK(has_warps_k)) + (has_warps_k ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.warp_tile.m: " + std::string(CHECK_MARK(has_warp_tile_m)) + (has_warp_tile_m ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.warp_tile.n: " + std::string(CHECK_MARK(has_warp_tile_n)) + (has_warp_tile_n ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.warp_tile.k: " + std::string(CHECK_MARK(has_warp_tile_k)) + (has_warp_tile_k ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.double_smem_buffer: " + std::string(CHECK_MARK(has_double_smem)) + (has_double_smem ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.num_wave_groups: " + std::string(CHECK_MARK(has_num_wave_groups)) + (has_num_wave_groups ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.pipeline_version: " + std::string(CHECK_MARK(has_pipeline)) + (has_pipeline ? "\n" : " (missing or wrong type)\n"); - msg += " → block_gemm.scheduler: " + std::string(CHECK_MARK(has_scheduler)) + (has_scheduler ? "\n" : " (missing or wrong type)\n"); - + + using BG = decltype(T::block_gemm); + constexpr bool has_warps_m = requires(BG t) { + { t.warps.m } -> std::convertible_to; + }; + constexpr bool has_warps_n = requires(BG t) { + { t.warps.n } -> std::convertible_to; + }; + constexpr bool has_warps_k = requires(BG t) { + { t.warps.k } -> std::convertible_to; + }; + constexpr bool has_warp_tile_m = requires(BG t) { + { t.warp_tile.m } -> std::convertible_to; + }; + constexpr bool has_warp_tile_n = requires(BG t) { + { t.warp_tile.n } -> std::convertible_to; + }; + constexpr bool has_warp_tile_k = requires(BG t) { + { t.warp_tile.k } -> std::convertible_to; + }; + constexpr bool has_double_smem = requires(BG t) { + { t.double_smem_buffer } -> std::convertible_to; + }; + constexpr bool has_num_wave_groups = requires(BG t) { + { t.num_wave_groups } -> std::convertible_to; + }; + constexpr bool has_pipeline = requires(BG t) { + { t.pipeline_version } -> std::convertible_to; + }; + constexpr bool has_scheduler = requires(BG t) { + { t.scheduler } -> std::convertible_to; + }; + + msg += " → block_gemm.warps.m: " + std::string(CHECK_MARK(has_warps_m)) + + (has_warps_m ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.warps.n: " + std::string(CHECK_MARK(has_warps_n)) + + (has_warps_n ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.warps.k: " + std::string(CHECK_MARK(has_warps_k)) + + (has_warps_k ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.warp_tile.m: " + std::string(CHECK_MARK(has_warp_tile_m)) + + (has_warp_tile_m ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.warp_tile.n: " + std::string(CHECK_MARK(has_warp_tile_n)) + + (has_warp_tile_n ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.warp_tile.k: " + std::string(CHECK_MARK(has_warp_tile_k)) + + (has_warp_tile_k ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.double_smem_buffer: " + std::string(CHECK_MARK(has_double_smem)) + + (has_double_smem ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.num_wave_groups: " + std::string(CHECK_MARK(has_num_wave_groups)) + + (has_num_wave_groups ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.pipeline_version: " + std::string(CHECK_MARK(has_pipeline)) + + (has_pipeline ? "\n" : " (missing or wrong type)\n"); + msg += " → block_gemm.scheduler: " + std::string(CHECK_MARK(has_scheduler)) + + (has_scheduler ? "\n" : " (missing or wrong type)\n"); + return msg; } template -consteval auto detailed_diagnostic_SpecifiesTileOptimizations() -> std::string { +consteval auto detailed_diagnostic_SpecifiesTileOptimizations() -> std::string +{ std::string msg; - constexpr bool has_optimizations = requires { { T::optimizations } -> TileOptimizationsDescriptor; }; + constexpr bool has_optimizations = requires { + { T::optimizations } -> TileOptimizationsDescriptor; + }; msg += " → T::optimizations member: " + std::string(CHECK_MARK(has_optimizations)) + "\n"; - - if constexpr (!has_optimizations) { + + if constexpr(!has_optimizations) + { return msg; } - - using OPT = decltype(T::optimizations); - constexpr bool has_num_groups = requires(OPT t) { { t.num_groups_to_merge } -> std::convertible_to; }; - constexpr bool has_split_image = requires(OPT t) { { t.split_image } -> std::convertible_to; }; - constexpr bool has_explicit_gemm = requires(OPT t) { { t.explicit_gemm } -> std::convertible_to; }; - - msg += " → optimizations.num_groups_to_merge: " + std::string(CHECK_MARK(has_num_groups)) + (has_num_groups ? "\n" : " (missing or wrong type)\n"); - msg += " → optimizations.split_image: " + std::string(CHECK_MARK(has_split_image)) + (has_split_image ? "\n" : " (missing or wrong type)\n"); - msg += " → optimizations.explicit_gemm: " + std::string(CHECK_MARK(has_explicit_gemm)) + (has_explicit_gemm ? "\n" : " (missing or wrong type)\n"); - + + using OPT = decltype(T::optimizations); + constexpr bool has_num_groups = requires(OPT t) { + { t.num_groups_to_merge } -> std::convertible_to; + }; + constexpr bool has_split_image = requires(OPT t) { + { t.split_image } -> std::convertible_to; + }; + constexpr bool has_explicit_gemm = requires(OPT t) { + { t.explicit_gemm } -> std::convertible_to; + }; + + msg += " → optimizations.num_groups_to_merge: " + std::string(CHECK_MARK(has_num_groups)) + + (has_num_groups ? "\n" : " (missing or wrong type)\n"); + msg += " → optimizations.split_image: " + std::string(CHECK_MARK(has_split_image)) + + (has_split_image ? "\n" : " (missing or wrong type)\n"); + msg += " → optimizations.explicit_gemm: " + std::string(CHECK_MARK(has_explicit_gemm)) + + (has_explicit_gemm ? "\n" : " (missing or wrong type)\n"); + return msg; } // DL-specific diagnostics template -consteval auto detailed_diagnostic_SpecifiesDlThreadConfig() -> std::string { +consteval auto detailed_diagnostic_SpecifiesDlThreadConfig() -> std::string +{ std::string msg; - constexpr bool has_thread_config = requires { { T::thread_config } -> DlThreadConfigDescriptor; }; + constexpr bool has_thread_config = requires { + { T::thread_config } -> DlThreadConfigDescriptor; + }; msg += " → T::thread_config member: " + std::string(CHECK_MARK(has_thread_config)) + "\n"; - - if constexpr (!has_thread_config) { + + if constexpr(!has_thread_config) + { return msg; } - - using TC = decltype(T::thread_config); - constexpr bool has_k0 = requires(TC t) { { t.k0_per_block } -> std::convertible_to; }; - constexpr bool has_k1 = requires(TC t) { { t.k1 } -> std::convertible_to; }; - constexpr bool has_m1 = requires(TC t) { { t.m1_per_thread } -> std::convertible_to; }; - constexpr bool has_n1 = requires(TC t) { { t.n1_per_thread } -> std::convertible_to; }; - constexpr bool has_k = requires(TC t) { { t.k_per_thread } -> std::convertible_to; }; - - msg += " → thread_config.k0_per_block: " + std::string(CHECK_MARK(has_k0)) + (has_k0 ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_config.k1: " + std::string(CHECK_MARK(has_k1)) + (has_k1 ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_config.m1_per_thread: " + std::string(CHECK_MARK(has_m1)) + (has_m1 ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_config.n1_per_thread: " + std::string(CHECK_MARK(has_n1)) + (has_n1 ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_config.k_per_thread: " + std::string(CHECK_MARK(has_k)) + (has_k ? "\n" : " (missing or wrong type)\n"); - + + using TC = decltype(T::thread_config); + constexpr bool has_k0 = requires(TC t) { + { t.k0_per_block } -> std::convertible_to; + }; + constexpr bool has_k1 = requires(TC t) { + { t.k1 } -> std::convertible_to; + }; + constexpr bool has_m1 = requires(TC t) { + { t.m1_per_thread } -> std::convertible_to; + }; + constexpr bool has_n1 = requires(TC t) { + { t.n1_per_thread } -> std::convertible_to; + }; + constexpr bool has_k = requires(TC t) { + { t.k_per_thread } -> std::convertible_to; + }; + + msg += " → thread_config.k0_per_block: " + std::string(CHECK_MARK(has_k0)) + + (has_k0 ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_config.k1: " + std::string(CHECK_MARK(has_k1)) + + (has_k1 ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_config.m1_per_thread: " + std::string(CHECK_MARK(has_m1)) + + (has_m1 ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_config.n1_per_thread: " + std::string(CHECK_MARK(has_n1)) + + (has_n1 ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_config.k_per_thread: " + std::string(CHECK_MARK(has_k)) + + (has_k ? "\n" : " (missing or wrong type)\n"); + return msg; } template -consteval auto detailed_diagnostic_SpecifiesDlThreadCluster() -> std::string { +consteval auto detailed_diagnostic_SpecifiesDlThreadCluster() -> std::string +{ std::string msg; - constexpr bool has_thread_cluster = requires { { T::thread_cluster } -> DlThreadClusterDescriptor; }; - msg += " → T::thread_cluster member: " + std::string(CHECK_MARK(has_thread_cluster)) + "\n"; - - if constexpr (!has_thread_cluster) { + constexpr bool has_thread_cluster = requires { + { T::thread_cluster } -> DlThreadClusterDescriptor; + }; + msg += + " → T::thread_cluster member: " + std::string(CHECK_MARK(has_thread_cluster)) + "\n"; + + if constexpr(!has_thread_cluster) + { return msg; } - - using TC = decltype(T::thread_cluster); - constexpr bool has_m1_xs = requires(TC t) { { t.m1_xs } -> std::convertible_to>; }; - constexpr bool has_n1_xs = requires(TC t) { { t.n1_xs } -> std::convertible_to>; }; - - msg += " → thread_cluster.m1_xs: " + std::string(CHECK_MARK(has_m1_xs)) + (has_m1_xs ? "\n" : " (missing or wrong type)\n"); - msg += " → thread_cluster.n1_xs: " + std::string(CHECK_MARK(has_n1_xs)) + (has_n1_xs ? "\n" : " (missing or wrong type)\n"); - + + using TC = decltype(T::thread_cluster); + constexpr bool has_m1_xs = requires(TC t) { + { t.m1_xs } -> std::convertible_to>; + }; + constexpr bool has_n1_xs = requires(TC t) { + { t.n1_xs } -> std::convertible_to>; + }; + + msg += " → thread_cluster.m1_xs: " + std::string(CHECK_MARK(has_m1_xs)) + + (has_m1_xs ? "\n" : " (missing or wrong type)\n"); + msg += " → thread_cluster.n1_xs: " + std::string(CHECK_MARK(has_n1_xs)) + + (has_n1_xs ? "\n" : " (missing or wrong type)\n"); + return msg; } template -consteval auto detailed_diagnostic_SpecifiesDlFwdBlockTransfer() -> std::string { +consteval auto detailed_diagnostic_SpecifiesDlFwdBlockTransfer() -> std::string +{ std::string msg; constexpr bool has_transfer = requires { T::transfer; }; msg += " → T::transfer member: " + std::string(CHECK_MARK(has_transfer)) + "\n"; - - if constexpr (!has_transfer) { + + if constexpr(!has_transfer) + { return msg; } - - constexpr bool has_a = requires { { T::transfer.a } -> DlBlockTransferDescriptor4D; }; - constexpr bool has_b = requires { { T::transfer.b } -> DlBlockTransferDescriptor4D; }; + + constexpr bool has_a = requires { + { T::transfer.a } -> DlBlockTransferDescriptor4D; + }; + constexpr bool has_b = requires { + { T::transfer.b } -> DlBlockTransferDescriptor4D; + }; msg += " → T::transfer.a: " + std::string(CHECK_MARK(has_a)) + "\n"; - - if constexpr (has_a) { - using ABT = decltype(T::transfer.a); - constexpr bool has_thread_slice = requires(ABT t) { { t.thread_slice_lengths } -> std::convertible_to>; }; - constexpr bool has_thread_cluster = requires(ABT t) { { t.thread_cluster_lengths } -> std::convertible_to>; }; - constexpr bool has_cluster_arrange = requires(ABT t) { { t.thread_cluster_arrange_order } -> std::convertible_to>; }; - constexpr bool has_src_access = requires(ABT t) { { t.src_access_order } -> std::convertible_to>; }; - constexpr bool has_src_vector = requires(ABT t) { { t.src_vector_tensor_lengths } -> std::convertible_to>; }; - constexpr bool has_src_contiguous = requires(ABT t) { { t.src_vector_tensor_contiguous_dim_order } -> std::convertible_to>; }; - constexpr bool has_dst_vector = requires(ABT t) { { t.dst_vector_tensor_lengths } -> std::convertible_to>; }; - - msg += " → transfer.a.thread_slice_lengths (4D): " + std::string(CHECK_MARK(has_thread_slice)) + (has_thread_slice ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.thread_cluster_lengths (4D): " + std::string(CHECK_MARK(has_thread_cluster)) + (has_thread_cluster ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.thread_cluster_arrange_order (4D): " + std::string(CHECK_MARK(has_cluster_arrange)) + (has_cluster_arrange ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.src_access_order (4D): " + std::string(CHECK_MARK(has_src_access)) + (has_src_access ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.src_vector_tensor_lengths (4D): " + std::string(CHECK_MARK(has_src_vector)) + (has_src_vector ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.src_vector_tensor_contiguous_dim_order (4D): " + std::string(CHECK_MARK(has_src_contiguous)) + (has_src_contiguous ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.dst_vector_tensor_lengths (4D): " + std::string(CHECK_MARK(has_dst_vector)) + (has_dst_vector ? "\n" : " (missing or wrong type)\n"); - } else { + + if constexpr(has_a) + { + using ABT = decltype(T::transfer.a); + constexpr bool has_thread_slice = requires(ABT t) { + { t.thread_slice_lengths } -> std::convertible_to>; + }; + constexpr bool has_thread_cluster = requires(ABT t) { + { t.thread_cluster_lengths } -> std::convertible_to>; + }; + constexpr bool has_cluster_arrange = requires(ABT t) { + { t.thread_cluster_arrange_order } -> std::convertible_to>; + }; + constexpr bool has_src_access = requires(ABT t) { + { t.src_access_order } -> std::convertible_to>; + }; + constexpr bool has_src_vector = requires(ABT t) { + { t.src_vector_tensor_lengths } -> std::convertible_to>; + }; + constexpr bool has_src_contiguous = requires(ABT t) { + { + t.src_vector_tensor_contiguous_dim_order + } -> std::convertible_to>; + }; + constexpr bool has_dst_vector = requires(ABT t) { + { t.dst_vector_tensor_lengths } -> std::convertible_to>; + }; + + msg += " → transfer.a.thread_slice_lengths (4D): " + + std::string(CHECK_MARK(has_thread_slice)) + + (has_thread_slice ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.thread_cluster_lengths (4D): " + + std::string(CHECK_MARK(has_thread_cluster)) + + (has_thread_cluster ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.thread_cluster_arrange_order (4D): " + + std::string(CHECK_MARK(has_cluster_arrange)) + + (has_cluster_arrange ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.src_access_order (4D): " + + std::string(CHECK_MARK(has_src_access)) + + (has_src_access ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.src_vector_tensor_lengths (4D): " + + std::string(CHECK_MARK(has_src_vector)) + + (has_src_vector ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.src_vector_tensor_contiguous_dim_order (4D): " + + std::string(CHECK_MARK(has_src_contiguous)) + + (has_src_contiguous ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.dst_vector_tensor_lengths (4D): " + + std::string(CHECK_MARK(has_dst_vector)) + + (has_dst_vector ? "\n" : " (missing or wrong type)\n"); + } + else + { msg += " → T::transfer.a (4D): [✗] (missing or wrong type)\n"; } - + msg += " → T::transfer.b: " + std::string(CHECK_MARK(has_b)) + "\n"; - - if constexpr (has_b) { + + if constexpr(has_b) + { msg += " → T::transfer.b (4D): [✓] (similar fields as transfer.a)\n"; - } else { + } + else + { msg += " → T::transfer.b (4D): [✗] (missing or wrong type)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesDlBwdBlockTransfer() -> std::string { +consteval auto detailed_diagnostic_SpecifiesDlBwdBlockTransfer() -> std::string +{ std::string msg; constexpr bool has_transfer = requires { T::transfer; }; msg += " → T::transfer member: " + std::string(CHECK_MARK(has_transfer)) + "\n"; - - if constexpr (!has_transfer) { + + if constexpr(!has_transfer) + { return msg; } - - constexpr bool has_a = requires { { T::transfer.a } -> DlBlockTransferDescriptor5D; }; - constexpr bool has_b = requires { { T::transfer.b } -> DlBlockTransferDescriptor5D; }; + + constexpr bool has_a = requires { + { T::transfer.a } -> DlBlockTransferDescriptor5D; + }; + constexpr bool has_b = requires { + { T::transfer.b } -> DlBlockTransferDescriptor5D; + }; msg += " → T::transfer.a: " + std::string(CHECK_MARK(has_a)) + "\n"; - - if constexpr (has_a) { - using ABT = decltype(T::transfer.a); - constexpr bool has_thread_slice = requires(ABT t) { { t.thread_slice_lengths } -> std::convertible_to>; }; - constexpr bool has_thread_cluster = requires(ABT t) { { t.thread_cluster_lengths } -> std::convertible_to>; }; - constexpr bool has_cluster_arrange = requires(ABT t) { { t.thread_cluster_arrange_order } -> std::convertible_to>; }; - constexpr bool has_src_access = requires(ABT t) { { t.src_access_order } -> std::convertible_to>; }; - constexpr bool has_src_vector = requires(ABT t) { { t.src_vector_tensor_lengths } -> std::convertible_to>; }; - constexpr bool has_src_contiguous = requires(ABT t) { { t.src_vector_tensor_contiguous_dim_order } -> std::convertible_to>; }; - constexpr bool has_dst_vector = requires(ABT t) { { t.dst_vector_tensor_lengths } -> std::convertible_to>; }; - - msg += " → transfer.a.thread_slice_lengths (5D): " + std::string(CHECK_MARK(has_thread_slice)) + (has_thread_slice ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.thread_cluster_lengths (5D): " + std::string(CHECK_MARK(has_thread_cluster)) + (has_thread_cluster ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.thread_cluster_arrange_order (5D): " + std::string(CHECK_MARK(has_cluster_arrange)) + (has_cluster_arrange ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.src_access_order (5D): " + std::string(CHECK_MARK(has_src_access)) + (has_src_access ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.src_vector_tensor_lengths (5D): " + std::string(CHECK_MARK(has_src_vector)) + (has_src_vector ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.src_vector_tensor_contiguous_dim_order (5D): " + std::string(CHECK_MARK(has_src_contiguous)) + (has_src_contiguous ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.a.dst_vector_tensor_lengths (5D): " + std::string(CHECK_MARK(has_dst_vector)) + (has_dst_vector ? "\n" : " (missing or wrong type)\n"); - } else { + + if constexpr(has_a) + { + using ABT = decltype(T::transfer.a); + constexpr bool has_thread_slice = requires(ABT t) { + { t.thread_slice_lengths } -> std::convertible_to>; + }; + constexpr bool has_thread_cluster = requires(ABT t) { + { t.thread_cluster_lengths } -> std::convertible_to>; + }; + constexpr bool has_cluster_arrange = requires(ABT t) { + { t.thread_cluster_arrange_order } -> std::convertible_to>; + }; + constexpr bool has_src_access = requires(ABT t) { + { t.src_access_order } -> std::convertible_to>; + }; + constexpr bool has_src_vector = requires(ABT t) { + { t.src_vector_tensor_lengths } -> std::convertible_to>; + }; + constexpr bool has_src_contiguous = requires(ABT t) { + { + t.src_vector_tensor_contiguous_dim_order + } -> std::convertible_to>; + }; + constexpr bool has_dst_vector = requires(ABT t) { + { t.dst_vector_tensor_lengths } -> std::convertible_to>; + }; + + msg += " → transfer.a.thread_slice_lengths (5D): " + + std::string(CHECK_MARK(has_thread_slice)) + + (has_thread_slice ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.thread_cluster_lengths (5D): " + + std::string(CHECK_MARK(has_thread_cluster)) + + (has_thread_cluster ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.thread_cluster_arrange_order (5D): " + + std::string(CHECK_MARK(has_cluster_arrange)) + + (has_cluster_arrange ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.src_access_order (5D): " + + std::string(CHECK_MARK(has_src_access)) + + (has_src_access ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.src_vector_tensor_lengths (5D): " + + std::string(CHECK_MARK(has_src_vector)) + + (has_src_vector ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.src_vector_tensor_contiguous_dim_order (5D): " + + std::string(CHECK_MARK(has_src_contiguous)) + + (has_src_contiguous ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.a.dst_vector_tensor_lengths (5D): " + + std::string(CHECK_MARK(has_dst_vector)) + + (has_dst_vector ? "\n" : " (missing or wrong type)\n"); + } + else + { msg += " → T::transfer.a (5D): [✗] (missing or wrong type)\n"; } - + msg += " → T::transfer.b: " + std::string(CHECK_MARK(has_b)) + "\n"; - - if constexpr (has_b) { + + if constexpr(has_b) + { msg += " → T::transfer.b (5D): [✓] (similar fields as transfer.a)\n"; - } else { + } + else + { msg += " → T::transfer.b (5D): [✗] (missing or wrong type)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesDlEpilogue() -> std::string { +consteval auto detailed_diagnostic_SpecifiesDlEpilogue() -> std::string +{ std::string msg; constexpr bool has_transfer = requires { T::transfer; }; - if constexpr (!has_transfer) { + if constexpr(!has_transfer) + { return " → T::transfer member: [✗] (not found)\n"; } - + constexpr bool has_c = requires { T::transfer.c; }; msg += " → T::transfer.c: " + std::string(CHECK_MARK(has_c)) + "\n"; - - if constexpr (has_c && requires { T::transfer.c.src_dst_access_order; }) { - using C = decltype(T::transfer.c); - constexpr bool has_src_dst_access = requires(C t) { { t.src_dst_access_order } -> std::convertible_to>; }; - constexpr bool has_src_dst_vector_dim = requires(C t) { { t.src_dst_vector_dim } -> std::convertible_to; }; - constexpr bool has_dst_scalar = requires(C t) { { t.dst_scalar_per_vector } -> std::convertible_to; }; - - msg += " → transfer.c.src_dst_access_order: " + std::string(CHECK_MARK(has_src_dst_access)) + (has_src_dst_access ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.c.src_dst_vector_dim: " + std::string(CHECK_MARK(has_src_dst_vector_dim)) + (has_src_dst_vector_dim ? "\n" : " (missing or wrong type)\n"); - msg += " → transfer.c.dst_scalar_per_vector: " + std::string(CHECK_MARK(has_dst_scalar)) + (has_dst_scalar ? "\n" : " (missing or wrong type)\n"); - } else if constexpr (has_c) { + + if constexpr(has_c && requires { T::transfer.c.src_dst_access_order; }) + { + using C = decltype(T::transfer.c); + constexpr bool has_src_dst_access = requires(C t) { + { t.src_dst_access_order } -> std::convertible_to>; + }; + constexpr bool has_src_dst_vector_dim = requires(C t) { + { t.src_dst_vector_dim } -> std::convertible_to; + }; + constexpr bool has_dst_scalar = requires(C t) { + { t.dst_scalar_per_vector } -> std::convertible_to; + }; + + msg += " → transfer.c.src_dst_access_order: " + + std::string(CHECK_MARK(has_src_dst_access)) + + (has_src_dst_access ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.c.src_dst_vector_dim: " + + std::string(CHECK_MARK(has_src_dst_vector_dim)) + + (has_src_dst_vector_dim ? "\n" : " (missing or wrong type)\n"); + msg += " → transfer.c.dst_scalar_per_vector: " + + std::string(CHECK_MARK(has_dst_scalar)) + + (has_dst_scalar ? "\n" : " (missing or wrong type)\n"); + } + else if constexpr(has_c) + { msg += " → T::transfer.c (DlEpilogue): [✗] (missing required fields)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifiesTileConvSpecialization() -> std::string { - if constexpr (requires { T::specialization; }) { - using SpecType = decltype(T::specialization); +consteval auto detailed_diagnostic_SpecifiesTileConvSpecialization() -> std::string +{ + if constexpr(requires { T::specialization; }) + { + using SpecType = decltype(T::specialization); constexpr bool convertible = std::convertible_to; - return " → T::specialization: " + std::string(CHECK_MARK(convertible)) + + return " → T::specialization: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::specialization: [✗] (missing member)\n"; } } template -consteval auto detailed_diagnostic_SpecifiesLdsTransfer() -> std::string { +consteval auto detailed_diagnostic_SpecifiesLdsTransfer() -> std::string +{ std::string msg; - + constexpr bool has_transfer = requires { T::transfer; }; msg += " → T::transfer member: " + std::string(CHECK_MARK(has_transfer)) + "\n"; - - if constexpr (!has_transfer) { + + if constexpr(!has_transfer) + { return msg; } - - constexpr bool has_a = requires { { T::transfer.a.lds_transfer } -> LdsTransferDescriptor; }; + + constexpr bool has_a = requires { + { T::transfer.a.lds_transfer } -> LdsTransferDescriptor; + }; msg += " → T::transfer.a: " + std::string(CHECK_MARK(has_a)) + "\n"; - if constexpr (!has_a) { + if constexpr(!has_a) + { msg += " → T::transfer.a.lds_transfer: [✗] (missing or wrong type)\n"; } - - constexpr bool has_b = requires { { T::transfer.b.lds_transfer } -> LdsTransferDescriptor; }; + + constexpr bool has_b = requires { + { T::transfer.b.lds_transfer } -> LdsTransferDescriptor; + }; msg += " → T::transfer.b: " + std::string(CHECK_MARK(has_b)) + "\n"; - if constexpr (!has_b) { + if constexpr(!has_b) + { msg += " → T::transfer.b.lds_transfer: [✗] (missing or wrong type)\n"; } - - constexpr bool has_c = requires { { T::transfer.c.epilogue } -> EpilogueDescriptor; }; + + constexpr bool has_c = requires { + { T::transfer.c.epilogue } -> EpilogueDescriptor; + }; msg += " → T::transfer.c: " + std::string(CHECK_MARK(has_c)) + "\n"; - if constexpr (!has_c) { + if constexpr(!has_c) + { msg += " → T::transfer.c.epilogue: [✗] (missing or wrong type)\n"; } - + return msg; } template -consteval auto detailed_diagnostic_SpecifieGridwiseGemmPipeline() -> std::string { - if constexpr (requires { T::pipeline_version; }) { - using PipelineType = decltype(T::pipeline_version); +consteval auto detailed_diagnostic_SpecifieGridwiseGemmPipeline() -> std::string +{ + if constexpr(requires { T::pipeline_version; }) + { + using PipelineType = decltype(T::pipeline_version); constexpr bool convertible = std::convertible_to; - return " → T::pipeline_version: " + std::string(CHECK_MARK(convertible)) + + return " → T::pipeline_version: " + std::string(CHECK_MARK(convertible)) + (convertible ? "" : std::string(detail::get_type_info())) + "\n"; - } else { + } + else + { return " → T::pipeline_version: [✗] (missing member)\n"; } } 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 eb02aae584..9b02b4c9fc 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -10,27 +10,28 @@ namespace ck_tile::builder::factory { using namespace ck_tile::builder::diagnostics; template -struct ReferenceAlgorithm { +struct ReferenceAlgorithm +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesReferenceAlgorithm) static constexpr bool c1 = c_ConvAlgorithmDescriptor; static constexpr bool c2 = c_SpecifiesReferenceAlgorithm; - static consteval bool is_valid() { - return c1 && c2; - } + static consteval bool is_valid() { return c1 && c2; } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Reference Algorithm Diagnostic (closest match) ===\n" - "Concepts for Reference Algorithm:\n") + + "Concepts for Reference Algorithm:\n") + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesReferenceAlgorithm); } }; template -struct FwdXdlV3Algorithm { +struct FwdXdlV3Algorithm +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer) @@ -42,39 +43,39 @@ struct FwdXdlV3Algorithm { CHECK_CONCEPT(T, SpecifiesGemmSpecialization) CHECK_CONCEPT(T, SpecifiesBlockGemm) - static constexpr bool c1 = c_ConvAlgorithmDescriptor; - static constexpr bool c2 = c_SpecifiesThreadBlock; - static constexpr bool c3 = c_SpecifiesBlockTransfer; - static constexpr bool c4 = c_SpecifiesLdsTransfer; - static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = c_SpecifiesSourceAccessOrder; - static constexpr bool c7 = c_SpecifiesGridwiseFwdXdlGemm; - static constexpr bool c8 = c_SpecifiesFwdConvSpecialization; - static constexpr bool c9 = c_SpecifiesGemmSpecialization; + static constexpr bool c1 = c_ConvAlgorithmDescriptor; + static constexpr bool c2 = c_SpecifiesThreadBlock; + static constexpr bool c3 = c_SpecifiesBlockTransfer; + static constexpr bool c4 = c_SpecifiesLdsTransfer; + static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; + static constexpr bool c6 = c_SpecifiesSourceAccessOrder; + static constexpr bool c7 = c_SpecifiesGridwiseFwdXdlGemm; + static constexpr bool c8 = c_SpecifiesFwdConvSpecialization; + static constexpr bool c9 = c_SpecifiesGemmSpecialization; static constexpr bool c10 = c_SpecifiesBlockGemm; - static consteval bool is_valid() { + static consteval bool is_valid() + { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10; } - static consteval auto message() -> std::string { - return std::string("\n=== Forward XDL V3 Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdXdlV3 Algorithm:\n") + - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + static consteval auto message() -> std::string + { + return std::string("\n=== Forward XDL V3 Algorithm Diagnostic (closest match) ===\n" + "Concepts for FwdXdlV3 Algorithm:\n") + + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseFwdXdlGemm) + DIAGNOSTIC_LINE(SpecifiesFwdConvSpecialization) + - DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + - DIAGNOSTIC_LINE(SpecifiesBlockGemm); + DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + DIAGNOSTIC_LINE(SpecifiesBlockGemm); } }; template -struct FwdXdlAlgorithmBase { +struct FwdXdlAlgorithmBase +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer) @@ -88,60 +89,58 @@ struct FwdXdlAlgorithmBase { CHECK_CONCEPT(T, SpecifiesNumGroupsToMerge) CHECK_CONCEPT(T, SpecifiesLoopScheduler) - static constexpr bool c1 = c_ConvAlgorithmDescriptor; - static constexpr bool c2 = c_SpecifiesThreadBlock; - static constexpr bool c3 = c_SpecifiesBlockTransfer; - static constexpr bool c4 = c_SpecifiesLdsTransfer; - static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = c_SpecifiesSourceAccessOrder; - static constexpr bool c7 = c_SpecifiesGridwiseFwdXdlGemm; - static constexpr bool c8 = c_SpecifiesFwdConvSpecialization; - static constexpr bool c9 = c_SpecifiesGemmSpecialization; + static constexpr bool c1 = c_ConvAlgorithmDescriptor; + static constexpr bool c2 = c_SpecifiesThreadBlock; + static constexpr bool c3 = c_SpecifiesBlockTransfer; + static constexpr bool c4 = c_SpecifiesLdsTransfer; + static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; + static constexpr bool c6 = c_SpecifiesSourceAccessOrder; + static constexpr bool c7 = c_SpecifiesGridwiseFwdXdlGemm; + static constexpr bool c8 = c_SpecifiesFwdConvSpecialization; + static constexpr bool c9 = c_SpecifiesGemmSpecialization; static constexpr bool c10 = c_SpecifiesNumPrefetchStages; static constexpr bool c11 = c_SpecifiesNumGroupsToMerge; static constexpr bool c12 = c_SpecifiesLoopScheduler; - static consteval bool is_valid() { + static consteval bool is_valid() + { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10 && c11 && c12; } - static consteval auto message() -> std::string { - return - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + static consteval auto message() -> std::string + { + return DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseFwdXdlGemm) + DIAGNOSTIC_LINE(SpecifiesFwdConvSpecialization) + DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + - DIAGNOSTIC_LINE(SpecifiesNumGroupsToMerge) + - DIAGNOSTIC_LINE(SpecifiesLoopScheduler); + DIAGNOSTIC_LINE(SpecifiesNumGroupsToMerge) + DIAGNOSTIC_LINE(SpecifiesLoopScheduler); } }; template -struct FwdXdlAlgorithm : public FwdXdlAlgorithmBase{ +struct FwdXdlAlgorithm : public FwdXdlAlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesGenericInstance) - + static constexpr bool c13 = c_SpecifiesGenericInstance; - static consteval bool is_valid() { - return c13 && FwdXdlAlgorithmBase::is_valid(); - } + static consteval bool is_valid() { return c13 && FwdXdlAlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Forward XDL Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdXdl Algorithm:\n") + - FwdXdlAlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesGenericInstance); + "Concepts for FwdXdl Algorithm:\n") + + FwdXdlAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesGenericInstance); } }; template -struct FwdWmmaAlgorithm { +struct FwdWmmaAlgorithm +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer) @@ -155,43 +154,44 @@ struct FwdWmmaAlgorithm { CHECK_CONCEPT(T, SpecifiesLoopScheduler) CHECK_CONCEPT(T, SpecifiesGridwiseGemmPipeline) - static constexpr bool c1 = c_ConvAlgorithmDescriptor; - static constexpr bool c2 = c_SpecifiesThreadBlock; - static constexpr bool c3 = c_SpecifiesBlockTransfer; - static constexpr bool c4 = c_SpecifiesLdsTransfer; - static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = c_SpecifiesSourceAccessOrder; - static constexpr bool c7 = c_SpecifiesGridwiseWmmaGemm; - static constexpr bool c8 = c_SpecifiesFwdConvSpecialization; - static constexpr bool c9 = c_SpecifiesGemmSpecialization; + static constexpr bool c1 = c_ConvAlgorithmDescriptor; + static constexpr bool c2 = c_SpecifiesThreadBlock; + static constexpr bool c3 = c_SpecifiesBlockTransfer; + static constexpr bool c4 = c_SpecifiesLdsTransfer; + static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; + static constexpr bool c6 = c_SpecifiesSourceAccessOrder; + static constexpr bool c7 = c_SpecifiesGridwiseWmmaGemm; + static constexpr bool c8 = c_SpecifiesFwdConvSpecialization; + static constexpr bool c9 = c_SpecifiesGemmSpecialization; static constexpr bool c10 = c_SpecifiesNumPrefetchStages; static constexpr bool c11 = c_SpecifiesLoopScheduler; static constexpr bool c12 = c_SpecifiesGridwiseGemmPipeline; - static consteval bool is_valid() { + static consteval bool is_valid() + { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10 && c11 && c12; } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Forward WMMA Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdWmma Algorithm:\n") + - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + "Concepts for FwdWmma Algorithm:\n") + + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseWmmaGemm) + DIAGNOSTIC_LINE(SpecifiesFwdConvSpecialization) + DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + - DIAGNOSTIC_LINE(SpecifiesLoopScheduler) + + DIAGNOSTIC_LINE(SpecifiesLoopScheduler) + DIAGNOSTIC_LINE(SpecifiesGridwiseGemmPipeline); } }; template -struct FwdDlAlgorithm { +struct FwdDlAlgorithm +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesFwdConvSpecialization) @@ -210,26 +210,24 @@ struct FwdDlAlgorithm { static constexpr bool c7 = c_SpecifiesDlFwdBlockTransfer; static constexpr bool c8 = c_SpecifiesDlEpilogue; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Forward DL Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdDl Algorithm:\n") + - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + "Concepts for FwdDl Algorithm:\n") + + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + DIAGNOSTIC_LINE(SpecifiesFwdConvSpecialization) + DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + DIAGNOSTIC_LINE(SpecifiesDlThreadConfig) + DIAGNOSTIC_LINE(SpecifiesDlThreadCluster) + - DIAGNOSTIC_LINE(SpecifiesDlFwdBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesDlEpilogue); + DIAGNOSTIC_LINE(SpecifiesDlFwdBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesDlEpilogue); } }; template -struct TileAlgorithm { +struct TileAlgorithm +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesTileThreadBlock) CHECK_CONCEPT(T, SpecifiesTileTransfer) @@ -244,16 +242,14 @@ struct TileAlgorithm { static constexpr bool c5 = c_SpecifiesTileBlockGemm; static constexpr bool c6 = c_SpecifiesTileOptimizations; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6; } - static consteval auto message() -> std::string { - return std::string("\n=== CK Tile Algorithm Diagnostic (closest match) ===\n" - "Concepts for CK Tile Conv Algorithm:\n") + + static consteval auto message() -> std::string + { + return std::string("\n=== CK Tile Algorithm Diagnostic (closest match) ===\n" + "Concepts for CK Tile Conv Algorithm:\n") + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesTileThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesTileTransfer) + + DIAGNOSTIC_LINE(SpecifiesTileThreadBlock) + DIAGNOSTIC_LINE(SpecifiesTileTransfer) + DIAGNOSTIC_LINE(SpecifiesTileConvSpecialization) + DIAGNOSTIC_LINE(SpecifiesTileBlockGemm) + DIAGNOSTIC_LINE(SpecifiesTileOptimizations); @@ -267,21 +263,24 @@ struct LargeTensorAlgorithm : public FwdXdlAlgorithmBase static constexpr bool c13 = c_SpecifiesLargeTensorSupport; - static consteval bool is_valid() { + static consteval bool is_valid() + { // Note: Check first if the specialization is set. return c13 && FwdXdlAlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { - return std::string("\n=== Forward XDL Large Tensor Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdLargeTensorXdl Algorithm:\n") + - FwdXdlAlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesLargeTensorSupport); + static consteval auto message() -> std::string + { + return std::string( + "\n=== Forward XDL Large Tensor Algorithm Diagnostic (closest match) ===\n" + "Concepts for FwdLargeTensorXdl Algorithm:\n") + + FwdXdlAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesLargeTensorSupport); } }; template -struct BwdXdlAlgorithmBase { +struct BwdXdlAlgorithmBase +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer4D) @@ -300,16 +299,12 @@ struct BwdXdlAlgorithmBase { static constexpr bool c7 = c_SpecifiesGridwiseBwdXdlGemm; static constexpr bool c8 = c_SpecifiesBwdWeightConvSpecialization; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; } - static consteval auto message() -> std::string { - return - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer4D) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + static consteval auto message() -> std::string + { + return DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer4D) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseBwdXdlGemm) + @@ -318,46 +313,45 @@ struct BwdXdlAlgorithmBase { }; template -struct BwdXdlAlgorithm : public BwdXdlAlgorithmBase{ +struct BwdXdlAlgorithm : public BwdXdlAlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesTransposeTransfer) CHECK_CONCEPT(T, SpecifiesGenericInstance) - static constexpr bool c9 = c_SpecifiesTransposeTransfer; + static constexpr bool c9 = c_SpecifiesTransposeTransfer; static constexpr bool c10 = c_SpecifiesGenericInstance; - static consteval bool is_valid() { - return c9 && c10 && BwdXdlAlgorithmBase::is_valid(); - } + static consteval bool is_valid() { return c9 && c10 && BwdXdlAlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward XDL Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdXdl Algorithm:\n") + - BwdXdlAlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + + "Concepts for BwdXdl Algorithm:\n") + + BwdXdlAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + DIAGNOSTIC_LINE(SpecifiesGenericInstance); } }; template -struct BwdMultiDXdlAlgorithm : public BwdXdlAlgorithmBase{ +struct BwdMultiDXdlAlgorithm : public BwdXdlAlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesMultipleDSupport) static constexpr bool c9 = c_SpecifiesMultipleDSupport; - static consteval bool is_valid() { - return c9 && BwdXdlAlgorithmBase::is_valid(); - } + static consteval bool is_valid() { return c9 && BwdXdlAlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward XDL Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdXdl Algorithm:\n") + - BwdXdlAlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesMultipleDSupport); + "Concepts for BwdXdl Algorithm:\n") + + BwdXdlAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesMultipleDSupport); } }; template -struct BwdXdlV3AlgorithmBase { +struct BwdXdlV3AlgorithmBase +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer) @@ -378,16 +372,12 @@ struct BwdXdlV3AlgorithmBase { static constexpr bool c8 = c_SpecifiesBwdWeightConvSpecialization; static constexpr bool c9 = c_SpecifiesBlockGemm; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9; } - static consteval auto message() -> std::string { - return - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + static consteval auto message() -> std::string + { + return DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseBwdXdlGemm) + @@ -397,25 +387,25 @@ struct BwdXdlV3AlgorithmBase { }; template -struct BwdXdlV3Algorithm : public BwdXdlV3AlgorithmBase{ +struct BwdXdlV3Algorithm : public BwdXdlV3AlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesGenericInstance) static constexpr bool c10 = c_SpecifiesGenericInstance; - static consteval bool is_valid() { - return c10 && BwdXdlV3AlgorithmBase::is_valid(); - } + static consteval bool is_valid() { return c10 && BwdXdlV3AlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward XDL V3 Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdXdlV3 Algorithm:\n") + - BwdXdlV3AlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesGenericInstance); + "Concepts for BwdXdlV3 Algorithm:\n") + + BwdXdlV3AlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesGenericInstance); } }; template -struct BwdTwoStageXdlAlgorithm : public BwdXdlV3AlgorithmBase{ +struct BwdTwoStageXdlAlgorithm : public BwdXdlV3AlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesTransposeTransfer) CHECK_CONCEPT(T, SpecifiesGemmBatchOptions) CHECK_CONCEPT(T, SpecifiesTwoStageSupport) @@ -424,22 +414,24 @@ struct BwdTwoStageXdlAlgorithm : public BwdXdlV3AlgorithmBase{ static constexpr bool c11 = c_SpecifiesGemmBatchOptions; static constexpr bool c12 = c_SpecifiesTwoStageSupport; - static consteval bool is_valid() { + static consteval bool is_valid() + { return c10 && c11 && c12 && BwdXdlV3AlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward two stage XDL Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdXdlV3 Algorithm:\n") + - BwdXdlV3AlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + + "Concepts for BwdXdlV3 Algorithm:\n") + + BwdXdlV3AlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + DIAGNOSTIC_LINE(SpecifiesGemmBatchOptions) + DIAGNOSTIC_LINE(SpecifiesTwoStageSupport); } }; template -struct BwdWmmaAlgorithmBase { +struct BwdWmmaAlgorithmBase +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer) @@ -458,16 +450,12 @@ struct BwdWmmaAlgorithmBase { static constexpr bool c7 = c_SpecifiesGridwiseWmmaGemm; static constexpr bool c8 = c_SpecifiesBwdWeightConvSpecialization; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; } - static consteval auto message() -> std::string { - return - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + static consteval auto message() -> std::string + { + return DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseWmmaGemm) + @@ -476,26 +464,28 @@ struct BwdWmmaAlgorithmBase { }; template -struct BwdWmmaAlgorithm : public BwdWmmaAlgorithmBase { +struct BwdWmmaAlgorithm : public BwdWmmaAlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesNumPrefetchStages) CHECK_CONCEPT(T, SpecifiesLoopScheduler) CHECK_CONCEPT(T, SpecifiesGridwiseGemmPipeline) CHECK_CONCEPT(T, SpecifiesGenericInstance) - static constexpr bool c9 = c_SpecifiesNumPrefetchStages; + static constexpr bool c9 = c_SpecifiesNumPrefetchStages; static constexpr bool c10 = c_SpecifiesLoopScheduler; static constexpr bool c11 = c_SpecifiesGridwiseGemmPipeline; static constexpr bool c12 = c_SpecifiesGenericInstance; - static consteval bool is_valid() { + static consteval bool is_valid() + { return c9 && c10 && c11 && c12 && BwdWmmaAlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward WMMA Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdWmma Algorithm:\n") + - BwdWmmaAlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + + "Concepts for BwdWmma Algorithm:\n") + + BwdWmmaAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + DIAGNOSTIC_LINE(SpecifiesLoopScheduler) + DIAGNOSTIC_LINE(SpecifiesGridwiseGemmPipeline) + DIAGNOSTIC_LINE(SpecifiesGenericInstance); @@ -503,7 +493,8 @@ struct BwdWmmaAlgorithm : public BwdWmmaAlgorithmBase { }; template -struct BwdWmmaV3AlgorithmBase { +struct BwdWmmaV3AlgorithmBase +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBlockTransfer) @@ -524,16 +515,12 @@ struct BwdWmmaV3AlgorithmBase { static constexpr bool c8 = c_SpecifiesBwdWeightConvSpecialization; static constexpr bool c9 = c_SpecifiesBlockGemm; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9; } - static consteval auto message() -> std::string { - return - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + static consteval auto message() -> std::string + { + return DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + DIAGNOSTIC_LINE(SpecifiesGridwiseWmmaGemm) + @@ -543,25 +530,24 @@ struct BwdWmmaV3AlgorithmBase { }; template -struct BwdMultiDWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase { +struct BwdMultiDWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase +{ CHECK_CONCEPT(T, SpecifiesMultipleDSupport) static constexpr bool c10 = c_SpecifiesMultipleDSupport; - static consteval bool is_valid() { - return c10 && BwdWmmaAlgorithmBase::is_valid(); - } + static consteval bool is_valid() { return c10 && BwdWmmaAlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward WMMA Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdMultiDWmma Algorithm:\n") + - BwdWmmaAlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesMultipleDSupport); + "Concepts for BwdMultiDWmma Algorithm:\n") + + BwdWmmaAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesMultipleDSupport); } }; template -struct BwdWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase +struct BwdWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase { CHECK_CONCEPT(T, SpecifiesTransposeTransfer) CHECK_CONCEPT(T, SpecifiesGenericInstance) @@ -569,21 +555,19 @@ struct BwdWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase static constexpr bool c10 = c_SpecifiesTransposeTransfer; static constexpr bool c11 = c_SpecifiesGenericInstance; - static consteval bool is_valid() { - return c10 && c11 && BwdWmmaV3AlgorithmBase::is_valid(); - } + static consteval bool is_valid() { return c10 && c11 && BwdWmmaV3AlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward WMMA V3 Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdWmmaV3 Algorithm:\n") + - BwdWmmaV3AlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + + "Concepts for BwdWmmaV3 Algorithm:\n") + + BwdWmmaV3AlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + DIAGNOSTIC_LINE(SpecifiesGenericInstance); } }; template -struct BwdTwoStageWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase +struct BwdTwoStageWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase { CHECK_CONCEPT(T, SpecifiesTransposeTransfer) CHECK_CONCEPT(T, SpecifiesTwoStageSupport) @@ -593,22 +577,25 @@ struct BwdTwoStageWmmaV3Algorithm : public BwdWmmaV3AlgorithmBase static constexpr bool c11 = c_SpecifiesTwoStageSupport; static constexpr bool c12 = c_SpecifiesGemmBatchOptions; - static consteval bool is_valid() { + static consteval bool is_valid() + { return c10 && c11 && c12 && BwdWmmaV3AlgorithmBase::is_valid(); } - static consteval auto message() -> std::string { - return std::string("\n=== Backward Two Stage WMMA V3 Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdTwoStageWmmaV3 Algorithm:\n") + - BwdWmmaV3AlgorithmBase::message() + - DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + + static consteval auto message() -> std::string + { + return std::string( + "\n=== Backward Two Stage WMMA V3 Algorithm Diagnostic (closest match) ===\n" + "Concepts for BwdTwoStageWmmaV3 Algorithm:\n") + + BwdWmmaV3AlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesTransposeTransfer) + DIAGNOSTIC_LINE(SpecifiesGemmBatchOptions) + DIAGNOSTIC_LINE(SpecifiesTwoStageSupport); } }; template -struct BwdDlAlgorithm { +struct BwdDlAlgorithm +{ CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) CHECK_CONCEPT(T, SpecifiesBwdWeightConvSpecialization) @@ -625,286 +612,345 @@ struct BwdDlAlgorithm { static constexpr bool c6 = c_SpecifiesDlBwdBlockTransfer; static constexpr bool c7 = c_SpecifiesDlEpilogue; - static consteval bool is_valid() { - return c1 && c2 && c3 && c4 && c5 && c6 && c7; - } + static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7; } - static consteval auto message() -> std::string { + static consteval auto message() -> std::string + { return std::string("\n=== Backward DL Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdDl Algorithm:\n") + - DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + - DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + "Concepts for BwdDl Algorithm:\n") + + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + DIAGNOSTIC_LINE(SpecifiesBwdWeightConvSpecialization) + DIAGNOSTIC_LINE(SpecifiesDlThreadConfig) + DIAGNOSTIC_LINE(SpecifiesDlThreadCluster) + - DIAGNOSTIC_LINE(SpecifiesDlBwdBlockTransfer) + - DIAGNOSTIC_LINE(SpecifiesDlEpilogue); + DIAGNOSTIC_LINE(SpecifiesDlBwdBlockTransfer) + DIAGNOSTIC_LINE(SpecifiesDlEpilogue); } }; template -consteval int count_matches_fwd_xdl_v3() { +consteval int count_matches_fwd_xdl_v3() +{ using Alg = FwdXdlV3Algorithm; - return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9 + Alg::c10; + return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9 + + Alg::c10; } template -consteval int count_matches_fwd_xdl() { +consteval int count_matches_fwd_xdl() +{ using Alg = FwdXdlAlgorithm; - 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 + Alg::c13; + 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 + Alg::c13; } template -consteval int count_matches_fwd_wmma() { +consteval int count_matches_fwd_wmma() +{ using Alg = FwdWmmaAlgorithm; - return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9 + Alg::c10 + Alg::c11; + 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_fwd_dl() { +consteval int count_matches_fwd_dl() +{ using Alg = FwdDlAlgorithm; return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8; } template -consteval int count_matches_bwd_xdl() { +consteval int count_matches_bwd_xdl() +{ using Alg = BwdXdlAlgorithm; return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9; } template -consteval int count_matches_bwd_multi_d_xdl() { +consteval int count_matches_bwd_multi_d_xdl() +{ using Alg = BwdMultiDXdlAlgorithm; return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9; } template -consteval int count_matches_bwd_xdl_v3() { +consteval int count_matches_bwd_xdl_v3() +{ using Alg = BwdXdlV3Algorithm; return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7 + Alg::c8 + Alg::c9; } template -consteval int count_matches_bwd_two_stage_xdl() { +consteval int count_matches_bwd_two_stage_xdl() +{ using Alg = BwdTwoStageXdlAlgorithm; - 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; + 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() { +consteval int count_matches_bwd_wmma() +{ using Alg = BwdWmmaAlgorithm; - 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; + 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_multi_d_wmma() { +consteval int count_matches_bwd_multi_d_wmma() +{ using Alg = BwdMultiDWmmaV3Algorithm; - 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; + 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() { +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; + 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_two_stage_wmma_v3() { +consteval int count_matches_bwd_two_stage_wmma_v3() +{ using Alg = BwdTwoStageWmmaV3Algorithm; - 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; + 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_dl() { +consteval int count_matches_bwd_dl() +{ using Alg = BwdDlAlgorithm; return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6 + Alg::c7; } template -consteval int count_matches_large_tensor() { +consteval int count_matches_large_tensor() +{ using Alg = LargeTensorAlgorithm; - 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 + Alg::c13; + 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 + Alg::c13; } template -consteval int count_matches_tile() { +consteval int count_matches_tile() +{ using Alg = TileAlgorithm; return Alg::c1 + Alg::c2 + Alg::c3 + Alg::c4 + Alg::c5 + Alg::c6; } template -consteval void diagnose_fwd_algorithm_signature() +consteval void diagnose_fwd_algorithm_signature() { // Find closest matching variant - constexpr int xdl_v3_matches = count_matches_fwd_xdl_v3(); - constexpr int xdl_matches = count_matches_fwd_xdl(); - constexpr int wmma_matches = count_matches_fwd_wmma(); - constexpr int dl_matches = count_matches_fwd_dl(); + constexpr int xdl_v3_matches = count_matches_fwd_xdl_v3(); + constexpr int xdl_matches = count_matches_fwd_xdl(); + constexpr int wmma_matches = count_matches_fwd_wmma(); + constexpr int dl_matches = count_matches_fwd_dl(); constexpr int large_tensor_matches = count_matches_large_tensor(); - constexpr int tile_matches = count_matches_tile(); - + constexpr int tile_matches = count_matches_tile(); + // Check whether we have XDL or WMMA algorithm - if constexpr (SpecifiesGridwiseFwdXdlGemm) + 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_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) { + if constexpr(max_matches == xdl_v3_matches) + { using Alg = FwdXdlV3Algorithm; static_assert(Alg::is_valid(), Alg::message()); - } else if constexpr(max_matches == xdl_matches) { + } + else if constexpr(max_matches == xdl_matches) + { using Alg = FwdXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } else if constexpr(max_matches == dl_matches) { + } + 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) { + } + else if constexpr(max_matches == large_tensor_matches) + { using Alg = LargeTensorAlgorithm; static_assert(Alg::is_valid(), Alg::message()); } } - else if constexpr (SpecifiesGridwiseWmmaGemm) + else if constexpr(SpecifiesGridwiseWmmaGemm) { using Alg = FwdWmmaAlgorithm; static_assert(Alg::is_valid(), Alg::message()); } - else + 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_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) { + if constexpr(max_matches == xdl_v3_matches) + { using Alg = FwdXdlV3Algorithm; static_assert(Alg::is_valid(), Alg::message()); - } else if constexpr(max_matches == xdl_matches) { + } + else if constexpr(max_matches == xdl_matches) + { using Alg = FwdXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } else if constexpr(max_matches == wmma_matches) { + } + else if constexpr(max_matches == wmma_matches) + { using Alg = FwdWmmaAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } else if constexpr(max_matches == dl_matches) { + } + 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) { + } + else if constexpr(max_matches == large_tensor_matches) + { using Alg = LargeTensorAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } else if constexpr (max_matches == tile_matches) { + } + else if constexpr(max_matches == tile_matches) + { using Alg = TileAlgorithm; static_assert(Alg::is_valid(), Alg::message()); } - else { + else + { // This should never happen - static_assert(false, "Internal Error: No matching algorithm variant found for diagnostics."); + static_assert(false, + "Internal Error: No matching algorithm variant found for diagnostics."); } } } template -consteval void diagnose_bwd_weight_algorithm_signature() +consteval void diagnose_bwd_weight_algorithm_signature() { - constexpr int xdl_matches = count_matches_bwd_xdl(); - constexpr int xdl_v3_matches = count_matches_fwd_xdl_v3(); - 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 xdl_matches = count_matches_bwd_xdl(); + constexpr int xdl_v3_matches = count_matches_fwd_xdl_v3(); + 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 two_stage_wmma_v3_matches = count_matches_bwd_two_stage_wmma_v3(); - constexpr int wmma_matches = count_matches_bwd_wmma(); - constexpr int multi_d_wmma_matches = count_matches_bwd_multi_d_wmma(); + constexpr int wmma_matches = count_matches_bwd_wmma(); + constexpr int multi_d_wmma_matches = count_matches_bwd_multi_d_wmma(); // Check whether we have XDL or WMMA algorithm - if constexpr (SpecifiesGridwiseBwdXdlGemm) + 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 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) { + if constexpr(max_matches == xdl_matches) + { using Alg = BwdXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } - else if constexpr (max_matches == xdl_v3_matches) { + } + 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) { + 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) { + 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) { + else if constexpr(max_matches == multi_d_xdl_matches) + { using Alg = BwdMultiDXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); } } - else if constexpr (SpecifiesGridwiseWmmaGemm) - { - constexpr int max_1 = wmma_v3_matches > two_stage_wmma_v3_matches ? wmma_v3_matches : two_stage_wmma_v3_matches; - constexpr int max_2 = max_1 > wmma_matches ? max_1 : wmma_matches; + else if constexpr(SpecifiesGridwiseWmmaGemm) + { + constexpr int max_1 = wmma_v3_matches > two_stage_wmma_v3_matches + ? wmma_v3_matches + : two_stage_wmma_v3_matches; + constexpr int max_2 = max_1 > wmma_matches ? max_1 : wmma_matches; constexpr int max_matches = multi_d_wmma_matches > max_2 ? multi_d_wmma_matches : max_2; - if constexpr (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 == two_stage_wmma_v3_matches) { + else if constexpr(max_matches == two_stage_wmma_v3_matches) + { using Alg = BwdTwoStageWmmaV3Algorithm; static_assert(Alg::is_valid(), Alg::message()); } - else if constexpr (max_matches == wmma_matches) { + else if constexpr(max_matches == wmma_matches) + { using Alg = BwdWmmaAlgorithm; static_assert(Alg::is_valid(), Alg::message()); } - else if constexpr (max_matches == multi_d_wmma_matches) { + else if constexpr(max_matches == multi_d_wmma_matches) + { using Alg = BwdMultiDWmmaV3Algorithm; static_assert(Alg::is_valid(), Alg::message()); } } - else + 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 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) { + if constexpr(max_matches == xdl_matches) + { using Alg = BwdXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); - } - else if constexpr (max_matches == xdl_v3_matches) { + } + 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) { + 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) { + 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) { + else if constexpr(max_matches == multi_d_xdl_matches) + { using Alg = BwdMultiDXdlAlgorithm; static_assert(Alg::is_valid(), Alg::message()); } - else { + else + { // This should never happen - static_assert(false, "Internal Error: No matching algorithm variant found for diagnostics."); + static_assert(false, + "Internal Error: No matching algorithm variant found for diagnostics."); } } } -} +} // namespace ck_tile::builder::factory 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 203280c063..80143427c7 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 @@ -23,12 +23,13 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); static constexpr auto BLOCK = internal::SetThreadBlockInfo(); 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 9d94404a88..6f5c679b59 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 @@ -26,15 +26,16 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + 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 = @@ -47,58 +48,63 @@ struct ConvBwdWeightMultiDWmmaV3Factory 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"); + 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::DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3< - SPATIAL_DIM, - typename Layouts::InLayout, - typename Layouts::WeiLayout, - typename Layouts::OutLayout, - typename Layouts::DsLayout, - typename Types::InDataType, - typename Types::WeiDataType, - typename Types::OutDataType, - typename Types::AccDataType, - typename Types::DsDataType, - 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::OutComputeType, - typename Types::InComputeType>; + using Instance = + ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3< + SPATIAL_DIM, + typename Layouts::InLayout, + typename Layouts::WeiLayout, + typename Layouts::OutLayout, + typename Layouts::DsLayout, + typename Types::InDataType, + typename Types::WeiDataType, + typename Types::OutDataType, + typename Types::AccDataType, + typename Types::DsDataType, + 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::OutComputeType, + typename Types::InComputeType>; }; } // namespace ck_tile::builder::factory 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 dae333d99e..9f76568ca8 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 @@ -26,16 +26,17 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; - static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; 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 8288874f0e..86c48fe322 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 @@ -26,15 +26,16 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + 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 = @@ -47,59 +48,64 @@ struct ConvBwdWeightTwoStageWmmaV3Factory 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"); + 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::DeviceGroupedConvBwdWeightTwoStage_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, - ALGORITHM.num_conv_groups_to_merge, - typename Types::OutComputeType, - typename Types::InComputeType, - ALGORITHM.max_transpose_transfer_src_scalar_per_vector, - ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; + using Instance = + ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_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, + ALGORITHM.num_conv_groups_to_merge, + typename Types::OutComputeType, + typename Types::InComputeType, + 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_two_stage_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp index df3e4c01b2..9c37beae46 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 @@ -26,16 +26,17 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; - static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = @@ -48,10 +49,14 @@ struct ConvBwdWeightTwoStageXdlFactory 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"); + 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::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle< 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 b8c3bc8f9b..817432081b 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 @@ -26,15 +26,16 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; static constexpr auto GRIDWISE_GEMM_PIPELINE_VERSION = internal::SetGridwiseGemmPipelineVersion(); static constexpr auto LOOP_SCHEDULER = internal::SetLoopScheduler(); @@ -50,10 +51,14 @@ struct ConvBwdWeightWmmaFactory 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(AccessOrderLimits4D, "Invalid A thread cluster access order"); - static_assert(AccessOrderLimits4D, "Invalid B thread cluster access order"); - static_assert(AccessOrderLimits4D, "Invalid A source access order"); - static_assert(AccessOrderLimits4D, "Invalid B source access order"); + static_assert(AccessOrderLimits4D, + "Invalid A thread cluster access order"); + static_assert(AccessOrderLimits4D, + "Invalid B thread cluster access order"); + static_assert(AccessOrderLimits4D, + "Invalid A source access order"); + static_assert(AccessOrderLimits4D, + "Invalid B source access order"); // The forward convolution kernel class instance. using Instance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Wmma_CShuffle< 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 d1cfbe9e8d..baf84402c3 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 @@ -26,15 +26,16 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + 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 = @@ -47,10 +48,14 @@ struct ConvBwdWeightWmmaV3Factory 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"); + 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< 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 e89d227f82..91c19d2bd0 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 @@ -26,16 +26,17 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; - static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; 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_xdl_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp index e93b1456cd..f3edd0e6d9 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 @@ -26,16 +26,17 @@ template ; - using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + 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 BWD_CONV_SPECIALIZATION = + internal::SetBwdWeightConvSpecialization(); - static constexpr auto BLOCK = internal::SetThreadBlockInfo(); - static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; - static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; + static constexpr auto BLOCK = internal::SetThreadBlockInfo(); + static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; + static constexpr auto XDL_PARAMS = GRIDWISE_GEMM.xdl_params; static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = @@ -48,10 +49,14 @@ struct ConvBwdWeightXdlV3Factory 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"); + 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_Xdl_CShuffleV3< 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 9bca017735..892ad832d8 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp @@ -155,41 +155,45 @@ constexpr auto make_conv_instance() // Backward weight direction (will expand with more algorithms in the future) else if constexpr(ConvDirectionIsBackwardWeight) { - if constexpr (BwdXdlAlgorithm::is_valid()) + if constexpr(BwdXdlAlgorithm::is_valid()) { return typename ConvBwdWeightXdlFactory::Instance{}; } - else if constexpr (BwdXdlV3Algorithm::is_valid()) + else if constexpr(BwdXdlV3Algorithm::is_valid()) { return typename ConvBwdWeightXdlV3Factory::Instance{}; } - else if constexpr (BwdTwoStageXdlAlgorithm::is_valid()) + else if constexpr(BwdTwoStageXdlAlgorithm::is_valid()) { - return typename ConvBwdWeightTwoStageXdlFactory::Instance{}; + return + typename ConvBwdWeightTwoStageXdlFactory::Instance{}; } - else if constexpr (BwdDlAlgorithm::is_valid()) + else if constexpr(BwdDlAlgorithm::is_valid()) { return typename ConvBwdWeightDlFactory::Instance{}; } - else if constexpr (BwdMultiDXdlAlgorithm::is_valid()) + else if constexpr(BwdMultiDXdlAlgorithm::is_valid()) { - return typename ConvBwdWeightMultiDXdlFactory::Instance{}; + return + typename ConvBwdWeightMultiDXdlFactory::Instance{}; } - else if constexpr (BwdWmmaV3Algorithm::is_valid()) + else if constexpr(BwdWmmaV3Algorithm::is_valid()) { return typename ConvBwdWeightWmmaV3Factory::Instance{}; } - else if constexpr (BwdTwoStageWmmaV3Algorithm::is_valid()) + else if constexpr(BwdTwoStageWmmaV3Algorithm::is_valid()) { - return typename ConvBwdWeightTwoStageWmmaV3Factory::Instance{}; + return typename ConvBwdWeightTwoStageWmmaV3Factory:: + Instance{}; } - else if constexpr (BwdWmmaAlgorithm::is_valid()) + else if constexpr(BwdWmmaAlgorithm::is_valid()) { return typename ConvBwdWeightWmmaFactory::Instance{}; } - else if constexpr (BwdMultiDWmmaV3Algorithm::is_valid()) + else if constexpr(BwdMultiDWmmaV3Algorithm::is_valid()) { - return typename ConvBwdWeightMultiDWmmaV3Factory::Instance{}; + return typename ConvBwdWeightMultiDWmmaV3Factory:: + Instance{}; } else { diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp index 07f0976b74..31246eb5a8 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp @@ -24,10 +24,10 @@ template ; - using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + using Layouts = internal::ConvTensorLayouts; + using Types = internal::FwdConvTensorDataTypes; + using Ops = internal::ElementwiseOps; + using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); static constexpr auto GEMM_SPECIALIZATION = internal::SetGemmSpecialization(); diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp index 77a930d1ce..9cd56ad7ad 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp @@ -26,14 +26,13 @@ template ; - using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + using Layouts = internal::ConvTensorLayouts; + using Types = internal::FwdConvTensorDataTypes; + using Ops = internal::ElementwiseOps; + using AlgorithmType = decltype(ALGORITHM); - static constexpr auto FWD_CONV_SPECIALIZATION = - internal::SetFwdConvSpecialization(); - static constexpr auto GEMM_SPECIALIZATION = internal::SetGemmSpecialization(); + static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); + static constexpr auto GEMM_SPECIALIZATION = internal::SetGemmSpecialization(); static constexpr internal::ConvSpec SPECIALIZATION{.conv_spec = FWD_CONV_SPECIALIZATION, .gemm_spec = GEMM_SPECIALIZATION}; @@ -45,8 +44,7 @@ struct ConvFwdLargeTensorFactory internal::SetFwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = internal::SetFwdConvBlockTransfer(); - static constexpr auto C_BLOCK_TRANSFER = - internal::SetCBlockTransfer(); + static constexpr auto C_BLOCK_TRANSFER = internal::SetCBlockTransfer(); // Check limits for the algorithm parameters. static_assert(InputVectorTransferLimits); 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 e34f39965f..7d889a0c01 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 @@ -26,10 +26,10 @@ template ; - using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + using Layouts = internal::ConvTensorLayouts; + using Types = internal::FwdConvTensorDataTypes; + using Ops = internal::ElementwiseOps; + using AlgorithmType = decltype(ALGORITHM); static_assert(ALGORITHM.transfer.a.lds_transfer.is_direct_load == ALGORITHM.transfer.b.lds_transfer.is_direct_load, diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp index dbaa8651eb..3506f5d1a9 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp @@ -26,10 +26,10 @@ template ; - using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + using Layouts = internal::ConvTensorLayouts; + using Types = internal::FwdConvTensorDataTypes; + using Ops = internal::ElementwiseOps; + using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); static constexpr auto GEMM_SPECIALIZATION = internal::SetGemmSpecialization(); diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp index 5a0084d6da..446ceceda2 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp @@ -26,10 +26,10 @@ template ; - using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; - using AlgorithmType = decltype(ALGORITHM); + using Layouts = internal::ConvTensorLayouts; + using Types = internal::FwdConvTensorDataTypes; + using Ops = internal::ElementwiseOps; + using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); static constexpr auto GEMM_SPECIALIZATION = internal::SetGemmSpecialization(); 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 69facce41b..4ef2f533c9 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 @@ -15,11 +15,11 @@ struct BlockTransfer ck::Array thread_cluster_dims{}; // k0, m, k1 ck::Array thread_cluster_order{}; ck::Array src_access_order{}; - 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 @@ -28,10 +28,10 @@ struct BwdBlockTransfer ck::Array thread_cluster_dims{}; ck::Array thread_cluster_order{}; ck::Array src_access_order{}; - size_t src_vector_dim = 0; - size_t src_scalar_per_vector = 0; - size_t lds_dst_scalar_per_vector = 0; - 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 lds_padding = false; }; template @@ -66,11 +66,13 @@ constexpr auto SetBwdConvBlockTransfer() static_assert(block_order.order.size() == src_order.order.size(), "Mismatched size between block order and src order"); - if constexpr (array_length == 3) + if constexpr(array_length == 3) { return BwdBlockTransfer<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]}, + .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]}, .src_vector_dim = lds_cfg.src_vector_dim, .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, @@ -78,14 +80,23 @@ constexpr auto SetBwdConvBlockTransfer() .lds_padding = lds_cfg.lds_padding, }; } - else if constexpr (array_length == 4) + else if constexpr(array_length == 4) { return BwdBlockTransfer<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]}, - .src_vector_dim = lds_cfg.src_vector_dim, - .src_scalar_per_vector = lds_cfg.src_scalar_per_vector, + .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]}, + .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, }; diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp index 00205d414e..c7d9ce0ac6 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp @@ -64,19 +64,19 @@ consteval auto GetElementwiseOp() template struct ElementwiseOps { -private: + private: static constexpr auto input_op = GetElementwiseOp(); static constexpr auto weight_op = GetElementwiseOp(); static constexpr auto output_op = GetElementwiseOp(); - static constexpr bool is_forward = ConvDirectionIsForward; + static constexpr bool is_forward = ConvDirectionIsForward; static constexpr bool is_bwd_weight = ConvDirectionIsBackwardWeight; using InputOp = typename decltype(input_op)::Op; using WeightOp = typename decltype(weight_op)::Op; using OutputOp = typename decltype(output_op)::Op; -public: + public: // Forward convolution elementwise ops using AElementwiseOp = std::conditional_t; using BElementwiseOp = std::conditional_t; diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp index fd6df5f09a..d08dddff83 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp @@ -222,8 +222,8 @@ template ValidConvOutputLayoutForSpatialDim) struct ConvTensorLayouts { -private: - static constexpr bool is_forward = ConvDirectionIsForward; + private: + static constexpr bool is_forward = ConvDirectionIsForward; static constexpr bool is_bwd_weight = ConvDirectionIsBackwardWeight; using InputLayout = decltype(TensorLayoutToCK()); @@ -231,12 +231,12 @@ private: using OutputLayout = decltype(TensorLayoutToCK()); using AuxLayout = decltype(GetAuxiliaryTensorLayouts())::type; -public: + public: // Forward convolution layouts - using ALayout = std::conditional_t; - using BLayout = std::conditional_t; - using ELayout = std::conditional_t; - + using ALayout = std::conditional_t; + using BLayout = std::conditional_t; + using ELayout = std::conditional_t; + // Backward weight convolution layouts using InLayout = std::conditional_t; using WeiLayout = std::conditional_t; diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp index 2ab1c40882..1cecb8d43b 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp @@ -186,8 +186,8 @@ struct BwdWeightConvTensorDataTypes static constexpr auto output_types = GetTensorDataAndComputeTypes(); - using InDataType = typename decltype(input_types.first)::type; - using InComputeType = typename decltype(input_types.second)::type; + using InDataType = typename decltype(input_types.first)::type; + using InComputeType = typename decltype(input_types.second)::type; using WeiDataType = typename decltype(weight_types.first)::type; using WeiComputeType = typename decltype(weight_types.second)::type; using OutDataType = typename decltype(output_types.first)::type; 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 dde8e19f43..9ed1eebc3c 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 @@ -160,17 +160,19 @@ consteval ck::tensor_operation::device::ConvolutionForwardSpecialization SetFwdC } template -consteval ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization SetBwdWeightConvSpecialization() +consteval ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization +SetBwdWeightConvSpecialization() { constexpr auto specialization = ALGORITHM.bwd_weight_specialization; - using ck_conv_spec = ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization; + using ck_conv_spec = ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization; switch(specialization) { case ConvSpecialization::DEFAULT: return ck_conv_spec::Default; case ConvSpecialization::FILTER_1X1_PAD0: return ck_conv_spec::Filter1x1Pad0; case ConvSpecialization::FILTER_1X1_STRIDE1_PAD0: return ck_conv_spec::Filter1x1Stride1Pad0; case ConvSpecialization::ODD_C: return ck_conv_spec::OddC; - case ConvSpecialization::FILTER_3x3: throw "FILTER_3x3 is not supported for backward weight convolution."; + case ConvSpecialization::FILTER_3x3: + throw "FILTER_3x3 is not supported for backward weight convolution."; default: throw "Unsupported ConvSpecialization"; } } diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_dl.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_dl.cpp index 99cadb6d20..584bce2f1b 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_dl.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_dl.cpp @@ -19,11 +19,11 @@ constexpr auto SIGNATURE = .output = {.config = {.layout = ckb::TensorLayout::GNHWK}}}; constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl{} - .with_thread_block(cku::ThreadBlock_256_128x128x16) - .with_bwd_specialization(cku::ConvSpecialization::DEFAULT) - .with_dl_thread_config(cku::DlThreadConfig_16x1x4x4x1) - .with_dl_thread_cluster(cku::DlThreadCluster_8x2) - .with_dl_transfer(cku::DlTransfer5D); + .with_thread_block(cku::ThreadBlock_256_128x128x16) + .with_bwd_specialization(cku::ConvSpecialization::DEFAULT) + .with_dl_thread_config(cku::DlThreadConfig_16x1x4x4x1) + .with_dl_thread_cluster(cku::DlThreadCluster_8x2) + .with_dl_transfer(cku::DlTransfer5D); using Builder = ckb::ConvBuilder; using Instance = Builder::Instance; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp index e581da4951..782f33f845 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp @@ -11,14 +11,13 @@ 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 = 2, - .direction = ckb::ConvDirection::BACKWARD_WEIGHT, - .data_type = ckb::DataType::FP16, - .accumulation_data_type = ckb::DataType::FP32, - .input = {.config = {.layout = NGCHW}}, - .weight = {.config = {.layout = GKYXC}}, - .output = {.config = {.layout = NGKHW}}}; +constexpr auto SIGNATURE = ckt::ConvSignature{.spatial_dim = 2, + .direction = ckb::ConvDirection::BACKWARD_WEIGHT, + .data_type = ckb::DataType::FP16, + .accumulation_data_type = ckb::DataType::FP32, + .input = {.config = {.layout = NGCHW}}, + .weight = {.config = {.layout = GKYXC}}, + .output = {.config = {.layout = NGKHW}}}; constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3{} .with_thread_block(cku::ThreadBlock_64_32x32x32) @@ -27,7 +26,7 @@ constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStag .with_bwd_specialization(ckb::ConvSpecialization::DEFAULT) .with_block_gemm(cku::BlockGemmDesc_v1_intrawave) .with_num_conv_groups_to_merge(2) - .with_transpose_params(2,2); + .with_transpose_params(2, 2); using Builder = ckb::ConvBuilder; using Instance = Builder::Instance; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp index 9a8b9573fa..a2a877dbcd 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp @@ -39,6 +39,6 @@ TEST(BwdWeight_2DBf16_TwoStage_CShuffle, Create) "Default", "GNHWC,GKYXC,GNHWK", "PassThrough,PassThrough,PassThrough", - "Intrawave,v2", // pipeline versions + "Intrawave,v2", // pipeline versions "bf16,bf16,2,4>"}); // compute types and transpose params } 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 47e07d0722..0981ea6c11 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 @@ -11,14 +11,13 @@ 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 = 3, - .direction = ckb::ConvDirection::BACKWARD_WEIGHT, - .data_type = ckb::DataType::BF16, - .accumulation_data_type = ckb::DataType::FP32, - .input = {.config = {.layout = NGCDHW}}, - .weight = {.config = {.layout = GKZYXC}}, - .output = {.config = {.layout = NGKDHW}}}; +constexpr auto SIGNATURE = ckt::ConvSignature{.spatial_dim = 3, + .direction = ckb::ConvDirection::BACKWARD_WEIGHT, + .data_type = ckb::DataType::BF16, + .accumulation_data_type = ckb::DataType::FP32, + .input = {.config = {.layout = NGCDHW}}, + .weight = {.config = {.layout = GKZYXC}}, + .output = {.config = {.layout = NGKDHW}}}; constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle{} .with_thread_block(cku::ThreadBlock_64_32x32x32) 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 index f58ec11129..60f7d5bd64 100644 --- 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 @@ -11,22 +11,22 @@ 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 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_16x16_2x1_per_wave) - .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3) - .with_bwd_specialization(ckb::ConvSpecialization::FILTER_1X1_STRIDE1_PAD0) - .with_block_gemm(cku::BlockGemmDesc_v1_intrawave) - .with_transpose_params(4,4); +constexpr auto ALGORITHM = + cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3{} + .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::FILTER_1X1_STRIDE1_PAD0) + .with_block_gemm(cku::BlockGemmDesc_v1_intrawave) + .with_transpose_params(4, 4); using Builder = ckb::ConvBuilder; using Instance = Builder::Instance; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp index fdc17fba2a..4ad97209e5 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp @@ -11,21 +11,21 @@ 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 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_Xdl_CShuffle_V3{} - .with_thread_block(cku::ThreadBlock_64_32x32x32) - .with_gemm_config(cku::BwdGemmParams_Xdl_1x1_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); +constexpr auto ALGORITHM = + cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle_V3{} + .with_thread_block(cku::ThreadBlock_64_32x32x32) + .with_gemm_config(cku::BwdGemmParams_Xdl_1x1_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); using Builder = ckb::ConvBuilder; using Instance = Builder::Instance; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp index 3eca32c897..8d85370b26 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp @@ -34,7 +34,7 @@ TEST(FwdConvInstances, .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) .with_fwd_specializations(ConvSpecialization::FILTER_1X1_STRIDE1_PAD0, - GemmSpecialization::MNKPadding) + GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v2_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp index a8df3c0d98..e4675c89a7 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp @@ -25,7 +25,7 @@ TEST(FwdConvInstances, .accumulation_data_type = INT32, .input = {.config = {.layout = GNWC}}, .weight = {.config = {.layout = GKXC}}, - .output = {.config = {.layout = GNWK}}}; + .output = {.config = {.layout = GNWK}}}; constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle{} 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 0ce530c1f8..610e2fad5f 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 @@ -67,7 +67,8 @@ TEST(FwdConvInstances, .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_fwd_specializations(ConvSpecialization::FILTER_3x3, + GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v5_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp index 79d4827fee..58171cd530 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp @@ -62,7 +62,7 @@ TEST(FwdConvInstances, ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK{} .with_thread_block(ThreadBlock_256_128x128x16) .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, - GemmSpecialization::MNKPadding) + GemmSpecialization::MNKPadding) .with_dl_thread_config(DlThreadConfig_16x2x4x4x1) .with_dl_thread_cluster(DlThreadCluster_8x2) .with_dl_transfer(DlTransfer4D); diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp index 5e114cb761..8ae88a0917 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp @@ -25,7 +25,7 @@ constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xd .with_gemm_config(cku::FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(cku::Transfer_4x64x1) .with_fwd_specializations(ckb::ConvSpecialization::DEFAULT, - ckb::GemmSpecialization::MNKPadding) + ckb::GemmSpecialization::MNKPadding) .with_block_gemm(cku::BlockGemmDesc_v3_intrawave); using Builder = ckb::ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp index a8f97c417e..bb35c53ba0 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) .with_fwd_specializations(ConvSpecialization::FILTER_1X1_STRIDE1_PAD0, - GemmSpecialization::MNKPadding) + GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v4_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp index 303875b348..97bc0a00e5 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp @@ -26,13 +26,12 @@ TEST(FwdConvInstances, constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor{} - .with_thread_block(ThreadBlock_256_256x128x32) - .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) - .with_transfer(Transfer_4x16x1) - .with_fwd_specializations(ConvSpecialization::DEFAULT, - GemmSpecialization::MNKPadding) - .with_prefetch_config(1, PipelineScheduler::DEFAULT) - .with_num_conv_groups_to_merge(1); + .with_thread_block(ThreadBlock_256_256x128x32) + .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) + .with_transfer(Transfer_4x16x1) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_prefetch_config(1, PipelineScheduler::DEFAULT) + .with_num_conv_groups_to_merge(1); using Builder = ConvBuilder; @@ -63,13 +62,13 @@ TEST( constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor{} - .with_thread_block(ThreadBlock_128_128x128x32) - .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) - .with_transfer(Transfer_4x16x1) - .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, - GemmSpecialization::MNKPadding) - .with_prefetch_config(1, PipelineScheduler::DEFAULT) - .with_num_conv_groups_to_merge(1); + .with_thread_block(ThreadBlock_128_128x128x32) + .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) + .with_transfer(Transfer_4x16x1) + .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, + GemmSpecialization::MNKPadding) + .with_prefetch_config(1, PipelineScheduler::DEFAULT) + .with_num_conv_groups_to_merge(1); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp index 46b82b2407..56d4b8be59 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp @@ -31,7 +31,7 @@ TEST(FwdConvInstances, .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) .with_transfer(Transfer_4x64x1) .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, - GemmSpecialization::MNKPadding) + GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v4_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp index 6b69ca89a6..df8339241b 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp @@ -31,7 +31,7 @@ TEST(FwdConvInstances, .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, - GemmSpecialization::MNKPadding) + GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v1_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_data_2d_fp16_v3.cpp b/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_data_2d_fp16_v3.cpp index 677a304371..b79fdf513a 100644 --- a/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_data_2d_fp16_v3.cpp +++ b/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_data_2d_fp16_v3.cpp @@ -10,13 +10,14 @@ using namespace ck_tile::builder::test_utils; TEST(BwdDataConvInstances, Create_ConvAlgorithm_Tile_GroupedConvolutionKernel_2D_FP16_NHWGC) { - constexpr ConvSignature BwdDataConvSignature{.spatial_dim = 2, - .direction = ConvDirection::BACKWARD_DATA, - .data_type = DataType::FP16, - .accumulation_data_type = DataType::FP32, - .input = {.config = {.layout = TensorLayout::NHWGC}}, - .weight = {.config = {.layout = TensorLayout::GKYXC}}, - .output = {.config = {.layout = TensorLayout::NHWGK}}}; + constexpr ConvSignature BwdDataConvSignature{ + .spatial_dim = 2, + .direction = ConvDirection::BACKWARD_DATA, + .data_type = DataType::FP16, + .accumulation_data_type = DataType::FP32, + .input = {.config = {.layout = TensorLayout::NHWGC}}, + .weight = {.config = {.layout = TensorLayout::GKYXC}}, + .output = {.config = {.layout = TensorLayout::NHWGK}}}; constexpr auto BwdDataConvAlgorithm = ConvAlgorithm_Tile_GroupedConvolutionKernel{} diff --git a/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_weight_2d_fp16_v3.cpp b/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_weight_2d_fp16_v3.cpp index f3de0bb762..a5801b0e85 100644 --- a/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_weight_2d_fp16_v3.cpp +++ b/experimental/builder/test/conv/ck_tile/test_ckb_conv_bwd_weight_2d_fp16_v3.cpp @@ -10,13 +10,14 @@ using namespace ck_tile::builder::test_utils; TEST(BwdWeightConvInstances, Create_ConvAlgorithm_Tile_GroupedConvolutionKernel_2D_FP16_NHWGC) { - constexpr ConvSignature BwdWeightConvSignature{.spatial_dim = 2, - .direction = ConvDirection::BACKWARD_WEIGHT, - .data_type = DataType::FP16, - .accumulation_data_type = DataType::FP32, - .input = {.config = {.layout = TensorLayout::NHWGC}}, - .weight = {.config = {.layout = TensorLayout::GKYXC}}, - .output = {.config = {.layout = TensorLayout::NHWGK}}}; + constexpr ConvSignature BwdWeightConvSignature{ + .spatial_dim = 2, + .direction = ConvDirection::BACKWARD_WEIGHT, + .data_type = DataType::FP16, + .accumulation_data_type = DataType::FP32, + .input = {.config = {.layout = TensorLayout::NHWGC}}, + .weight = {.config = {.layout = TensorLayout::GKYXC}}, + .output = {.config = {.layout = TensorLayout::NHWGK}}}; constexpr auto BwdWeightConvAlgorithm = ConvAlgorithm_Tile_GroupedConvolutionKernel{} diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp index 57ef018ae3..797f440bdc 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -41,15 +41,15 @@ static_assert(ckb::GridwiseXdlGemmDescriptor); struct GridwiseFwdXdlGemm { // NOTE: ak1 and bk1 are difficult to verify in the kernel instantiation!!! - size_t ak1 = 0; - size_t bk1 = 0; + size_t ak1 = 0; + size_t bk1 = 0; XdlParams xdl_params; }; static_assert(ckb::GridwiseFwdXdlGemmDescriptor); struct GridwiseBwdXdlGemm { - size_t k1 = 0; + size_t k1 = 0; XdlParams xdl_params; }; static_assert(ckb::GridwiseBwdXdlGemmDescriptor); @@ -284,17 +284,20 @@ struct DlTransfer_ struct TwoStageSpecialization_ { - static constexpr ConvAlgorithmSpecialization specialization = ConvAlgorithmSpecialization::TWO_STAGE; + static constexpr ConvAlgorithmSpecialization specialization = + ConvAlgorithmSpecialization::TWO_STAGE; }; struct MultipleDSpecialization_ { - static constexpr ConvAlgorithmSpecialization specialization = ConvAlgorithmSpecialization::MULTIPLE_D; + static constexpr ConvAlgorithmSpecialization specialization = + ConvAlgorithmSpecialization::MULTIPLE_D; }; struct LargeTensorSpecialization_ { - static constexpr ConvAlgorithmSpecialization specialization = ConvAlgorithmSpecialization::LARGE_TENSOR; + static constexpr ConvAlgorithmSpecialization specialization = + ConvAlgorithmSpecialization::LARGE_TENSOR; }; // Specify thread block dimensions for a GEMM (CK Tile). @@ -395,7 +398,8 @@ struct ConvAlgorithmTemplate : Components... { result.gridwise_gemm = gemm; } - else { + else + { static_assert(false, "Unrecognized GemmConfig type"); } return result; @@ -412,7 +416,7 @@ struct ConvAlgorithmTemplate : Components... } constexpr auto with_fwd_specializations(ConvSpecialization fwd_spec, - GemmSpecialization gemm_spec) const + GemmSpecialization gemm_spec) const { static_assert(std::is_base_of_v); auto result = *this; @@ -424,7 +428,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_bwd_specialization(ConvSpecialization bwd_spec) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.bwd_weight_specialization = bwd_spec; return result; } @@ -442,7 +446,7 @@ struct ConvAlgorithmTemplate : Components... size_t max_dst_scalar_per_vector) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.max_transpose_transfer_src_scalar_per_vector = max_src_scalar_per_vector; result.max_transpose_transfer_dst_scalar_per_vector = max_dst_scalar_per_vector; return result; @@ -451,7 +455,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.num_conv_groups_to_merge = num_groups_to_merge; return result; } @@ -460,7 +464,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_block_gemm(const BG& bg) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.block_gemm_pipeline = bg; return result; } @@ -468,7 +472,7 @@ struct ConvAlgorithmTemplate : Components... constexpr auto with_gridwise_gemm_pipeline(const PipelineVersion plv) const { static_assert(std::is_base_of_v); - auto result = *this; + auto result = *this; result.pipeline_version = plv; return result; } @@ -550,13 +554,28 @@ struct ConvAlgorithmTemplate : Components... // Fwd algorithm types using ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle = - ConvAlgorithmTemplate, ConvSpecializationFwd_, Prefetch_, GemmBatchOptions_>; + ConvAlgorithmTemplate, + ConvSpecializationFwd_, + Prefetch_, + GemmBatchOptions_>; using ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 = - ConvAlgorithmTemplate, ConvSpecializationFwd_, BlockGemm_>; + ConvAlgorithmTemplate, + ConvSpecializationFwd_, + BlockGemm_>; using ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle = - ConvAlgorithmTemplate, ConvSpecializationFwd_, GridGemm_, Prefetch_, GemmBatchOptions_>; + ConvAlgorithmTemplate, + ConvSpecializationFwd_, + GridGemm_, + Prefetch_, + GemmBatchOptions_>; using ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK = ConvAlgorithmTemplate>; using ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor = - ConvAlgorithmTemplate, ConvSpecializationFwd_, Prefetch_, GemmBatchOptions_, LargeTensorSpecialization_>; + ConvAlgorithmTemplate, + ConvSpecializationFwd_, + Prefetch_, + GemmBatchOptions_, + LargeTensorSpecialization_>; // CK Tile algorithm using ConvAlgorithm_Tile_GroupedConvolutionKernel = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, TransposeParams_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + TransposeParams_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Xdl_CShuffle = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, BlockGemm_, TransposeParams_, GemmBatchOptions_, TwoStageSpecialization_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Xdl_CShuffle = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + BlockGemm_, + TransposeParams_, + GemmBatchOptions_, + TwoStageSpecialization_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle_V3 = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, BlockGemm_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle_V3 = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + BlockGemm_>; using ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_>; + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, MultipleDSpecialization_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + MultipleDSpecialization_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, BlockGemm_, TransposeParams_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + BlockGemm_, + TransposeParams_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, BlockGemm_, TransposeParams_, GemmBatchOptions_, TwoStageSpecialization_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + BlockGemm_, + TransposeParams_, + GemmBatchOptions_, + TwoStageSpecialization_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, GridGemm_, Prefetch_>; - -using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 = - ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, BlockGemm_, MultipleDSpecialization_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + GridGemm_, + Prefetch_>; +using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 = + ConvAlgorithmTemplate, + ConvSpecializationBwdWeight_, + BlockGemm_, + MultipleDSpecialization_>; } // namespace ck_tile::builder::test diff --git a/experimental/builder/test/test_concept_diagnostics_sync.cpp b/experimental/builder/test/test_concept_diagnostics_sync.cpp index 0bc786dbdf..ca08ae92ed 100644 --- a/experimental/builder/test/test_concept_diagnostics_sync.cpp +++ b/experimental/builder/test/test_concept_diagnostics_sync.cpp @@ -22,40 +22,40 @@ namespace ck_tile::builder::test { -using ck_tile::builder::ThreadBlockDescriptor; -using ck_tile::builder::GridwiseXdlGemmDescriptor; -using ck_tile::builder::BlockTransferDescriptor; -using ck_tile::builder::ThreadClusterDescriptor; -using ck_tile::builder::LdsTransferDescriptor; -using ck_tile::builder::EpilogueDescriptor; using ck_tile::builder::AccessOrderDescriptor; using ck_tile::builder::BlockGemmDescriptor; -using ck_tile::builder::GridwiseWmmaGemmDescriptor; -using ck_tile::builder::TileThreadBlockDescriptor; -using ck_tile::builder::TileTransferDescriptor; -using ck_tile::builder::TileBlockGemmDescriptor; -using ck_tile::builder::TileOptimizationsDescriptor; -using ck_tile::builder::DlThreadConfigDescriptor; -using ck_tile::builder::DlThreadClusterDescriptor; +using ck_tile::builder::BlockTransferDescriptor; +using ck_tile::builder::ConvAlgorithmDescriptor; using ck_tile::builder::DlBlockTransferDescriptor; using ck_tile::builder::DlEpilogueDescriptor; -using ck_tile::builder::ConvAlgorithmDescriptor; -using ck_tile::builder::SpecifiesThreadBlock; -using ck_tile::builder::SpecifiesGridwiseFwdXdlGemm; -using ck_tile::builder::SpecifiesGridwiseBwdXdlGemm; +using ck_tile::builder::DlThreadClusterDescriptor; +using ck_tile::builder::DlThreadConfigDescriptor; +using ck_tile::builder::EpilogueDescriptor; +using ck_tile::builder::GridwiseWmmaGemmDescriptor; +using ck_tile::builder::GridwiseXdlGemmDescriptor; +using ck_tile::builder::LdsTransferDescriptor; using ck_tile::builder::SpecifiesBlockGemm; -using ck_tile::builder::SpecifiesFwdConvSpecialization; using ck_tile::builder::SpecifiesBwdWeightConvSpecialization; +using ck_tile::builder::SpecifiesDlThreadCluster; +using ck_tile::builder::SpecifiesDlThreadConfig; +using ck_tile::builder::SpecifiesFwdConvSpecialization; using ck_tile::builder::SpecifiesGemmSpecialization; -using ck_tile::builder::SpecifiesNumPrefetchStages; +using ck_tile::builder::SpecifiesGridwiseBwdXdlGemm; +using ck_tile::builder::SpecifiesGridwiseFwdXdlGemm; using ck_tile::builder::SpecifiesLoopScheduler; +using ck_tile::builder::SpecifiesNumPrefetchStages; +using ck_tile::builder::SpecifiesThreadBlock; +using ck_tile::builder::SpecifiesTileBlockGemm; +using ck_tile::builder::SpecifiesTileConvSpecialization; +using ck_tile::builder::SpecifiesTileOptimizations; using ck_tile::builder::SpecifiesTileThreadBlock; using ck_tile::builder::SpecifiesTileTransfer; -using ck_tile::builder::SpecifiesTileBlockGemm; -using ck_tile::builder::SpecifiesTileOptimizations; -using ck_tile::builder::SpecifiesTileConvSpecialization; -using ck_tile::builder::SpecifiesDlThreadConfig; -using ck_tile::builder::SpecifiesDlThreadCluster; +using ck_tile::builder::ThreadBlockDescriptor; +using ck_tile::builder::ThreadClusterDescriptor; +using ck_tile::builder::TileBlockGemmDescriptor; +using ck_tile::builder::TileOptimizationsDescriptor; +using ck_tile::builder::TileThreadBlockDescriptor; +using ck_tile::builder::TileTransferDescriptor; // Helper to check if a string contains a substring bool contains(const std::string& str, const std::string& substr) @@ -331,18 +331,24 @@ TEST(ConceptDiagnosticsSync, LdsTransferDescriptor_Invalid) TEST(ConceptDiagnosticsSync, CompleteAlgorithmTypes) { // Test that complete algorithm types satisfy their concepts - static_assert(ConvAlgorithmDescriptor); - static_assert(ConvAlgorithmDescriptor); - static_assert(ConvAlgorithmDescriptor); + static_assert( + ConvAlgorithmDescriptor); + static_assert( + ConvAlgorithmDescriptor); + static_assert( + ConvAlgorithmDescriptor); static_assert(ConvAlgorithmDescriptor); static_assert(ConvAlgorithmDescriptor); - + // Test specific requirements for each algorithm type static_assert(SpecifiesThreadBlock); - static_assert(SpecifiesGridwiseFwdXdlGemm); - static_assert(SpecifiesFwdConvSpecialization); - static_assert(SpecifiesNumPrefetchStages); - + static_assert( + SpecifiesGridwiseFwdXdlGemm); + static_assert( + SpecifiesFwdConvSpecialization); + static_assert( + SpecifiesNumPrefetchStages); + static_assert(SpecifiesTileThreadBlock); static_assert(SpecifiesTileBlockGemm); static_assert(SpecifiesTileOptimizations); @@ -356,9 +362,12 @@ TEST(ConceptDiagnosticsSync, DiagnosticMessages) { // Test that diagnostics can be called (even if messages may be empty at compile-time) // The key is that the diagnostic functions exist and compile - std::string diag1 = ck_tile::builder::diagnostics::detailed_diagnostic_SpecifiesThreadBlock(); - std::string diag2 = ck_tile::builder::diagnostics::detailed_diagnostic_SpecifiesGridwiseFwdXdlGemm(); - + std::string diag1 = ck_tile::builder::diagnostics::detailed_diagnostic_SpecifiesThreadBlock< + invalid_types::MissingBlockSize>(); + std::string diag2 = + ck_tile::builder::diagnostics::detailed_diagnostic_SpecifiesGridwiseFwdXdlGemm< + invalid_types::MissingMPerXdl>(); + // These may be empty depending on the implementation, but they should compile EXPECT_TRUE(diag1.empty() || contains(diag1, "thread_block") || contains(diag1, "missing")); EXPECT_TRUE(diag2.empty() || contains(diag2, "gridwise_gemm") || contains(diag2, "missing")); @@ -370,7 +379,7 @@ TEST(ConceptDiagnosticsSync, DiagnosticMessages) /** * @brief Verify that all concepts defined in conv_algorithm_concepts.hpp have tests - * + * * This test serves as documentation of which concepts are tested. If new concepts * are added, this test should be updated to include them. */ @@ -386,13 +395,13 @@ TEST(ConceptDiagnosticsSync, ConceptCoverage) EXPECT_TRUE((LdsTransferDescriptor)); EXPECT_TRUE((EpilogueDescriptor)); EXPECT_TRUE((AccessOrderDescriptor)); - - // Tile Descriptor Concepts + + // Tile Descriptor Concepts EXPECT_TRUE((TileThreadBlockDescriptor)); EXPECT_TRUE((TileTransferDescriptor)); EXPECT_TRUE((TileBlockGemmDescriptor)); EXPECT_TRUE((TileOptimizationsDescriptor)); - + // DL Descriptor Concepts EXPECT_TRUE((DlThreadConfigDescriptor)); EXPECT_TRUE((DlThreadClusterDescriptor)); diff --git a/experimental/builder/test/test_conv_description.cpp b/experimental/builder/test/test_conv_description.cpp index 8808f9bee7..498de9a42f 100644 --- a/experimental/builder/test/test_conv_description.cpp +++ b/experimental/builder/test/test_conv_description.cpp @@ -120,14 +120,10 @@ struct DefaultAlgorithm ckb::test::ThreadBlock thread_block{.block_size = 256, .tile_size = {.m = 256, .n = 256, .k = 32}}; - ckb::test::GridwiseFwdXdlGemm gridwise_gemm{.ak1 = 8, - .bk1 = 8, - .xdl_params = - { - .m_per_xdl = 16, - .n_per_xdl = 16, - .m_xdl_per_wave = 8, - .n_xdl_per_wave = 8}}; + ckb::test::GridwiseFwdXdlGemm gridwise_gemm{ + .ak1 = 8, + .bk1 = 8, + .xdl_params = {.m_per_xdl = 16, .n_per_xdl = 16, .m_xdl_per_wave = 8, .n_xdl_per_wave = 8}}; ckb::test::Transfer<> transfer{ .a = @@ -163,8 +159,8 @@ struct DefaultAlgorithm }, }; - ckb::ConvSpecialization fwd_specialization = ckb::ConvSpecialization::DEFAULT; - ckb::GemmSpecialization gemm_specialization = ckb::GemmSpecialization::Default; + ckb::ConvSpecialization fwd_specialization = ckb::ConvSpecialization::DEFAULT; + ckb::GemmSpecialization gemm_specialization = ckb::GemmSpecialization::Default; ckb::test::BlockGemm block_gemm{.pipeline_version = ckb::PipelineVersion::V4, .scheduler = ckb::PipelineScheduler::INTRAWAVE}; }; diff --git a/experimental/builder/test/utils/ckb_conv_test_configs.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp index e1f5a34e20..3b83ead2d0 100644 --- a/experimental/builder/test/utils/ckb_conv_test_configs.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp @@ -20,37 +20,35 @@ constexpr DlThreadConfig DlThreadConfig_16x1x4x4x1{ constexpr DlThreadCluster DlThreadCluster_8x2{.m1_xs = {8, 2}, .n1_xs = {8, 2}}; -constexpr DlBlockTransfer<4> DlBlockTransfer_8x1x1x2 - {.thread_slice_lengths = {8, 1, 1, 2}, - .thread_cluster_lengths = {2, 1, 128, 1}, - .thread_cluster_arrange_order = {1, 2, 0, 3}, - .src_access_order = {1, 2, 0, 3}, - .src_vector_tensor_lengths = {4, 1, 1, 2}, - .src_vector_tensor_contiguous_dim_order = {1, 2, 0, 3}, - .dst_vector_tensor_lengths = {1, 1, 1, 2}}; +constexpr DlBlockTransfer<4> DlBlockTransfer_8x1x1x2{ + .thread_slice_lengths = {8, 1, 1, 2}, + .thread_cluster_lengths = {2, 1, 128, 1}, + .thread_cluster_arrange_order = {1, 2, 0, 3}, + .src_access_order = {1, 2, 0, 3}, + .src_vector_tensor_lengths = {4, 1, 1, 2}, + .src_vector_tensor_contiguous_dim_order = {1, 2, 0, 3}, + .dst_vector_tensor_lengths = {1, 1, 1, 2}}; -constexpr DlTransfer<4> DlTransfer4D {.a = DlBlockTransfer_8x1x1x2, - .b = DlBlockTransfer_8x1x1x2, - .c = { - .src_dst_access_order = {0, 1, 2, 3, 4, 5}, - .src_dst_vector_dim = 5, - .dst_scalar_per_vector = 4}}; +constexpr DlTransfer<4> DlTransfer4D{.a = DlBlockTransfer_8x1x1x2, + .b = DlBlockTransfer_8x1x1x2, + .c = {.src_dst_access_order = {0, 1, 2, 3, 4, 5}, + .src_dst_vector_dim = 5, + .dst_scalar_per_vector = 4}}; -constexpr DlBlockTransfer<5> DlBlockTransfer_1x8x1x1x1 - {.thread_slice_lengths = {1, 8, 1, 1, 1}, - .thread_cluster_lengths = {1, 2, 1, 128, 1}, - .thread_cluster_arrange_order = {0, 2, 3, 1, 4}, - .src_access_order = {0, 2, 3, 1, 4}, - .src_vector_tensor_lengths = {1, 1, 1, 1, 1}, - .src_vector_tensor_contiguous_dim_order = {0, 2, 3, 1, 4}, - .dst_vector_tensor_lengths = {1, 1, 1, 1, 1}}; +constexpr DlBlockTransfer<5> DlBlockTransfer_1x8x1x1x1{ + .thread_slice_lengths = {1, 8, 1, 1, 1}, + .thread_cluster_lengths = {1, 2, 1, 128, 1}, + .thread_cluster_arrange_order = {0, 2, 3, 1, 4}, + .src_access_order = {0, 2, 3, 1, 4}, + .src_vector_tensor_lengths = {1, 1, 1, 1, 1}, + .src_vector_tensor_contiguous_dim_order = {0, 2, 3, 1, 4}, + .dst_vector_tensor_lengths = {1, 1, 1, 1, 1}}; -constexpr DlTransfer<5> DlTransfer5D {.a = DlBlockTransfer_1x8x1x1x1, - .b = DlBlockTransfer_1x8x1x1x1, - .c = { - .src_dst_access_order = {0, 1, 2, 3, 4, 5}, - .src_dst_vector_dim = 5, - .dst_scalar_per_vector = 1}}; +constexpr DlTransfer<5> DlTransfer5D{.a = DlBlockTransfer_1x8x1x1x1, + .b = DlBlockTransfer_1x8x1x1x1, + .c = {.src_dst_access_order = {0, 1, 2, 3, 4, 5}, + .src_dst_vector_dim = 5, + .dst_scalar_per_vector = 1}}; constexpr Transfer<> Transfer_4x64x1{ .a = @@ -252,40 +250,38 @@ constexpr Transfer<> Transfer_4x32x1{ }; constexpr GridwiseBwdXdlGemm BwdGemmParams_Xdl_4x4_per_wave{ - .k1 = 8, + .k1 = 8, .xdl_params = {.m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4}}; constexpr GridwiseBwdXdlGemm BwdGemmParams_Xdl_1x1_per_wave{ - .k1 = 8, + .k1 = 8, .xdl_params = {.m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 1, .n_xdl_per_wave = 1}}; constexpr GridwiseFwdXdlGemm FwdGemmParams_Xdl_4x4_per_wave{ - .ak1 = 8, .bk1 = 8, + .ak1 = 8, + .bk1 = 8, .xdl_params = {.m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4}}; constexpr GridwiseFwdXdlGemm FwdGemmParams_Xdl_4x2_per_wave{ - .ak1 = 8, .bk1 = 8, + .ak1 = 8, + .bk1 = 8, .xdl_params = {.m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 4, .n_xdl_per_wave = 2}}; constexpr GridwiseFwdXdlGemm FwdGemmParams_Xdl_2x2_per_wave{ - .ak1 = 8, .bk1 = 8, + .ak1 = 8, + .bk1 = 8, .xdl_params = {.m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 2, .n_xdl_per_wave = 2}}; constexpr GridwiseFwdXdlGemm FwdGemmParams_Xdl_2x1_per_wave{ - .ak1 = 8, .bk1 = 8, + .ak1 = 8, + .bk1 = 8, .xdl_params = {.m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 2, .n_xdl_per_wave = 1}}; -constexpr GridwiseWmmaGemm GemmParams_Wmma_2x1_per_wave{.k1 = 8, - .m_per_wmma = 32, - .n_per_wmma = 32, - .m_wmma_per_wave = 2, - .n_wmma_per_wave = 1}; +constexpr GridwiseWmmaGemm GemmParams_Wmma_2x1_per_wave{ + .k1 = 8, .m_per_wmma = 32, .n_per_wmma = 32, .m_wmma_per_wave = 2, .n_wmma_per_wave = 1}; -constexpr GridwiseWmmaGemm GemmParams_Wmma_16x16_2x1_per_wave{.k1 = 8, - .m_per_wmma = 16, - .n_per_wmma = 16, - .m_wmma_per_wave = 2, - .n_wmma_per_wave = 1}; +constexpr GridwiseWmmaGemm GemmParams_Wmma_16x16_2x1_per_wave{ + .k1 = 8, .m_per_wmma = 16, .n_per_wmma = 16, .m_wmma_per_wave = 2, .n_wmma_per_wave = 1}; constexpr ThreadBlock ThreadBlock_256_256x256x32{.block_size = 256, .tile_size = {.m = 256, .n = 256, .k = 32}}; @@ -300,13 +296,13 @@ constexpr ThreadBlock ThreadBlock_256_128x128x16{.block_size = 256, .tile_size = {.m = 128, .n = 128, .k = 16}}; constexpr ThreadBlock ThreadBlock_256_128x128x8{.block_size = 256, - .tile_size = {.m = 128, .n = 128, .k = 8}}; + .tile_size = {.m = 128, .n = 128, .k = 8}}; constexpr ThreadBlock ThreadBlock_64_64x32x32{.block_size = 64, .tile_size = {.m = 64, .n = 32, .k = 32}}; constexpr ThreadBlock ThreadBlock_64_32x32x32{.block_size = 64, - .tile_size = {.m = 32, .n = 32, .k = 32}}; + .tile_size = {.m = 32, .n = 32, .k = 32}}; constexpr ThreadBlock ThreadBlock_128_128x128x32{.block_size = 128, .tile_size = {.m = 128, .n = 128, .k = 32}}; @@ -314,19 +310,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 BlockGemmPipeline BlockGemmDesc_v1_intrawave = {.pipeline_version = PipelineVersion::V1, - .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr BlockGemmPipeline BlockGemmDesc_v1_intrawave = { + .pipeline_version = PipelineVersion::V1, .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr BlockGemmPipeline BlockGemmDesc_v2_intrawave = {.pipeline_version = PipelineVersion::V2, - .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr BlockGemmPipeline BlockGemmDesc_v2_intrawave = { + .pipeline_version = PipelineVersion::V2, .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr BlockGemmPipeline BlockGemmDesc_v3_intrawave = {.pipeline_version = PipelineVersion::V3, - .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr BlockGemmPipeline BlockGemmDesc_v3_intrawave = { + .pipeline_version = PipelineVersion::V3, .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr BlockGemmPipeline BlockGemmDesc_v4_intrawave = {.pipeline_version = PipelineVersion::V4, - .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr BlockGemmPipeline BlockGemmDesc_v4_intrawave = { + .pipeline_version = PipelineVersion::V4, .scheduler = PipelineScheduler::INTRAWAVE}; -constexpr BlockGemmPipeline BlockGemmDesc_v5_intrawave = {.pipeline_version = PipelineVersion::V5, - .scheduler = PipelineScheduler::INTRAWAVE}; +constexpr BlockGemmPipeline 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 6ad6179780..b705ab1e47 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -98,8 +98,8 @@ template <> inline std::string to_string(GridwiseFwdXdlGemm t) { std::ostringstream oss; - oss << t.ak1 << "," << t.bk1 << "," << t.xdl_params.m_per_xdl << "," << t.xdl_params.n_per_xdl << "," - << t.xdl_params.m_xdl_per_wave << "," << t.xdl_params.n_xdl_per_wave; + oss << t.ak1 << "," << t.bk1 << "," << t.xdl_params.m_per_xdl << "," << t.xdl_params.n_per_xdl + << "," << t.xdl_params.m_xdl_per_wave << "," << t.xdl_params.n_xdl_per_wave; return oss.str(); } @@ -123,15 +123,15 @@ inline std::string to_string(BlockGemmPipeline t) template inline std::string to_string(BlockTransfer t) { - if constexpr (ThreadSliceDim == 4) + if constexpr(ThreadSliceDim == 4) { return array_to_seq(std::array{t.k_batch_size, t.k0, t.m_n, t.k1}); } - else if constexpr (ThreadSliceDim == 3) + else if constexpr(ThreadSliceDim == 3) { return array_to_seq(std::array{t.k0, t.m_n, t.k1}); } - else + else { static_assert(ThreadSliceDim == 3 || ThreadSliceDim == 4, "Unsupported ThreadSliceDim"); } @@ -160,7 +160,7 @@ inline std::string to_string(AccessOrder t) return array_to_seq(t.order); } -template +template inline std::string to_string(InputTransfer t) { std::ostringstream oss; @@ -314,8 +314,7 @@ template <> inline std::string to_string(Prefetch_ t) { std::ostringstream oss; - oss << t.num_gemm_k_prefetch_stages << "," - << to_string(t.loop_scheduler); + oss << t.num_gemm_k_prefetch_stages << "," << to_string(t.loop_scheduler); return oss.str(); } diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_wmma_cshuffle_v3.hpp index dbce8e8ccf..80e4ea983e 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_wmma_cshuffle_v3.hpp @@ -861,30 +861,32 @@ struct DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3 { if(gemm_arg.KBatch > 1) { - const auto kernel = kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< - GridwiseGemm, - remove_reference_t, - remove_reference_t, - remove_reference_t< - DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, - ComputePtrOffsetOfStridedBatch, - true, - InMemoryDataOperationEnum::AtomicAdd, - minimum_occupancy>; + const auto kernel = + kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< + GridwiseGemm, + remove_reference_t, + remove_reference_t, + remove_reference_t< + DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, + ComputePtrOffsetOfStridedBatch, + true, + InMemoryDataOperationEnum::AtomicAdd, + minimum_occupancy>; Run(kernel); } else { - const auto kernel = kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< - GridwiseGemm, - remove_reference_t, - remove_reference_t, - remove_reference_t< - DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, - ComputePtrOffsetOfStridedBatch, - true, - InMemoryDataOperationEnum::Set, - minimum_occupancy>; + const auto kernel = + kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< + GridwiseGemm, + remove_reference_t, + remove_reference_t, + remove_reference_t< + DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, + ComputePtrOffsetOfStridedBatch, + true, + InMemoryDataOperationEnum::Set, + minimum_occupancy>; Run(kernel); } } @@ -900,30 +902,32 @@ struct DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3 { if(gemm_arg.KBatch > 1) { - const auto kernel = kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< - GridwiseGemm, - remove_reference_t, - remove_reference_t, - remove_reference_t< - DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, - ComputePtrOffsetOfStridedBatch, - false, - InMemoryDataOperationEnum::AtomicAdd, - minimum_occupancy>; + const auto kernel = + kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< + GridwiseGemm, + remove_reference_t, + remove_reference_t, + remove_reference_t< + DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, + ComputePtrOffsetOfStridedBatch, + false, + InMemoryDataOperationEnum::AtomicAdd, + minimum_occupancy>; Run(kernel); } else { - const auto kernel = kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< - GridwiseGemm, - remove_reference_t, - remove_reference_t, - remove_reference_t< - DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, - ComputePtrOffsetOfStridedBatch, - false, - InMemoryDataOperationEnum::Set, - minimum_occupancy>; + const auto kernel = + kernel_grouped_conv_bwd_weight_wmma_cshuffle_v3_multiple_d< + GridwiseGemm, + remove_reference_t, + remove_reference_t, + remove_reference_t< + DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>, + ComputePtrOffsetOfStridedBatch, + false, + InMemoryDataOperationEnum::Set, + minimum_occupancy>; Run(kernel); } } diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp index 0910da35a6..d3bf2a364a 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp @@ -52,19 +52,20 @@ __global__ void #if CK_USE_LAUNCH_BOUNDS __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) #endif - kernel_batched_gemm_xdlops_bwd_weight_multiple_d(const FloatA* __restrict__ p_a_grid, - const FloatB* __restrict__ p_b_grid, - FloatC* __restrict__ p_c_grid, - const AElementwiseOperation a_element_op, - const BElementwiseOperation b_element_op, - const CElementwiseOperation c_element_op, - const index_t batch_count, - const AGridDesc_B_K0_M_K1 a_b_k0_m_k1_grid_desc, - const BGridDesc_B_K0_N_K1 b_b_k0_n_k1_grid_desc, - const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock - c_grid_desc_mblock_mperblock_nblock_nperblock, - const Block2CTileMap block_2_ctile_map, - const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) + kernel_batched_gemm_xdlops_bwd_weight_multiple_d( + const FloatA* __restrict__ p_a_grid, + const FloatB* __restrict__ p_b_grid, + FloatC* __restrict__ p_c_grid, + const AElementwiseOperation a_element_op, + const BElementwiseOperation b_element_op, + const CElementwiseOperation c_element_op, + const index_t batch_count, + const AGridDesc_B_K0_M_K1 a_b_k0_m_k1_grid_desc, + const BGridDesc_B_K0_N_K1 b_b_k0_n_k1_grid_desc, + const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock + c_grid_desc_mblock_mperblock_nblock_nperblock, + const Block2CTileMap block_2_ctile_map, + const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { #if defined(__gfx9__) || defined(__gfx11__) || defined(__gfx12__) if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())