From 4d20cc6b4df097af54d3faef1f7f12ceb2814ba4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Mon, 22 Dec 2025 07:36:13 -0500 Subject: [PATCH] Use amcro to ensure automatic macthing between concepts are their string representations. --- .../builder/factory/conv_algorithms.hpp | 327 +++++++++++------- 1 file changed, 200 insertions(+), 127 deletions(-) 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 a192a34df1..f96fc2c86b 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -9,162 +9,223 @@ namespace ck_tile::builder::factory { #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; \ + static constexpr const char* s_##Concept = #Concept; + +// Helper to create diagnostic message line +#define DIAGNOSTIC_LINE(Concept) \ + " " + std::string(s_##Concept) + ": " + std::string(CHECK_MARK(c_##Concept)) + "\n" + template struct FwdXdlV3Algorithm { - static constexpr bool c1 = ConvAlgorithmDescriptor; - static constexpr bool c2 = SpecifiesThreadBlock; - static constexpr bool c3 = SpecifiesBlockTransfer; - static constexpr bool c4 = SpecifiesLdsTransfer; - static constexpr bool c5 = SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = SpecifiesSourceAccessOrder; - static constexpr bool c7 = SpecifiesGridwiseFwdXdlGemm; - static constexpr bool c8 = SpecifiesFwdConvSpecialization; - static constexpr bool c9 = SpecifiesGemmSpecialization; - static constexpr bool c10 = SpecifiesBlockGemm; + CHECK_CONCEPT(T, ConvAlgorithmDescriptor) + CHECK_CONCEPT(T, SpecifiesThreadBlock) + CHECK_CONCEPT(T, SpecifiesBlockTransfer) + CHECK_CONCEPT(T, SpecifiesLdsTransfer) + CHECK_CONCEPT(T, SpecifiesThreadClusterAccessOrder) + CHECK_CONCEPT(T, SpecifiesSourceAccessOrder) + CHECK_CONCEPT(T, SpecifiesGridwiseFwdXdlGemm) + CHECK_CONCEPT(T, SpecifiesFwdConvSpecialization) + 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 c10 = c_SpecifiesBlockGemm; static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10; } static consteval const std::string message() { - return "\n=== Forward XDL V3 Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdXdlV3 Algorithm:\n" - " ConvAlgorithmDescriptor: " + std::string(CHECK_MARK(c1)) + "\n" - " SpecifiesThreadBlock: " + std::string(CHECK_MARK(c2)) + "\n" - " SpecifiesBlockTransfer: " + std::string(CHECK_MARK(c3)) + "\n" - " SpecifiesLdsTransfer: " + std::string(CHECK_MARK(c4)) + "\n" - " SpecifiesThreadClusterAccessOrder: " + std::string(CHECK_MARK(c5)) + "\n" - " SpecifiesSourceAccessOrder: " + std::string(CHECK_MARK(c6)) + "\n" - " SpecifiesGridwiseFwdXdlGemm: " + std::string(CHECK_MARK(c7)) + "\n" - " SpecifiesFwdConvSpecialization: " + std::string(CHECK_MARK(c8)) + "\n" - " SpecifiesGemmSpecialization: " + std::string(CHECK_MARK(c9)) + "\n" - " SpecifiesBlockGemm: " + std::string(CHECK_MARK(c10)) + "\n"; + 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); } }; template struct FwdXdlAlgorithm { - static constexpr bool c1 = ConvAlgorithmDescriptor; - static constexpr bool c2 = SpecifiesThreadBlock; - static constexpr bool c3 = SpecifiesBlockTransfer; - static constexpr bool c4 = SpecifiesLdsTransfer; - static constexpr bool c5 = SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = SpecifiesSourceAccessOrder; - static constexpr bool c7 = SpecifiesGridwiseFwdXdlGemm; - static constexpr bool c8 = SpecifiesFwdConvSpecialization; - static constexpr bool c9 = SpecifiesGemmSpecialization; - static constexpr bool c10 = SpecifiesNumPrefetchStages; - static constexpr bool c11 = SpecifiesNumGroupsToMerge; - static constexpr bool c12 = SpecifiesLoopScheduler; + CHECK_CONCEPT(T, ConvAlgorithmDescriptor) + CHECK_CONCEPT(T, SpecifiesThreadBlock) + CHECK_CONCEPT(T, SpecifiesBlockTransfer) + CHECK_CONCEPT(T, SpecifiesLdsTransfer) + CHECK_CONCEPT(T, SpecifiesThreadClusterAccessOrder) + CHECK_CONCEPT(T, SpecifiesSourceAccessOrder) + CHECK_CONCEPT(T, SpecifiesGridwiseFwdXdlGemm) + CHECK_CONCEPT(T, SpecifiesFwdConvSpecialization) + CHECK_CONCEPT(T, SpecifiesGemmSpecialization) + CHECK_CONCEPT(T, SpecifiesNumPrefetchStages) + 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 c10 = c_SpecifiesNumPrefetchStages; + static constexpr bool c11 = c_SpecifiesNumGroupsToMerge; + static constexpr bool c12 = c_SpecifiesLoopScheduler; static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10 && c11 && c12; } static consteval const std::string message() { - return "\n=== Forward XDL Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdXdl Algorithm:\n" - " ConvAlgorithmDescriptor: " + std::string(CHECK_MARK(c1)) + "\n" - " SpecifiesThreadBlock: " + std::string(CHECK_MARK(c2)) + "\n" - " SpecifiesBlockTransfer: " + std::string(CHECK_MARK(c3)) + "\n" - " SpecifiesLdsTransfer: " + std::string(CHECK_MARK(c4)) + "\n" - " SpecifiesThreadClusterAccessOrder: " + std::string(CHECK_MARK(c5)) + "\n" - " SpecifiesSourceAccessOrder: " + std::string(CHECK_MARK(c6)) + "\n" - " SpecifiesGridwiseFwdXdlGemm: " + std::string(CHECK_MARK(c7)) + "\n" - " SpecifiesFwdConvSpecialization: " + std::string(CHECK_MARK(c8)) + "\n" - " SpecifiesGemmSpecialization: " + std::string(CHECK_MARK(c9)) + "\n" - " SpecifiesNumPrefetchStages: " + std::string(CHECK_MARK(c10)) + "\n" - " SpecifiesNumGroupsToMerge: " + std::string(CHECK_MARK(c11)) + "\n" - " SpecifiesLoopScheduler: " + std::string(CHECK_MARK(c12)) + "\n"; + return std::string("\n=== Forward XDL Algorithm Diagnostic (closest match) ===\n" + "Concepts for FwdXdl 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(SpecifiesNumPrefetchStages) + + DIAGNOSTIC_LINE(SpecifiesNumGroupsToMerge) + + DIAGNOSTIC_LINE(SpecifiesLoopScheduler); } }; template struct FwdWmmaAlgorithm { - static constexpr bool c1 = ConvAlgorithmDescriptor; - static constexpr bool c2 = SpecifiesThreadBlock; - static constexpr bool c3 = SpecifiesBlockTransfer; - static constexpr bool c4 = SpecifiesLdsTransfer; - static constexpr bool c5 = SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = SpecifiesSourceAccessOrder; - static constexpr bool c7 = SpecifiesGridwiseWmmaGemm; - static constexpr bool c8 = SpecifiesFwdConvSpecialization; - static constexpr bool c9 = SpecifiesGemmSpecialization; - static constexpr bool c10 = SpecifiesNumPrefetchStages; - static constexpr bool c11 = SpecifiesLoopScheduler; + CHECK_CONCEPT(T, ConvAlgorithmDescriptor) + CHECK_CONCEPT(T, SpecifiesThreadBlock) + CHECK_CONCEPT(T, SpecifiesBlockTransfer) + CHECK_CONCEPT(T, SpecifiesLdsTransfer) + CHECK_CONCEPT(T, SpecifiesThreadClusterAccessOrder) + CHECK_CONCEPT(T, SpecifiesSourceAccessOrder) + CHECK_CONCEPT(T, SpecifiesGridwiseWmmaGemm) + CHECK_CONCEPT(T, SpecifiesFwdConvSpecialization) + CHECK_CONCEPT(T, SpecifiesGemmSpecialization) + CHECK_CONCEPT(T, SpecifiesNumPrefetchStages) + 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_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 consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10 && c11; } static consteval const std::string message() { - return "\n=== Forward WMMA Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdWmma Algorithm:\n" - " ConvAlgorithmDescriptor: " + std::string(CHECK_MARK(c1)) + "\n" - " SpecifiesThreadBlock: " + std::string(CHECK_MARK(c2)) + "\n" - " SpecifiesBlockTransfer: " + std::string(CHECK_MARK(c3)) + "\n" - " SpecifiesLdsTransfer: " + std::string(CHECK_MARK(c4)) + "\n" - " SpecifiesThreadClusterAccessOrder: " + std::string(CHECK_MARK(c5)) + "\n" - " SpecifiesSourceAccessOrder: " + std::string(CHECK_MARK(c6)) + "\n" - " SpecifiesGridwiseWmmaGemm: " + std::string(CHECK_MARK(c7)) + "\n" - " SpecifiesFwdConvSpecialization: " + std::string(CHECK_MARK(c8)) + "\n" - " SpecifiesGemmSpecialization: " + std::string(CHECK_MARK(c9)) + "\n" - " SpecifiesNumPrefetchStages: " + std::string(CHECK_MARK(c10)) + "\n" - " SpecifiesLoopScheduler: " + std::string(CHECK_MARK(c11)) + "\n"; + 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) + + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + + DIAGNOSTIC_LINE(SpecifiesGridwiseWmmaGemm) + + DIAGNOSTIC_LINE(SpecifiesFwdConvSpecialization) + + DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + + DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + + DIAGNOSTIC_LINE(SpecifiesLoopScheduler); } }; template struct FwdDlAlgorithm { - static constexpr bool c1 = ConvAlgorithmDescriptor; - static constexpr bool c2 = SpecifiesThreadBlock; - static constexpr bool c3 = SpecifiesFwdConvSpecialization; - static constexpr bool c4 = SpecifiesGemmSpecialization; - static constexpr bool c5 = SpecifiesDlThreadConfig; - static constexpr bool c6 = SpecifiesDlThreadCluster; - static constexpr bool c7 = SpecifiesDlBlockTransfer; - static constexpr bool c8 = SpecifiesDlEpilogue; + CHECK_CONCEPT(T, ConvAlgorithmDescriptor) + CHECK_CONCEPT(T, SpecifiesThreadBlock) + CHECK_CONCEPT(T, SpecifiesFwdConvSpecialization) + CHECK_CONCEPT(T, SpecifiesGemmSpecialization) + CHECK_CONCEPT(T, SpecifiesDlThreadConfig) + CHECK_CONCEPT(T, SpecifiesDlThreadCluster) + CHECK_CONCEPT(T, SpecifiesDlBlockTransfer) + CHECK_CONCEPT(T, SpecifiesDlEpilogue) + + static constexpr bool c1 = c_ConvAlgorithmDescriptor; + static constexpr bool c2 = c_SpecifiesThreadBlock; + static constexpr bool c3 = c_SpecifiesFwdConvSpecialization; + static constexpr bool c4 = c_SpecifiesGemmSpecialization; + static constexpr bool c5 = c_SpecifiesDlThreadConfig; + static constexpr bool c6 = c_SpecifiesDlThreadCluster; + static constexpr bool c7 = c_SpecifiesDlBlockTransfer; + static constexpr bool c8 = c_SpecifiesDlEpilogue; static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8; } static consteval const std::string message() { - return "\n=== Forward DL Algorithm Diagnostic (closest match) ===\n" - "Concepts for FwdDl Algorithm:\n" - " ConvAlgorithmDescriptor: " + std::string(CHECK_MARK(c1)) + "\n" - " SpecifiesThreadBlock: " + std::string(CHECK_MARK(c2)) + "\n" - " SpecifiesFwdConvSpecialization: " + std::string(CHECK_MARK(c3)) + "\n" - " SpecifiesGemmSpecialization: " + std::string(CHECK_MARK(c4)) + "\n" - " SpecifiesDlThreadConfig: " + std::string(CHECK_MARK(c5)) + "\n" - " SpecifiesDlThreadCluster: " + std::string(CHECK_MARK(c6)) + "\n" - " SpecifiesDlBlockTransfer: " + std::string(CHECK_MARK(c7)) + "\n" - " SpecifiesDlEpilogue: " + std::string(CHECK_MARK(c8)) + "\n"; + return std::string("\n=== Forward DL Algorithm Diagnostic (closest match) ===\n" + "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(SpecifiesDlBlockTransfer) + + DIAGNOSTIC_LINE(SpecifiesDlEpilogue); } }; template struct TileAlgorithm { - static constexpr bool c1 = ConvAlgorithmDescriptor; - static constexpr bool c2 = SpecifiesTileThreadBlock; - static constexpr bool c3 = SpecifiesTileTransfer; - static constexpr bool c4 = SpecifiesTileConvSpecialization; - static constexpr bool c5 = SpecifiesTileBlockGemm; - static constexpr bool c6 = SpecifiesTileOptimizations; + CHECK_CONCEPT(T, ConvAlgorithmDescriptor) + CHECK_CONCEPT(T, SpecifiesTileThreadBlock) + CHECK_CONCEPT(T, SpecifiesTileTransfer) + CHECK_CONCEPT(T, SpecifiesTileConvSpecialization) + CHECK_CONCEPT(T, SpecifiesTileBlockGemm) + CHECK_CONCEPT(T, SpecifiesTileOptimizations) + + static constexpr bool c1 = c_ConvAlgorithmDescriptor; + static constexpr bool c2 = c_SpecifiesTileThreadBlock; + static constexpr bool c3 = c_SpecifiesTileTransfer; + static constexpr bool c4 = c_SpecifiesTileConvSpecialization; + 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 const std::string message() { - return "\n=== CK Tile Algorithm Diagnostic (closest match) ===\n" - "Concepts for CK Tile Conv Algorithm:\n" - " ConvAlgorithmDescriptor: " + std::string(CHECK_MARK(c1)) + "\n" - " SpecifiesTileThreadBlock: " + std::string(CHECK_MARK(c2)) + "\n" - " SpecifiesTileTransfer: " + std::string(CHECK_MARK(c3)) + "\n" - " SpecifiesTileConvSpecialization: " + std::string(CHECK_MARK(c4)) + "\n" - " SpecifiesTileBlockGemm: " + std::string(CHECK_MARK(c5)) + "\n" - " SpecifiesTileOptimizations: " + std::string(CHECK_MARK(c6)) + "\n"; + 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(SpecifiesTileConvSpecialization) + + DIAGNOSTIC_LINE(SpecifiesTileBlockGemm) + + DIAGNOSTIC_LINE(SpecifiesTileOptimizations); } }; @@ -172,7 +233,9 @@ template struct LargeTensorAlgorithm : public FwdXdlAlgorithm { using BaseAlgorithmType = decltype(T::base_algorithm); - static constexpr bool c13 = SpecifiesLargeTensorSupport; + CHECK_CONCEPT(T, SpecifiesLargeTensorSupport) + + static constexpr bool c13 = c_SpecifiesLargeTensorSupport; static consteval bool is_valid() { return FwdXdlAlgorithm::is_valid() && c13; @@ -180,38 +243,48 @@ struct LargeTensorAlgorithm : public FwdXdlAlgorithm::message() + - " SpecifiesLargeTensorSupport: " + std::string(CHECK_MARK(c13)) + "\n"; + DIAGNOSTIC_LINE(SpecifiesLargeTensorSupport); } }; template struct BwdXdlAlgorithm { - static constexpr bool c1 = ConvAlgorithmDescriptor; - static constexpr bool c2 = SpecifiesThreadBlock; - static constexpr bool c3 = SpecifiesBlockTransfer; - static constexpr bool c4 = SpecifiesLdsTransfer; - static constexpr bool c5 = SpecifiesThreadClusterAccessOrder; - static constexpr bool c6 = SpecifiesSourceAccessOrder; - static constexpr bool c7 = SpecifiesGridwiseBwdXdlGemm; - static constexpr bool c8 = SpecifiesBwdWeightConvSpecialization; - static constexpr bool c9 = SpecifiesTransposeTransfer; + CHECK_CONCEPT(T, ConvAlgorithmDescriptor) + CHECK_CONCEPT(T, SpecifiesThreadBlock) + CHECK_CONCEPT(T, SpecifiesBlockTransfer) + CHECK_CONCEPT(T, SpecifiesLdsTransfer) + CHECK_CONCEPT(T, SpecifiesThreadClusterAccessOrder) + CHECK_CONCEPT(T, SpecifiesSourceAccessOrder) + CHECK_CONCEPT(T, SpecifiesGridwiseBwdXdlGemm) + CHECK_CONCEPT(T, SpecifiesBwdWeightConvSpecialization) + CHECK_CONCEPT(T, SpecifiesTransposeTransfer) + + 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_SpecifiesGridwiseBwdXdlGemm; + static constexpr bool c8 = c_SpecifiesBwdWeightConvSpecialization; + static constexpr bool c9 = c_SpecifiesTransposeTransfer; static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9; } static consteval const std::string message() { - return "\n=== Backward XDL Algorithm Diagnostic (closest match) ===\n" - "Concepts for BwdXdl Algorithm:\n" - " ConvAlgorithmDescriptor: " + std::string(CHECK_MARK(c1)) + "\n" - " SpecifiesThreadBlock: " + std::string(CHECK_MARK(c2)) + "\n" - " SpecifiesBlockTransfer: " + std::string(CHECK_MARK(c3)) + "\n" - " SpecifiesLdsTransfer: " + std::string(CHECK_MARK(c4)) + "\n" - " SpecifiesThreadClusterAccessOrder: " + std::string(CHECK_MARK(c5)) + "\n" - " SpecifiesSourceAccessOrder: " + std::string(CHECK_MARK(c6)) + "\n" - " SpecifiesGridwiseBwdXdlGemm: " + std::string(CHECK_MARK(c7)) + "\n" - " SpecifiesBwdWeightConvSpecialization: " + std::string(CHECK_MARK(c8)) + "\n" - " SpecifiesTransposeTransfer: " + std::string(CHECK_MARK(c9)) + "\n"; + return std::string("\n=== Backward XDL Algorithm Diagnostic (closest match) ===\n" + "Concepts for BwdXdl Algorithm:\n") + + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + + DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) + + DIAGNOSTIC_LINE(SpecifiesGridwiseBwdXdlGemm) + + DIAGNOSTIC_LINE(SpecifiesBwdWeightConvSpecialization) + + DIAGNOSTIC_LINE(SpecifiesTransposeTransfer); } }; @@ -315,4 +388,4 @@ consteval void diagnose_bwd_weight_algorithm_signature() } } -} +}