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 4d81becfb5..4feb09cfe9 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -52,23 +52,20 @@ concept GridwiseWmmaGemmDescriptor = requires(T t) { }; -template -concept HasGemmKBatch = requires(T t) { - { t.k_batch_size}; -}; - -// Concept to check if GEMM k batch size is specified. -template -concept GemmKBatchSizeWellDefinedIfProvided = - !HasGemmKBatch || requires(T t) { {t.k_batch_size} -> std::convertible_to; }; - // Concept for vectorized data transfer for convolution input tensors. template concept BlockTransferDescriptor = requires(T t) { { t.k0 } -> std::convertible_to; { t.m_n } -> std::convertible_to; { t.k1 } -> std::convertible_to; - GemmKBatchSizeWellDefinedIfProvided; +}; + +template +concept BlockTransferDescriptorBwd = requires(T t) { + { t.k0 } -> std::convertible_to; + { t.m_n } -> std::convertible_to; + { t.k1 } -> std::convertible_to; + { t.k_batch_size } -> std::convertible_to;; }; // Concept for thread cluster dimensions for GEMM output tensor. @@ -210,6 +207,14 @@ 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 (Bwd direction). +template +concept SpecifiesBlockTransferBwd = requires(T t) { + { T::transfer.a.block_transfer } -> BlockTransferDescriptorBwd; + { T::transfer.b.block_transfer } -> BlockTransferDescriptorBwd; + { T::transfer.c.thread_cluster_dims } -> ThreadClusterDescriptor; +}; + // Concept to check if a struct specifies convolution scalar per vector infor for A, B and C. template concept SpecifiesTileTransfer = requires(T t) { 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 4f268b05e9..71e3d03da3 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -242,7 +242,7 @@ template struct BwdXdlAlgorithm { CHECK_CONCEPT(T, ConvAlgorithmDescriptor) CHECK_CONCEPT(T, SpecifiesThreadBlock) - CHECK_CONCEPT(T, SpecifiesBlockTransfer) + CHECK_CONCEPT(T, SpecifiesBlockTransferBwd) CHECK_CONCEPT(T, SpecifiesLdsTransfer) CHECK_CONCEPT(T, SpecifiesThreadClusterAccessOrder) CHECK_CONCEPT(T, SpecifiesSourceAccessOrder) @@ -252,7 +252,7 @@ struct BwdXdlAlgorithm { static constexpr bool c1 = c_ConvAlgorithmDescriptor; static constexpr bool c2 = c_SpecifiesThreadBlock; - static constexpr bool c3 = c_SpecifiesBlockTransfer; + static constexpr bool c3 = c_SpecifiesBlockTransferBwd; static constexpr bool c4 = c_SpecifiesLdsTransfer; static constexpr bool c5 = c_SpecifiesThreadClusterAccessOrder; static constexpr bool c6 = c_SpecifiesSourceAccessOrder; @@ -269,7 +269,7 @@ struct BwdXdlAlgorithm { "Concepts for BwdXdl Algorithm:\n") + DIAGNOSTIC_LINE(ConvAlgorithmDescriptor) + DIAGNOSTIC_LINE(SpecifiesThreadBlock) + - DIAGNOSTIC_LINE(SpecifiesBlockTransfer) + + DIAGNOSTIC_LINE(SpecifiesBlockTransferBwd) + DIAGNOSTIC_LINE(SpecifiesLdsTransfer) + DIAGNOSTIC_LINE(SpecifiesThreadClusterAccessOrder) + DIAGNOSTIC_LINE(SpecifiesSourceAccessOrder) +