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 b3c3a9a59f..120590e2d7 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -315,24 +315,6 @@ concept SpecifiesLoopScheduler = requires { { T::loop_scheduler } -> std::convertible_to; }; -template -concept SpecifiesLargeTensorSupport = requires { - { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR; -}; - -template -concept SpecifiesTwoStageSupport = requires { - { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::TWO_STAGE; -}; - -template -concept SpecifiesMultipleDSupport = requires { - { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D; -}; - template concept SpecifiesGenericInstance = !requires { { T::specialization }; @@ -359,6 +341,33 @@ concept SpecifiesGemmBatchOptions = requires { { T::num_conv_groups_to_merge } -> SizeType; }; +/******************************************** */ +/* Algorithm specialization concepts */ +/******************************************** */ +template +concept SpecifiesLargeTensorSupport = requires { + { T::specialization } -> std::convertible_to; + requires T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR; +}; + +template +concept SpecifiesReferenceAlgorithm = requires { + { T::specialization } -> std::convertible_to; + requires T::specialization == ConvAlgorithmSpecialization::REFERENCE; +}; + +template +concept SpecifiesTwoStageSupport = requires { + { T::specialization } -> std::convertible_to; + requires T::specialization == ConvAlgorithmSpecialization::TWO_STAGE; +}; + +template +concept SpecifiesMultipleDSupport = requires { + { T::specialization } -> std::convertible_to; + requires T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D; +}; + /******************************************** */ /* DL-specific descriptors and requirements */ /******************************************** */ 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 d979b77809..1f7bf6565d 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp @@ -712,6 +712,26 @@ consteval auto detailed_diagnostic_SpecifiesLargeTensorSupport() -> std::string return msg; } +template +consteval auto detailed_diagnostic_SpecifiesReferenceAlgorithm() -> std::string { + std::string msg; + if constexpr (requires { T::specialization; }) { + using SpecType = decltype(T::specialization); + constexpr bool convertible = std::convertible_to; + 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"; + } + } else { + msg += " → T::specialization: [✗] (missing member)\n"; + } + + return msg; +} + template consteval auto detailed_diagnostic_SpecifiesTwoStageSupport() -> std::string { std::string msg; 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 3f237ae744..15c383fb7a 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -9,6 +9,26 @@ namespace ck_tile::builder::factory { using namespace ck_tile::builder::diagnostics; +template +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 auto message() -> std::string { + return std::string("\n=== Reference Algorithm Diagnostic (closest match) ===\n" + "Concepts for Reference Algorithm:\n") + + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + + DIAGNOSTIC_LINE(SpecifiesReferenceAlgorithm); + } +}; + template struct FwdXdlV3Algorithm { CHECK_CONCEPT(T, ConvAlgorithmDescriptor)