From 73b1e59c75a3cd602d800e4db2757241490d2e41 Mon Sep 17 00:00:00 2001 From: John Shumway Date: Tue, 6 Jan 2026 12:50:42 -0500 Subject: [PATCH] Create a non-templated conv_traits struct --- .../ck_tile/builder/reflect/conv_describe.hpp | 42 ++--- .../ck_tile/builder/reflect/conv_traits.hpp | 111 +++--------- .../reflect/instance_to_conv_traits.hpp | 140 +++++++++++++++ .../builder/test/conv/ck/test_conv_traits.cpp | 159 +++++++++--------- .../conv/ck/unit_instance_to_conv_traits.cpp | 127 +++++++------- 5 files changed, 331 insertions(+), 248 deletions(-) create mode 100644 experimental/builder/include/ck_tile/builder/reflect/instance_to_conv_traits.hpp diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_describe.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_describe.hpp index fdbfa7c4e1..3cfbc122c8 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_describe.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_describe.hpp @@ -7,7 +7,7 @@ #pragma once #include "ck_tile/builder/reflect/conv_description.hpp" -#include "ck_tile/builder/reflect/conv_traits.hpp" +#include "ck_tile/builder/reflect/instance_to_conv_traits.hpp" namespace ck_tile::reflect { @@ -17,31 +17,31 @@ namespace ck_tile::reflect { template conv::ConvDescription describe() { - using Traits = conv::ConvTraits; + const auto traits = conv::instance_to_conv_traits(); return conv::ConvDescription( conv::ConvSignatureInfo{ - .spatial_dim = Traits::spatial_dim, - .direction = Traits::direction, - .input_layout = Traits::layout[0], - .weight_layout = Traits::layout[1], - .output_layout = Traits::layout[2], - .data_type = Traits::data_type, - .input_element_op = Traits::input_element_op, - .weight_element_op = Traits::weight_element_op, - .output_element_op = Traits::output_element_op, + .spatial_dim = traits.spatial_dim, + .direction = traits.direction, + .input_layout = traits.layout[0], + .weight_layout = traits.layout[1], + .output_layout = traits.layout[2], + .data_type = traits.data_type, + .input_element_op = traits.input_element_op, + .weight_element_op = traits.weight_element_op, + .output_element_op = traits.output_element_op, }, conv::GemmAlgorithmInfo{ - .thread_block_size = Traits::thread_block_size, - .tile_dims = Traits::tile_dims, - .warp_gemm = Traits::warp_gemm, - .a_tile_transfer = Traits::a_tile_transfer, - .b_tile_transfer = Traits::b_tile_transfer, - .c_tile_transfer = Traits::c_tile_transfer, - .pipeline_version = Traits::pipeline_version, - .pipeline_scheduler = Traits::pipeline_scheduler, - .conv_specialization = Traits::conv_specialization, - .padding = Traits::gemm_padding, + .thread_block_size = traits.thread_block_size, + .tile_dims = traits.tile_dims, + .warp_gemm = traits.warp_gemm, + .a_tile_transfer = traits.a_tile_transfer, + .b_tile_transfer = traits.b_tile_transfer, + .c_tile_transfer = traits.c_tile_transfer, + .pipeline_version = traits.pipeline_version, + .pipeline_scheduler = traits.pipeline_scheduler, + .conv_specialization = traits.conv_specialization, + .padding = traits.gemm_padding, }, []() { return reflect::instance_string(); }); } diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp index b9cb20d675..3c0fad49d2 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp @@ -9,102 +9,41 @@ namespace ck_tile::reflect::conv { -/// @brief Primary template for extracting convolution traits. -/// @details This struct is the main entry point for reflecting on a convolution -/// kernel's properties. It is specialized to handle different kinds of input types. -template -struct ConvTraits; - -/// @brief Specialization of `ConvTraits` for a direct device kernel `Instance`. -/// @details This is the primary specialization used to extract a comprehensive -/// set of traits directly from a fully-formed device kernel `Instance` type. -/// It uses `InstanceTraits` to access the kernel's template parameters. -template - requires IsXdlFwdConv> -struct ConvTraits +/// @brief Pure data struct holding convolution traits without template parameters or static +/// members. +/// @details This struct can hold the data from any ConvTraitsImpl class, allowing runtime storage +/// and manipulation of convolution configuration information. +struct ConvTraits { - using InstTraits = InstanceTraits; - // --- Signature Information --- - /// @brief The number of spatial dimensions in the convolution (1, 2, or 3). - static constexpr int spatial_dim = InstTraits::kSpatialDim; - /// @brief The direction of the convolution (Forward, Backward Data, or Backward Weight). - static constexpr builder::ConvDirection direction = conv_direction(); - /// @brief The memory layout of the convolution tensors (e.g., GNHWC_GKYXC_GNHWK). - static constexpr auto layout = conv_layout(); - /// @brief The primary data type used in the computation (e.g., FP16, FP32). - static constexpr builder::DataType data_type = conv_data_type(); + int spatial_dim; + builder::ConvDirection direction; + std::array layout; // [input, weight, output] + builder::DataType data_type; - static constexpr builder::ElementwiseOperation input_element_op = - elementwise_op(); - static constexpr builder::ElementwiseOperation weight_element_op = - elementwise_op(); - static constexpr builder::ElementwiseOperation output_element_op = - elementwise_op(); + builder::ElementwiseOperation input_element_op; + builder::ElementwiseOperation weight_element_op; + builder::ElementwiseOperation output_element_op; - /// @brief The GEMM specialization used by the kernel - padding - static constexpr auto gemm_padding = gemm_spec(); - /// @brief The convolution-specific specialization (e.g., Default, 1x1). - static constexpr auto conv_specialization = conv_spec(); + builder::GemmPadding gemm_padding; + std::variant + conv_specialization; // --- Algorithm Information --- - /// @brief The total number of threads in a thread block (workgroup). - static constexpr int thread_block_size = InstTraits::kBlockSize; - /// @brief The dimensions of the data tile processed by the thread block. - static constexpr DataTileInfo tile_dims = { - .m = InstTraits::kMPerBlock, .n = InstTraits::kNPerBlock, .k = InstTraits::kKPerBlock}; + int thread_block_size; + DataTileInfo tile_dims; - /// @brief Configuration for the A-matrix (input) tile transfer. - static constexpr InputTileTransferInfo a_tile_transfer = { - .tile_dimensions = {.k0 = InstTraits::kKPerBlock / InstTraits::kAK1, - .m_or_n = InstTraits::kMPerBlock, - .k1 = InstTraits::kAK1}, - .transfer_params = {.k1 = InstTraits::kAK1, - .thread_cluster_dims = InstTraits::kAThreadClusterLengths, - .thread_cluster_order = InstTraits::kAThreadClusterArrangeOrder, - .src_access_order = InstTraits::kABlockTransferSrcAccessOrder, - .src_vector_dim = InstTraits::kABlockTransferSrcVectorDim, - .src_scalar_per_vector = InstTraits::kABlockTransferSrcScalarPerVector, - .dst_scalar_per_vector_k1 = - InstTraits::kABlockTransferDstScalarPerVectorK1, - .lds_padding = static_cast(InstTraits::kABlockLdsExtraM)}}; + InputTileTransferInfo a_tile_transfer; + InputTileTransferInfo b_tile_transfer; - /// @brief Configuration for the B-matrix (weights) tile transfer. - static constexpr InputTileTransferInfo b_tile_transfer = { - .tile_dimensions = {.k0 = InstTraits::kKPerBlock / InstTraits::kBK1, - .m_or_n = InstTraits::kNPerBlock, - .k1 = InstTraits::kBK1}, - .transfer_params = {.k1 = InstTraits::kBK1, - .thread_cluster_dims = InstTraits::kBThreadClusterLengths, - .thread_cluster_order = InstTraits::kBThreadClusterArrangeOrder, - .src_access_order = InstTraits::kBBlockTransferSrcAccessOrder, - .src_vector_dim = InstTraits::kBBlockTransferSrcVectorDim, - .src_scalar_per_vector = InstTraits::kBBlockTransferSrcScalarPerVector, - .dst_scalar_per_vector_k1 = - InstTraits::kBBlockTransferDstScalarPerVectorK1, - .lds_padding = static_cast(InstTraits::kBBlockLdsExtraN)}}; + WarpGemmParams warp_gemm; - /// @brief Parameters for the warp-level GEMM computation. - static constexpr WarpGemmParams warp_gemm = {.gemm_m = InstTraits::kMPerXDL, - .gemm_n = InstTraits::kNPerXDL, - .m_iter = InstTraits::kMXdlPerWave, - .n_iter = InstTraits::kNXdlPerWave}; + OutputTileTransferInfo c_tile_transfer; - /// @brief Configuration for the C-matrix (output) tile transfer. - static constexpr OutputTileTransferInfo c_tile_transfer = { - .shuffle_params = {.m_gemms_per_shuffle = InstTraits::kCShuffleMXdlPerWavePerShuffle, - .n_gemms_per_shuffle = InstTraits::kCShuffleNXdlPerWavePerShuffle}, - .thread_cluster_dims = {InstTraits::kCThreadClusterLengths[0], - InstTraits::kCThreadClusterLengths[1], - InstTraits::kCThreadClusterLengths[2], - InstTraits::kCThreadClusterLengths[3]}, - .scalar_per_vector = InstTraits::kCBlockTransferScalarPerVector}; - - /// @brief The block GEMM pipeline version used by the kernel. - static constexpr auto pipeline_version = get_pipeline_version(); - - /// @brief The pipeline scheduler used by the kernel. - static constexpr auto pipeline_scheduler = get_pipeline_scheduler(); + builder::PipelineVersion pipeline_version; + builder::PipelineScheduler pipeline_scheduler; }; } // namespace ck_tile::reflect::conv diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_to_conv_traits.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_to_conv_traits.hpp new file mode 100644 index 0000000000..302173b23e --- /dev/null +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_to_conv_traits.hpp @@ -0,0 +1,140 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "ck_tile/builder/reflect/conv_traits.hpp" +#include "ck_tile/builder/reflect/conv_traits_helpers.hpp" +#include "ck_tile/builder/reflect/instance_traits.hpp" + +namespace ck_tile::reflect::conv { + +/// @brief Primary template for extracting convolution traits. +/// @details This struct is the main entry point for reflecting on a convolution +/// kernel's properties. It is specialized to handle different kinds of input types. +template +struct ConvTraitsTmpl; + +/// @brief Specialization of `ConvTraitsTmpl` for a direct device kernel `Instance`. +/// @details This is the primary specialization used to extract a comprehensive +/// set of traits directly from a fully-formed device kernel `Instance` type. +/// It uses `InstanceTraits` to access the kernel's template parameters. +template + requires IsXdlFwdConv> +struct ConvTraitsTmpl +{ + using InstTraits = InstanceTraits; + + // --- Signature Information --- + /// @brief The number of spatial dimensions in the convolution (1, 2, or 3). + static constexpr int spatial_dim = InstTraits::kSpatialDim; + /// @brief The direction of the convolution (Forward, Backward Data, or Backward Weight). + static constexpr builder::ConvDirection direction = conv_direction(); + /// @brief The memory layout of the convolution tensors (e.g., GNHWC_GKYXC_GNHWK). + static constexpr auto layout = conv_layout(); + /// @brief The primary data type used in the computation (e.g., FP16, FP32). + static constexpr builder::DataType data_type = conv_data_type(); + + static constexpr builder::ElementwiseOperation input_element_op = + elementwise_op(); + static constexpr builder::ElementwiseOperation weight_element_op = + elementwise_op(); + static constexpr builder::ElementwiseOperation output_element_op = + elementwise_op(); + + /// @brief The GEMM specialization used by the kernel - padding + static constexpr auto gemm_padding = gemm_spec(); + /// @brief The convolution-specific specialization (e.g., Default, 1x1). + static constexpr auto conv_specialization = conv_spec(); + + // --- Algorithm Information --- + /// @brief The total number of threads in a thread block (workgroup). + static constexpr int thread_block_size = InstTraits::kBlockSize; + /// @brief The dimensions of the data tile processed by the thread block. + static constexpr DataTileInfo tile_dims = { + .m = InstTraits::kMPerBlock, .n = InstTraits::kNPerBlock, .k = InstTraits::kKPerBlock}; + + /// @brief Configuration for the A-matrix (input) tile transfer. + static constexpr InputTileTransferInfo a_tile_transfer = { + .tile_dimensions = {.k0 = InstTraits::kKPerBlock / InstTraits::kAK1, + .m_or_n = InstTraits::kMPerBlock, + .k1 = InstTraits::kAK1}, + .transfer_params = {.k1 = InstTraits::kAK1, + .thread_cluster_dims = InstTraits::kAThreadClusterLengths, + .thread_cluster_order = InstTraits::kAThreadClusterArrangeOrder, + .src_access_order = InstTraits::kABlockTransferSrcAccessOrder, + .src_vector_dim = InstTraits::kABlockTransferSrcVectorDim, + .src_scalar_per_vector = InstTraits::kABlockTransferSrcScalarPerVector, + .dst_scalar_per_vector_k1 = + InstTraits::kABlockTransferDstScalarPerVectorK1, + .lds_padding = static_cast(InstTraits::kABlockLdsExtraM)}}; + + /// @brief Configuration for the B-matrix (weights) tile transfer. + static constexpr InputTileTransferInfo b_tile_transfer = { + .tile_dimensions = {.k0 = InstTraits::kKPerBlock / InstTraits::kBK1, + .m_or_n = InstTraits::kNPerBlock, + .k1 = InstTraits::kBK1}, + .transfer_params = {.k1 = InstTraits::kBK1, + .thread_cluster_dims = InstTraits::kBThreadClusterLengths, + .thread_cluster_order = InstTraits::kBThreadClusterArrangeOrder, + .src_access_order = InstTraits::kBBlockTransferSrcAccessOrder, + .src_vector_dim = InstTraits::kBBlockTransferSrcVectorDim, + .src_scalar_per_vector = InstTraits::kBBlockTransferSrcScalarPerVector, + .dst_scalar_per_vector_k1 = + InstTraits::kBBlockTransferDstScalarPerVectorK1, + .lds_padding = static_cast(InstTraits::kBBlockLdsExtraN)}}; + + /// @brief Parameters for the warp-level GEMM computation. + static constexpr WarpGemmParams warp_gemm = {.gemm_m = InstTraits::kMPerXDL, + .gemm_n = InstTraits::kNPerXDL, + .m_iter = InstTraits::kMXdlPerWave, + .n_iter = InstTraits::kNXdlPerWave}; + + /// @brief Configuration for the C-matrix (output) tile transfer. + static constexpr OutputTileTransferInfo c_tile_transfer = { + .shuffle_params = {.m_gemms_per_shuffle = InstTraits::kCShuffleMXdlPerWavePerShuffle, + .n_gemms_per_shuffle = InstTraits::kCShuffleNXdlPerWavePerShuffle}, + .thread_cluster_dims = {InstTraits::kCThreadClusterLengths[0], + InstTraits::kCThreadClusterLengths[1], + InstTraits::kCThreadClusterLengths[2], + InstTraits::kCThreadClusterLengths[3]}, + .scalar_per_vector = InstTraits::kCBlockTransferScalarPerVector}; + + /// @brief The block GEMM pipeline version used by the kernel. + static constexpr auto pipeline_version = get_pipeline_version(); + + /// @brief The pipeline scheduler used by the kernel. + static constexpr auto pipeline_scheduler = get_pipeline_scheduler(); +}; + +/// @brief Converts ConvTraitsTmpl to runtime ConvTraits struct. +/// @details This function extracts compile-time traits from a device kernel instance +/// and packages them into a runtime-accessible ConvTraits struct. +template + requires HasConvTraits +constexpr ConvTraits instance_to_conv_traits() +{ + using Traits = ConvTraitsTmpl; + + return ConvTraits{ + .spatial_dim = Traits::spatial_dim, + .direction = Traits::direction, + .layout = Traits::layout, + .data_type = Traits::data_type, + .input_element_op = Traits::input_element_op, + .weight_element_op = Traits::weight_element_op, + .output_element_op = Traits::output_element_op, + .gemm_padding = Traits::gemm_padding, + .conv_specialization = Traits::conv_specialization, + .thread_block_size = Traits::thread_block_size, + .tile_dims = Traits::tile_dims, + .a_tile_transfer = Traits::a_tile_transfer, + .b_tile_transfer = Traits::b_tile_transfer, + .warp_gemm = Traits::warp_gemm, + .c_tile_transfer = Traits::c_tile_transfer, + .pipeline_version = Traits::pipeline_version, + .pipeline_scheduler = Traits::pipeline_scheduler, + }; +} + +} // namespace ck_tile::reflect::conv diff --git a/experimental/builder/test/conv/ck/test_conv_traits.cpp b/experimental/builder/test/conv/ck/test_conv_traits.cpp index d5661ad67b..e43b2f4504 100644 --- a/experimental/builder/test/conv/ck/test_conv_traits.cpp +++ b/experimental/builder/test/conv/ck/test_conv_traits.cpp @@ -6,7 +6,7 @@ #include #include -#include +#include #include #include #include @@ -86,72 +86,73 @@ TEST_F(ConvTraitsTest, ConvFwdTraitsExtraction) ck::half_t, // BComputeDataType false>; // DirectLoad - // Use ConvTraits to extract compile-time information - using Traits = ck_tile::reflect::conv::ConvTraits; + // Use ConvTraitsTmpl to extract compile-time information + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); // Verify signature information - EXPECT_EQ(Traits::spatial_dim, 2); - EXPECT_EQ(Traits::direction, ConvDirection::FORWARD); - EXPECT_THAT(Traits::layout, + EXPECT_EQ(traits.spatial_dim, 2); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::GNHWC, TensorLayout::GKYXC, TensorLayout::GNHWK)); - EXPECT_EQ(Traits::data_type, DataType::FP16); - EXPECT_EQ(Traits::input_element_op, ElementwiseOperation::PASS_THROUGH); - EXPECT_EQ(Traits::weight_element_op, ElementwiseOperation::PASS_THROUGH); - EXPECT_EQ(Traits::output_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.data_type, DataType::FP16); + EXPECT_EQ(traits.input_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.weight_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.output_element_op, ElementwiseOperation::PASS_THROUGH); // Verify specializations - EXPECT_EQ(Traits::gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(traits.gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); + EXPECT_EQ(std::get(traits.conv_specialization), + ck_tile::builder::ConvFwdSpecialization::DEFAULT); // Verify algorithm information - EXPECT_EQ(Traits::thread_block_size, 256); + EXPECT_EQ(traits.thread_block_size, 256); // Verify tile dimensions - EXPECT_EQ(Traits::tile_dims.m, 128); - EXPECT_EQ(Traits::tile_dims.n, 128); - EXPECT_EQ(Traits::tile_dims.k, 16); + EXPECT_EQ(traits.tile_dims.m, 128); + EXPECT_EQ(traits.tile_dims.n, 128); + EXPECT_EQ(traits.tile_dims.k, 16); // Verify A tile transfer info - EXPECT_EQ(Traits::a_tile_transfer.tile_dimensions.k0, 2); - EXPECT_EQ(Traits::a_tile_transfer.tile_dimensions.m_or_n, 128); - EXPECT_EQ(Traits::a_tile_transfer.tile_dimensions.k1, 8); - EXPECT_EQ(Traits::a_tile_transfer.transfer_params.k1, 8); - EXPECT_THAT(Traits::a_tile_transfer.transfer_params.thread_cluster_dims, ElementsAre(4, 64, 1)); - EXPECT_THAT(Traits::a_tile_transfer.transfer_params.thread_cluster_order, ElementsAre(1, 0, 2)); - EXPECT_THAT(Traits::a_tile_transfer.transfer_params.src_access_order, ElementsAre(1, 0, 2)); - EXPECT_EQ(Traits::a_tile_transfer.transfer_params.src_vector_dim, 2); - EXPECT_EQ(Traits::a_tile_transfer.transfer_params.src_scalar_per_vector, 8); - EXPECT_EQ(Traits::a_tile_transfer.transfer_params.dst_scalar_per_vector_k1, 8); - EXPECT_TRUE(Traits::a_tile_transfer.transfer_params.lds_padding); + EXPECT_EQ(traits.a_tile_transfer.tile_dimensions.k0, 2); + EXPECT_EQ(traits.a_tile_transfer.tile_dimensions.m_or_n, 128); + EXPECT_EQ(traits.a_tile_transfer.tile_dimensions.k1, 8); + EXPECT_EQ(traits.a_tile_transfer.transfer_params.k1, 8); + EXPECT_THAT(traits.a_tile_transfer.transfer_params.thread_cluster_dims, ElementsAre(4, 64, 1)); + EXPECT_THAT(traits.a_tile_transfer.transfer_params.thread_cluster_order, ElementsAre(1, 0, 2)); + EXPECT_THAT(traits.a_tile_transfer.transfer_params.src_access_order, ElementsAre(1, 0, 2)); + EXPECT_EQ(traits.a_tile_transfer.transfer_params.src_vector_dim, 2); + EXPECT_EQ(traits.a_tile_transfer.transfer_params.src_scalar_per_vector, 8); + EXPECT_EQ(traits.a_tile_transfer.transfer_params.dst_scalar_per_vector_k1, 8); + EXPECT_TRUE(traits.a_tile_transfer.transfer_params.lds_padding); // Verify B tile transfer info - EXPECT_EQ(Traits::b_tile_transfer.tile_dimensions.k0, 2); - EXPECT_EQ(Traits::b_tile_transfer.tile_dimensions.m_or_n, 128); - EXPECT_EQ(Traits::b_tile_transfer.tile_dimensions.k1, 8); - EXPECT_EQ(Traits::b_tile_transfer.transfer_params.k1, 8); - EXPECT_THAT(Traits::b_tile_transfer.transfer_params.thread_cluster_dims, ElementsAre(4, 64, 1)); - EXPECT_THAT(Traits::b_tile_transfer.transfer_params.thread_cluster_order, ElementsAre(1, 0, 2)); - EXPECT_THAT(Traits::b_tile_transfer.transfer_params.src_access_order, ElementsAre(1, 0, 2)); - EXPECT_EQ(Traits::b_tile_transfer.transfer_params.src_vector_dim, 2); - EXPECT_EQ(Traits::b_tile_transfer.transfer_params.src_scalar_per_vector, 8); - EXPECT_EQ(Traits::b_tile_transfer.transfer_params.dst_scalar_per_vector_k1, 8); - EXPECT_TRUE(Traits::b_tile_transfer.transfer_params.lds_padding); + EXPECT_EQ(traits.b_tile_transfer.tile_dimensions.k0, 2); + EXPECT_EQ(traits.b_tile_transfer.tile_dimensions.m_or_n, 128); + EXPECT_EQ(traits.b_tile_transfer.tile_dimensions.k1, 8); + EXPECT_EQ(traits.b_tile_transfer.transfer_params.k1, 8); + EXPECT_THAT(traits.b_tile_transfer.transfer_params.thread_cluster_dims, ElementsAre(4, 64, 1)); + EXPECT_THAT(traits.b_tile_transfer.transfer_params.thread_cluster_order, ElementsAre(1, 0, 2)); + EXPECT_THAT(traits.b_tile_transfer.transfer_params.src_access_order, ElementsAre(1, 0, 2)); + EXPECT_EQ(traits.b_tile_transfer.transfer_params.src_vector_dim, 2); + EXPECT_EQ(traits.b_tile_transfer.transfer_params.src_scalar_per_vector, 8); + EXPECT_EQ(traits.b_tile_transfer.transfer_params.dst_scalar_per_vector_k1, 8); + EXPECT_TRUE(traits.b_tile_transfer.transfer_params.lds_padding); // Verify warp GEMM params - EXPECT_EQ(Traits::warp_gemm.gemm_m, 32); - EXPECT_EQ(Traits::warp_gemm.gemm_n, 32); - EXPECT_EQ(Traits::warp_gemm.m_iter, 4); - EXPECT_EQ(Traits::warp_gemm.n_iter, 4); + EXPECT_EQ(traits.warp_gemm.gemm_m, 32); + EXPECT_EQ(traits.warp_gemm.gemm_n, 32); + EXPECT_EQ(traits.warp_gemm.m_iter, 4); + EXPECT_EQ(traits.warp_gemm.n_iter, 4); // Verify output tile transfer info - EXPECT_EQ(Traits::c_tile_transfer.shuffle_params.m_gemms_per_shuffle, 1); - EXPECT_EQ(Traits::c_tile_transfer.shuffle_params.n_gemms_per_shuffle, 1); - EXPECT_THAT(Traits::c_tile_transfer.thread_cluster_dims, ElementsAre(1, 32, 1, 8)); - EXPECT_EQ(Traits::c_tile_transfer.scalar_per_vector, 8); + EXPECT_EQ(traits.c_tile_transfer.shuffle_params.m_gemms_per_shuffle, 1); + EXPECT_EQ(traits.c_tile_transfer.shuffle_params.n_gemms_per_shuffle, 1); + EXPECT_THAT(traits.c_tile_transfer.thread_cluster_dims, ElementsAre(1, 32, 1, 8)); + EXPECT_EQ(traits.c_tile_transfer.scalar_per_vector, 8); // Verify pipeline configuration - EXPECT_EQ(Traits::pipeline_scheduler, PipelineScheduler::INTRAWAVE); - EXPECT_EQ(Traits::pipeline_version, PipelineVersion::V1); + EXPECT_EQ(traits.pipeline_scheduler, PipelineScheduler::INTRAWAVE); + EXPECT_EQ(traits.pipeline_version, PipelineVersion::V1); } // Test ConvTraits with DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle @@ -214,30 +215,31 @@ TEST_F(ConvTraitsTest, ConvFwdBaseTraitsExtraction) ck::LoopScheduler::Default, // LoopSched 1>; // NumGroupsToMerge - // Use ConvTraits to extract compile-time information - using Traits = ck_tile::reflect::conv::ConvTraits; + // Use ConvTraitsTmpl to extract compile-time information + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); // Verify signature information - EXPECT_EQ(Traits::spatial_dim, 2); - EXPECT_EQ(Traits::direction, ConvDirection::FORWARD); - EXPECT_THAT(Traits::layout, + EXPECT_EQ(traits.spatial_dim, 2); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::GNHWC, TensorLayout::GKYXC, TensorLayout::GNHWK)); - EXPECT_EQ(Traits::data_type, DataType::FP16); - EXPECT_EQ(Traits::input_element_op, ElementwiseOperation::PASS_THROUGH); - EXPECT_EQ(Traits::weight_element_op, ElementwiseOperation::PASS_THROUGH); - EXPECT_EQ(Traits::output_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.data_type, DataType::FP16); + EXPECT_EQ(traits.input_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.weight_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.output_element_op, ElementwiseOperation::PASS_THROUGH); // Verify specializations - EXPECT_EQ(Traits::gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(traits.gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); + EXPECT_EQ(std::get(traits.conv_specialization), + ck_tile::builder::ConvFwdSpecialization::DEFAULT); // Verify algorithm information - EXPECT_EQ(Traits::thread_block_size, 256); + EXPECT_EQ(traits.thread_block_size, 256); // Verify tile dimensions - EXPECT_EQ(Traits::tile_dims.m, 128); - EXPECT_EQ(Traits::tile_dims.n, 128); - EXPECT_EQ(Traits::tile_dims.k, 16); + EXPECT_EQ(traits.tile_dims.m, 128); + EXPECT_EQ(traits.tile_dims.n, 128); + EXPECT_EQ(traits.tile_dims.k, 16); } // Test ConvTraits with DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor TEST_F(ConvTraitsTest, ConvFwdLargeTensorTraitsExtraction) @@ -298,29 +300,30 @@ TEST_F(ConvTraitsTest, ConvFwdLargeTensorTraitsExtraction) ck::half_t, // BComputeDataType ck::LoopScheduler::Default>; // LoopSched - // Use ConvTraits to extract compile-time information - using Traits = ck_tile::reflect::conv::ConvTraits; + // Use ConvTraitsTmpl to extract compile-time information + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); // Verify signature information - EXPECT_EQ(Traits::spatial_dim, 2); - EXPECT_EQ(Traits::direction, ConvDirection::FORWARD); - EXPECT_THAT(Traits::layout, + EXPECT_EQ(traits.spatial_dim, 2); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::GNHWC, TensorLayout::GKYXC, TensorLayout::GNHWK)); - EXPECT_EQ(Traits::data_type, DataType::FP16); - EXPECT_EQ(Traits::input_element_op, ElementwiseOperation::PASS_THROUGH); - EXPECT_EQ(Traits::weight_element_op, ElementwiseOperation::PASS_THROUGH); - EXPECT_EQ(Traits::output_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.data_type, DataType::FP16); + EXPECT_EQ(traits.input_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.weight_element_op, ElementwiseOperation::PASS_THROUGH); + EXPECT_EQ(traits.output_element_op, ElementwiseOperation::PASS_THROUGH); // Verify specializations - EXPECT_EQ(Traits::gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(traits.gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); + EXPECT_EQ(std::get(traits.conv_specialization), + ck_tile::builder::ConvFwdSpecialization::DEFAULT); // Verify algorithm information - EXPECT_EQ(Traits::thread_block_size, 256); + EXPECT_EQ(traits.thread_block_size, 256); // Verify tile dimensions - EXPECT_EQ(Traits::tile_dims.m, 128); - EXPECT_EQ(Traits::tile_dims.n, 128); - EXPECT_EQ(Traits::tile_dims.k, 16); + EXPECT_EQ(traits.tile_dims.m, 128); + EXPECT_EQ(traits.tile_dims.n, 128); + EXPECT_EQ(traits.tile_dims.k, 16); } } // anonymous namespace diff --git a/experimental/builder/test/conv/ck/unit_instance_to_conv_traits.cpp b/experimental/builder/test/conv/ck/unit_instance_to_conv_traits.cpp index de2a4fdd14..fcf22440da 100644 --- a/experimental/builder/test/conv/ck/unit_instance_to_conv_traits.cpp +++ b/experimental/builder/test/conv/ck/unit_instance_to_conv_traits.cpp @@ -2,19 +2,19 @@ // SPDX-License-Identifier: MIT // ============================================================================ -// Unit Tests for InstanceTraits to ConvTraits Conversion +// Unit Tests for InstanceTraits to ConvTraitsTmpl Conversion // ============================================================================ // // PURPOSE: // -------- // These tests verify the conversion layer between InstanceTraits (low-level -// template parameter extraction) and ConvTraits (high-level semantic traits). +// template parameter extraction) and ConvTraitsTmpl (high-level semantic traits). // The conversion transforms raw CK kernel parameters into builder-friendly // enums and structures. // // DESIGN RATIONALE: // ----------------- -// ConvTraits uses a single generic specialization that works with any Device +// ConvTraitsTmpl uses a single generic specialization that works with any Device // class satisfying the IsXdlFwdConv concept. This use of concepts is fragile // and introduces extra complexity. We want to refector to just use functions // for this conversion. @@ -42,7 +42,7 @@ #include #include -#include +#include #include #include #include @@ -165,9 +165,9 @@ TEST(InstanceToConvTraits, DetectsForwardDirection) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::direction, ConvDirection::FORWARD); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); } // ============================================================================ @@ -228,9 +228,10 @@ TEST(InstanceToConvTraits, ExtractsDefaultSpecialization) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(std::get(traits.conv_specialization), + ck_tile::builder::ConvFwdSpecialization::DEFAULT); } TEST(InstanceToConvTraits, ExtractsFilter1x1Pad0Specialization) @@ -287,9 +288,9 @@ TEST(InstanceToConvTraits, ExtractsFilter1x1Pad0Specialization) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::conv_specialization, + EXPECT_EQ(std::get(traits.conv_specialization), ck_tile::builder::ConvFwdSpecialization::FILTER_1X1_PAD0); } @@ -351,9 +352,9 @@ TEST(InstanceToConvTraits, ExtractsGnhwcLayout) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_THAT(Traits::layout, + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::GNHWC, TensorLayout::GKYXC, TensorLayout::GNHWK)); } @@ -411,9 +412,9 @@ TEST(InstanceToConvTraits, ExtractsNhwgcLayout) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_THAT(Traits::layout, + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::NHWGC, TensorLayout::GKYXC, TensorLayout::NHWGK)); } @@ -471,9 +472,9 @@ TEST(InstanceToConvTraits, ExtractsNgchwGkyxcLayout) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_THAT(Traits::layout, + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::NGCHW, TensorLayout::GKYXC, TensorLayout::NGKHW)); } @@ -531,9 +532,9 @@ TEST(InstanceToConvTraits, ExtractsNgchwGkcyxLayout) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_THAT(Traits::layout, + EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::NGCHW, TensorLayout::GKCYX, TensorLayout::NGKHW)); } @@ -595,9 +596,9 @@ TEST(InstanceToConvTraits, ExtractsFp16DataType) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::data_type, DataType::FP16); + EXPECT_EQ(traits.data_type, DataType::FP16); } TEST(InstanceToConvTraits, ExtractsBf16DataType) @@ -654,9 +655,9 @@ TEST(InstanceToConvTraits, ExtractsBf16DataType) ck::bhalf_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::data_type, DataType::BF16); + EXPECT_EQ(traits.data_type, DataType::BF16); } TEST(InstanceToConvTraits, ExtractsFp32DataType) @@ -713,9 +714,9 @@ TEST(InstanceToConvTraits, ExtractsFp32DataType) float, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::data_type, DataType::FP32); + EXPECT_EQ(traits.data_type, DataType::FP32); } TEST(InstanceToConvTraits, ExtractsI8DataType) @@ -772,9 +773,9 @@ TEST(InstanceToConvTraits, ExtractsI8DataType) int8_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::data_type, DataType::I8); + EXPECT_EQ(traits.data_type, DataType::I8); } // ============================================================================ @@ -835,9 +836,9 @@ TEST(InstanceToConvTraits, ExtractsDefaultGemmPadding) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::gemm_padding, GemmPadding::DEFAULT); + EXPECT_EQ(traits.gemm_padding, GemmPadding::DEFAULT); } TEST(InstanceToConvTraits, ExtractsMnkGemmPadding) @@ -894,9 +895,9 @@ TEST(InstanceToConvTraits, ExtractsMnkGemmPadding) ck::half_t, false>; - using Traits = ck_tile::reflect::conv::ConvTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); - EXPECT_EQ(Traits::gemm_padding, GemmPadding::MNK_PADDING); + EXPECT_EQ(traits.gemm_padding, GemmPadding::MNK_PADDING); } // ============================================================================ @@ -961,23 +962,23 @@ TEST(InstanceToConvTraits, TransformsFwdMultipleAbdXdlCShuffleV3) ck::half_t, // BComputeDataType false>; // DirectLoad - using InstTraits = ck_tile::reflect::InstanceTraits; - using ConvTraits = ck_tile::reflect::conv::ConvTraits; + using InstTraits = ck_tile::reflect::InstanceTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); // Verify signature information - EXPECT_EQ(ConvTraits::spatial_dim, InstTraits::kSpatialDim); - EXPECT_EQ(ConvTraits::direction, ConvDirection::FORWARD); - EXPECT_EQ(ConvTraits::data_type, DataType::FP16); - EXPECT_EQ(ConvTraits::gemm_padding, GemmPadding::DEFAULT); + EXPECT_EQ(traits.spatial_dim, InstTraits::kSpatialDim); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_EQ(traits.data_type, DataType::FP16); + EXPECT_EQ(traits.gemm_padding, GemmPadding::DEFAULT); // Verify tile dimensions - EXPECT_EQ(ConvTraits::tile_dims.m, InstTraits::kMPerBlock); - EXPECT_EQ(ConvTraits::tile_dims.n, InstTraits::kNPerBlock); - EXPECT_EQ(ConvTraits::tile_dims.k, InstTraits::kKPerBlock); + EXPECT_EQ(traits.tile_dims.m, InstTraits::kMPerBlock); + EXPECT_EQ(traits.tile_dims.n, InstTraits::kNPerBlock); + EXPECT_EQ(traits.tile_dims.k, InstTraits::kKPerBlock); // Verify pipeline configuration - EXPECT_EQ(ConvTraits::pipeline_scheduler, PipelineScheduler::INTRAWAVE); - EXPECT_EQ(ConvTraits::pipeline_version, PipelineVersion::V1); + EXPECT_EQ(traits.pipeline_scheduler, PipelineScheduler::INTRAWAVE); + EXPECT_EQ(traits.pipeline_version, PipelineVersion::V1); } TEST(InstanceToConvTraits, TransformsFwdMultipleAbdXdlCShuffle) @@ -1034,23 +1035,23 @@ TEST(InstanceToConvTraits, TransformsFwdMultipleAbdXdlCShuffle) ck::LoopScheduler::Default, // LoopSched 1>; // NumGroupsToMerge - using InstTraits = ck_tile::reflect::InstanceTraits; - using ConvTraits = ck_tile::reflect::conv::ConvTraits; + using InstTraits = ck_tile::reflect::InstanceTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); // Verify signature information - EXPECT_EQ(ConvTraits::spatial_dim, InstTraits::kSpatialDim); - EXPECT_EQ(ConvTraits::direction, ConvDirection::FORWARD); - EXPECT_EQ(ConvTraits::data_type, DataType::FP16); - EXPECT_EQ(ConvTraits::gemm_padding, GemmPadding::DEFAULT); + EXPECT_EQ(traits.spatial_dim, InstTraits::kSpatialDim); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_EQ(traits.data_type, DataType::FP16); + EXPECT_EQ(traits.gemm_padding, GemmPadding::DEFAULT); // Verify tile dimensions - EXPECT_EQ(ConvTraits::tile_dims.m, InstTraits::kMPerBlock); - EXPECT_EQ(ConvTraits::tile_dims.n, InstTraits::kNPerBlock); - EXPECT_EQ(ConvTraits::tile_dims.k, InstTraits::kKPerBlock); + EXPECT_EQ(traits.tile_dims.m, InstTraits::kMPerBlock); + EXPECT_EQ(traits.tile_dims.n, InstTraits::kNPerBlock); + EXPECT_EQ(traits.tile_dims.k, InstTraits::kKPerBlock); // Verify pipeline configuration (uses LoopScheduler instead of BlockGemmPipelineScheduler) - EXPECT_EQ(ConvTraits::pipeline_scheduler, PipelineScheduler::DEFAULT); - EXPECT_EQ(ConvTraits::pipeline_version, PipelineVersion::V1); + EXPECT_EQ(traits.pipeline_scheduler, PipelineScheduler::DEFAULT); + EXPECT_EQ(traits.pipeline_version, PipelineVersion::V1); } TEST(InstanceToConvTraits, TransformsFwdMultipleDXdlLargeTensor) @@ -1106,23 +1107,23 @@ TEST(InstanceToConvTraits, TransformsFwdMultipleDXdlLargeTensor) ck::half_t, // BComputeDataType ck::LoopScheduler::Default>; // LoopSched - using InstTraits = ck_tile::reflect::InstanceTraits; - using ConvTraits = ck_tile::reflect::conv::ConvTraits; + using InstTraits = ck_tile::reflect::InstanceTraits; + const auto traits = ck_tile::reflect::conv::instance_to_conv_traits(); // Verify signature information - EXPECT_EQ(ConvTraits::spatial_dim, InstTraits::kSpatialDim); - EXPECT_EQ(ConvTraits::direction, ConvDirection::FORWARD); - EXPECT_EQ(ConvTraits::data_type, DataType::FP16); - EXPECT_EQ(ConvTraits::gemm_padding, GemmPadding::DEFAULT); + EXPECT_EQ(traits.spatial_dim, InstTraits::kSpatialDim); + EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_EQ(traits.data_type, DataType::FP16); + EXPECT_EQ(traits.gemm_padding, GemmPadding::DEFAULT); // Verify tile dimensions - EXPECT_EQ(ConvTraits::tile_dims.m, InstTraits::kMPerBlock); - EXPECT_EQ(ConvTraits::tile_dims.n, InstTraits::kNPerBlock); - EXPECT_EQ(ConvTraits::tile_dims.k, InstTraits::kKPerBlock); + EXPECT_EQ(traits.tile_dims.m, InstTraits::kMPerBlock); + EXPECT_EQ(traits.tile_dims.n, InstTraits::kNPerBlock); + EXPECT_EQ(traits.tile_dims.k, InstTraits::kKPerBlock); // Verify pipeline configuration - EXPECT_EQ(ConvTraits::pipeline_scheduler, PipelineScheduler::DEFAULT); - EXPECT_EQ(ConvTraits::pipeline_version, PipelineVersion::V1); + EXPECT_EQ(traits.pipeline_scheduler, PipelineScheduler::DEFAULT); + EXPECT_EQ(traits.pipeline_version, PipelineVersion::V1); } } // anonymous namespace