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

Concept improvements.

Browse Source
This commit is contained in:
Ville Pietilä
2025-12-23 10:27:38 -05:00
parent a1740c614b
commit 77e10c7b08
2 changed files with 19 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -52,23 +52,20 @@ concept GridwiseWmmaGemmDescriptor = requires(T t) {
};
template <typename T>
concept HasGemmKBatch = requires(T t) {
{ t.k_batch_size};
};
// Concept to check if GEMM k batch size is specified.
template <typename T>
concept GemmKBatchSizeWellDefinedIfProvided =
!HasGemmKBatch<T> || requires(T t) { {t.k_batch_size} -> std::convertible_to<size_t>; };
// Concept for vectorized data transfer for convolution input tensors.
template <typename T>
concept BlockTransferDescriptor = requires(T t) {
{ t.k0 } -> std::convertible_to<size_t>;
{ t.m_n } -> std::convertible_to<size_t>;
{ t.k1 } -> std::convertible_to<size_t>;
GemmKBatchSizeWellDefinedIfProvided<T>;
};
template <typename T>
concept BlockTransferDescriptorBwd = requires(T t) {
{ t.k0 } -> std::convertible_to<size_t>;
{ t.m_n } -> std::convertible_to<size_t>;
{ t.k1 } -> std::convertible_to<size_t>;
{ t.k_batch_size } -> std::convertible_to<size_t>;;
};
// 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 <typename T>
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 <typename T>
concept SpecifiesTileTransfer = requires(T t) {

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

@@ -242,7 +242,7 @@ template <typename T>
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) +