ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-12 18:17:58 +00:00
Code Issues Packages Projects Releases Wiki Activity

Treat ref algorithm the same way as real algorithms in the dispatcher.

Browse Source
This commit is contained in:
Ville Pietilä
2026-01-02 02:28:05 -05:00
parent 5be1ed65fb
commit 09e188f2a8
3 changed files with 67 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

45
experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
View File

@@ -315,24 +315,6 @@ concept SpecifiesLoopScheduler = requires {
{ T::loop_scheduler } -> std::convertible_to<PipelineScheduler>;
};
template <typename T>
concept SpecifiesLargeTensorSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR;
};
template <typename T>
concept SpecifiesTwoStageSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::TWO_STAGE;
};
template <typename T>
concept SpecifiesMultipleDSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D;
};
template <typename T>
concept SpecifiesGenericInstance = !requires {
{ T::specialization };
@@ -359,6 +341,33 @@ concept SpecifiesGemmBatchOptions = requires {
{ T::num_conv_groups_to_merge } -> SizeType;
};
/******************************************** */
/* Algorithm specialization concepts */
/******************************************** */
template <typename T>
concept SpecifiesLargeTensorSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR;
};
template <typename T>
concept SpecifiesReferenceAlgorithm = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::REFERENCE;
};
template <typename T>
concept SpecifiesTwoStageSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::TWO_STAGE;
};
template <typename T>
concept SpecifiesMultipleDSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D;
};
/******************************************** */
/* DL-specific descriptors and requirements */
/******************************************** */

20
experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp
View File

@@ -712,6 +712,26 @@ consteval auto detailed_diagnostic_SpecifiesLargeTensorSupport() -> std::string
return msg;
}
template <typename T>
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<SpecType, ConvAlgorithmSpecialization>;
msg += " → T::specialization: " + std::string(CHECK_MARK(convertible)) +
(convertible ? "" : std::string(detail::get_type_info<SpecType>())) + "\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 <typename T>
consteval auto detailed_diagnostic_SpecifiesTwoStageSupport() -> std::string {
std::string msg;

20
experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
View File

@@ -9,6 +9,26 @@ namespace ck_tile::builder::factory {
using namespace ck_tile::builder::diagnostics;
template <typename T>
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 <typename T>
struct FwdXdlV3Algorithm {
CHECK_CONCEPT(T, ConvAlgorithmDescriptor)