From afbc4223ae06a8a2a91bab085b86820051acc153 Mon Sep 17 00:00:00 2001 From: John Shumway Date: Sun, 7 Sep 2025 18:40:33 +0000 Subject: [PATCH] Update concept names. Concepts should represent semantic categories. Use "Descriptor" suffix and "Specifies" prefix. --- .../ck_tile/builder/conv_algorithm.hpp | 63 ++++++++++--------- .../include/ck_tile/builder/conv_factory.hpp | 10 +-- .../builder/test/test_conv_instances.cpp | 16 ++--- 3 files changed, 45 insertions(+), 44 deletions(-) diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm.hpp index c3df8902e0..d8039c112c 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm.hpp @@ -15,16 +15,16 @@ struct MNK T k{}; }; -// Concept for thread block info for a GEMM problem. +// Concept for thread block dimensions for a GEMM problem. template -concept ThreadBlockInfo = requires(T t) { +concept ThreadBlockDescriptor = requires(T t) { { t.block_size } -> std::convertible_to; { t.submatrix.m } -> std::convertible_to; { t.submatrix.n } -> std::convertible_to; { t.submatrix.k } -> std::convertible_to; }; -// Describe a thread block for a GEMM. +// Specifiy thread block dimensions for a GEMM. struct ThreadBlock { // Thread block size. @@ -32,17 +32,17 @@ struct ThreadBlock // Size of the submatrix problem in a thread block. MNK submatrix; }; -static_assert(ThreadBlockInfo); +static_assert(ThreadBlockDescriptor); -// Concept to check if struct provides thread block info. +// Concept to check if struct specifies thread block info. template -concept HasThreadBlockInfo = requires { - { T::thread_block } -> ThreadBlockInfo; +concept SpecifiesThreadBlock = requires { + { T::thread_block } -> ThreadBlockDescriptor; }; // Concept for tuning parameters for a convolution problem. template -concept ConvTuningInfo = requires(T t) { +concept ConvTuningDescriptor = requires(T t) { { t.ak1 } -> std::convertible_to; { t.bk1 } -> std::convertible_to; { t.m_xdl_per_wave } -> std::convertible_to; @@ -58,51 +58,51 @@ struct ConvTuningParams int m_xdl_per_wave = 0; int n_xdl_per_wave = 0; }; -static_assert(ConvTuningInfo); +static_assert(ConvTuningDescriptor); -// Concept to check if a struct provides convolution tuning info. +// Concept to check if a struct specifies convolution tuning info. template -concept HasConvTuningInfo = requires { - { T::tuning_params } -> ConvTuningInfo; +concept SpecifiesConvTuning = requires { + { T::tuning_params } -> ConvTuningDescriptor; }; // Concept for A block transfer thread cluster lengths. template -concept BlockATransferLengths = requires(T t) { +concept BlockATransferDescriptor = requires(T t) { { t.k0 } -> std::convertible_to; { t.m } -> std::convertible_to; { t.k1 } -> std::convertible_to; }; // Describe A block transfer thread cluster lengths. -struct BlockATransferLengthsInfo +struct BlockATransferLengths { int k0; int m; int k1; }; -static_assert(BlockATransferLengths); +static_assert(BlockATransferDescriptor); // Concept for B block transfer thread cluster lengths. template -concept BlockBTransferLengths = requires(T t) { +concept BlockBTransferDescriptor = requires(T t) { { t.k0 } -> std::convertible_to; { t.n } -> std::convertible_to; { t.k1 } -> std::convertible_to; }; // Describe B block transfer thread cluster lengths. -struct BlockBTransferLengthsInfo +struct BlockBTransferLengths { int k0; int n; int k1; }; -static_assert(BlockBTransferLengths); +static_assert(BlockBTransferDescriptor); // Concept for C block transfer thread cluster lengths. template -concept BlockCTransferLengths = requires(T t) { +concept BlockCTransferDescriptor = requires(T t) { { t.m_block } -> std::convertible_to; { t.m_wave_per_xdl } -> std::convertible_to; { t.n_block } -> std::convertible_to; @@ -110,33 +110,34 @@ concept BlockCTransferLengths = requires(T t) { }; // Describe C block transfer thread cluster lengths. -struct BlockCTransferLengthsInfo +struct BlockCTransferLengths { int m_block; int m_wave_per_xdl; int n_block; int n_wave_per_xdl; }; -static_assert(BlockCTransferLengths); +static_assert(BlockCTransferDescriptor); -// Concept to check if a struct provides A Block tranfer info. +// Concept to check if a struct specifies A Block tranfer info. template -concept HasABlockTransferInfo = requires(T t) { - { T::block_transfer.thread_cluster_dims_a } -> BlockATransferLengths; +concept SpecifiesBlockATransfer = requires(T t) { + { T::block_transfer.thread_cluster_dims_a } -> BlockATransferDescriptor; }; -// Concept to check if a struct provides B Block tranfer info. +// Concept to check if a struct specifies B Block tranfer info. template -concept HasBBlockTransferInfo = requires(T t) { - { T::block_transfer.thread_cluster_dims_b } -> BlockBTransferLengths; +concept SpecifiesBlockBTransfer = requires(T t) { + { T::block_transfer.thread_cluster_dims_b } -> BlockBTransferDescriptor; }; -// Concept to check if a struct provides C Block tranfer info. +// Concept to check if a struct specifies C Block tranfer info. template -concept HasCBlockTransferInfo = requires(T t) { - { T::block_transfer.thread_cluster_dims_c } -> BlockCTransferLengths; +concept SpecifiesBlockCTransfer = requires(T t) { + { T::block_transfer.thread_cluster_dims_c } -> BlockCTransferDescriptor; }; +// Enums for the current block GEMM pipeline versions. enum class BlockGemmPipelineVersion { V1, @@ -145,7 +146,7 @@ enum class BlockGemmPipelineVersion V5 }; -// Concept to check if struct provides block_gemm_pipeline_version. +// Concept to check if struct specifies block_gemm_pipeline_version. template concept ProvidesBlockGemmPipelineVersion = requires { { T::pipeline_version } -> std::convertible_to; diff --git a/experimental/builder/include/ck_tile/builder/conv_factory.hpp b/experimental/builder/include/ck_tile/builder/conv_factory.hpp index 0fab2e6110..bc432784fb 100644 --- a/experimental/builder/include/ck_tile/builder/conv_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_factory.hpp @@ -77,7 +77,7 @@ template constexpr ConvBlock SetThreadBlockInfo() { using AlgorithmType = decltype(ALGORITHM); - if constexpr(HasThreadBlockInfo) + if constexpr(SpecifiesThreadBlock) { constexpr auto& TB = ALGORITHM.thread_block; return ConvBlock{ @@ -106,7 +106,7 @@ template constexpr ConvTuning SetConvTuningInfo() { using AlgorithmType = decltype(ALGORITHM); - if constexpr(HasConvTuningInfo) + if constexpr(SpecifiesConvTuning) { constexpr auto& TP = ALGORITHM.tuning_params; return ConvTuning{ @@ -163,7 +163,7 @@ constexpr BlockTransfer SetABlockTransfer() .add_extra = 0, }; using AlgorithmType = decltype(ALGORITHM); - if constexpr(HasABlockTransferInfo) + if constexpr(SpecifiesBlockATransfer) { constexpr auto& TCL = ALGORITHM.block_transfer.thread_cluster_dims_a; block_transfer.thread_cluster_dims = {TCL.k0, TCL.m, TCL.k1}; @@ -185,7 +185,7 @@ constexpr BlockTransfer SetBBlockTransfer() .add_extra = 0, }; using AlgorithmType = decltype(ALGORITHM); - if constexpr(HasBBlockTransferInfo) + if constexpr(SpecifiesBlockBTransfer) { constexpr auto& TCL = ALGORITHM.block_transfer.thread_cluster_dims_b; block_transfer.thread_cluster_dims = {TCL.k0, TCL.n, TCL.k1}; @@ -204,7 +204,7 @@ constexpr CBlockTransfer SetCBlockTransfer() .scaler_per_vector = 8, }; using AlgorithmType = decltype(ALGORITHM); - if constexpr(HasCBlockTransferInfo) + if constexpr(SpecifiesBlockCTransfer) { constexpr auto& TCL = ALGORITHM.block_transfer.thread_cluster_dims_c; block_transfer.thread_cluster_dims = { diff --git a/experimental/builder/test/test_conv_instances.cpp b/experimental/builder/test/test_conv_instances.cpp index 1c4fbff00b..26c46aba49 100644 --- a/experimental/builder/test/test_conv_instances.cpp +++ b/experimental/builder/test/test_conv_instances.cpp @@ -24,18 +24,18 @@ struct FwdConvAlgorithm ckb::ConvTuningParams tuning_params; struct BlockTransfer { - ckb::BlockATransferLengthsInfo thread_cluster_dims_a; - ckb::BlockBTransferLengthsInfo thread_cluster_dims_b; - ckb::BlockCTransferLengthsInfo thread_cluster_dims_c; + ckb::BlockATransferLengths thread_cluster_dims_a; + ckb::BlockBTransferLengths thread_cluster_dims_b; + ckb::BlockCTransferLengths thread_cluster_dims_c; } block_transfer; ckb::BlockGemmPipelineVersion pipeline_version; }; static_assert(ckb::ConvAlgorithm); -static_assert(ckb::HasThreadBlockInfo); -static_assert(ckb::HasConvTuningInfo); -static_assert(ckb::HasABlockTransferInfo); -static_assert(ckb::HasBBlockTransferInfo); -static_assert(ckb::HasCBlockTransferInfo); +static_assert(ckb::SpecifiesThreadBlock); +static_assert(ckb::SpecifiesConvTuning); +static_assert(ckb::SpecifiesBlockATransfer); +static_assert(ckb::SpecifiesBlockBTransfer); +static_assert(ckb::SpecifiesBlockCTransfer); static_assert(ckb::ProvidesBlockGemmPipelineVersion); struct TestCase