From b70a58fcddc6ffb5ac9c43a9d44487ff963aab61 Mon Sep 17 00:00:00 2001 From: John Shumway Date: Thu, 4 Sep 2025 14:02:34 +0000 Subject: [PATCH] Split builder tests and instance tests. We have a typed test suite of all the instance we want to create, and also we have simple test of the builder. Split those into two different test suites. --- experimental/builder/test/CMakeLists.txt | 2 +- .../builder/test/test_conv_builder.cpp | 124 --------------- .../builder/test/test_conv_instances.cpp | 145 ++++++++++++++++++ 3 files changed, 146 insertions(+), 125 deletions(-) create mode 100644 experimental/builder/test/test_conv_instances.cpp diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index f9c0b0cb84..3560c22146 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -2,7 +2,7 @@ set(CMAKE_CXX_STANDARD 20) include(gtest) -add_executable(test_conv_builder test_conv_builder.cpp) +add_executable(test_conv_builder test_conv_builder.cpp test_conv_instances.cpp) target_include_directories(test_conv_builder PRIVATE "${PROJECT_SOURCE_DIR}/experimental/builder/include" "${PROJECT_SOURCE_DIR}/include" diff --git a/experimental/builder/test/test_conv_builder.cpp b/experimental/builder/test/test_conv_builder.cpp index 4be3949643..d0a6ddaa0d 100644 --- a/experimental/builder/test/test_conv_builder.cpp +++ b/experimental/builder/test/test_conv_builder.cpp @@ -32,128 +32,4 @@ TEST(ConvBuilderTest, TestDefaultInstance) "8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>"); } -struct FwdConvAlgorithm -{ - ckb::ThreadBlock thread_block; - ckb::ConvTuningParams tuning_params; - struct BlockTransfer - { - ckb::BlockATransferLengthsInfo thread_cluster_lengths_a; - ckb::BlockBTransferLengthsInfo thread_cluster_lengths_b; - ckb::BlockCTransferLengthsInfo thread_cluster_lengths_c; - } block_transfer; -}; -static_assert(ckb::ConvAlgorithm); -static_assert(ckb::HasThreadBlockInfo); -static_assert(ckb::HasConvTuningInfo); -static_assert(ckb::HasABlockTransferInfo); -static_assert(ckb::HasBBlockTransferInfo); -static_assert(ckb::HasCBlockTransferInfo); - -struct TestCase -{ - std::string_view name; - FwdConvAlgorithm algorithm; - std::string_view expected_type; -}; - -constexpr std::array TEST_CASES = { - TestCase{ - .name = "ConvFwdXdlBf16CompInstances2x_0", - .algorithm = - {.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 128, .k = 64}}, - .tuning_params{.ak1 = 16, .bk1 = 16, .m_xdl_per_wave = 2, .n_xdl_per_wave = 2}, - .block_transfer{ - .thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1}, - .thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1}, - .thread_cluster_lengths_c = - {.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8}, - }}, - .expected_type = - "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 128, 64, Default, 32, 32, " - "2, 2, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>", - }, - TestCase{ - .name = "GroupedConvFwdXdlBf16CompInstance0", - .algorithm = {.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 256, .k = 32}}, - .tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4}, - .block_transfer{ - .thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1}, - .thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1}, - .thread_cluster_lengths_c = {.m_block = 1, - .m_wave_per_xdl = 32, - .n_block = 1, - .n_wave_per_xdl = 8}, - }}, - .expected_type = - "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, " - "4, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>", - }}; - -static constexpr int NUM_TEST_CASES = std::tuple_size_v; - -// Helper to generate testing::Types, TestIndex<1>, ..., TestIndex>. -template -struct TestingIndices -{ - template - struct TestIndex - { - static constexpr int index = INDEX; - }; - template - static auto GenerateTypes(std::integer_sequence) - { - return ::testing::Types...>{}; - } - // testing::Types sequence of TestIndex types. - using Types = decltype(GenerateTypes(std::make_integer_sequence{})); -}; - -// A typed test suite so we can instantiate all the kernel builders. -template -class ConvBuilderTest : public ::testing::Test -{ - protected: - static constexpr int N = T::index; - static constexpr const std::string_view& NAME = TEST_CASES[N].name; - static constexpr auto& ALGORITHM = TEST_CASES[N].algorithm; - static constexpr const std::string_view& EXPECTED_TYPE = TEST_CASES[N].expected_type; -}; - -struct TestNameGenerator -{ - template - static std::string GetName(int index) - { - return std::to_string(index) + "." + std::string(TEST_CASES[index].name); - } -}; - -TYPED_TEST_SUITE(ConvBuilderTest, TestingIndices::Types, TestNameGenerator); - -// General test case, instantiated for each test case. -TYPED_TEST(ConvBuilderTest, TestInstance) -{ - static constexpr const FwdConvAlgorithm& ALGORITHM = ConvBuilderTest::ALGORITHM; - using Builder = ckb::ConvBuilder; - EXPECT_EQ(Builder::Instance::TypeString(), ConvBuilderTest::EXPECTED_TYPE); - const auto& tp = ALGORITHM.tuning_params; - EXPECT_EQ(Builder::factory::TUNING.ak1, tp.ak1); - EXPECT_EQ(Builder::factory::TUNING.bk1, tp.bk1); - const auto& tcla = ALGORITHM.block_transfer.thread_cluster_lengths_a; - EXPECT_EQ(Builder::factory::A_BLOCK_TRANSFER.thread_cluster_lengths[0], tcla.k0); - EXPECT_EQ(Builder::factory::A_BLOCK_TRANSFER.thread_cluster_lengths[1], tcla.m); - EXPECT_EQ(Builder::factory::A_BLOCK_TRANSFER.thread_cluster_lengths[2], tcla.k1); - const auto& tclb = ALGORITHM.block_transfer.thread_cluster_lengths_b; - EXPECT_EQ(Builder::factory::B_BLOCK_TRANSFER.thread_cluster_lengths[0], tclb.k0); - EXPECT_EQ(Builder::factory::B_BLOCK_TRANSFER.thread_cluster_lengths[1], tclb.n); - EXPECT_EQ(Builder::factory::B_BLOCK_TRANSFER.thread_cluster_lengths[2], tclb.k1); - const auto& tclc = ALGORITHM.block_transfer.thread_cluster_lengths_c; - EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[0], tclc.m_block); - EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[1], tclc.m_wave_per_xdl); - EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[2], tclc.n_block); - EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[3], tclc.n_wave_per_xdl); -} - } // namespace diff --git a/experimental/builder/test/test_conv_instances.cpp b/experimental/builder/test/test_conv_instances.cpp new file mode 100644 index 0000000000..e02a7034df --- /dev/null +++ b/experimental/builder/test/test_conv_instances.cpp @@ -0,0 +1,145 @@ +#include + +#include + +namespace { + +namespace ckb = ck_tile::builder; + +struct FwdConvSignature +{ + static constexpr int SPATIAL_DIM = 2; + static constexpr auto DIRECTION = ckb::ConvDirection::Forward; + static constexpr auto LAYOUT = ckb::GroupConvLayout::NHWGC_GKYXC_NHWGK; + static constexpr auto DATA_TYPE = ckb::DataType::FP16; +}; +static_assert(ckb::ConvSignature); + +static constexpr char API_VERSION[] = "0.1.0"; + +struct FwdConvAlgorithm +{ + ckb::ThreadBlock thread_block; + ckb::ConvTuningParams tuning_params; + struct BlockTransfer + { + ckb::BlockATransferLengthsInfo thread_cluster_lengths_a; + ckb::BlockBTransferLengthsInfo thread_cluster_lengths_b; + ckb::BlockCTransferLengthsInfo thread_cluster_lengths_c; + } block_transfer; +}; +static_assert(ckb::ConvAlgorithm); +static_assert(ckb::HasThreadBlockInfo); +static_assert(ckb::HasConvTuningInfo); +static_assert(ckb::HasABlockTransferInfo); +static_assert(ckb::HasBBlockTransferInfo); +static_assert(ckb::HasCBlockTransferInfo); + +struct TestCase +{ + std::string_view name; + FwdConvAlgorithm algorithm; + std::string_view expected_type; +}; + +// Test cases to drive the typed test suite. +constexpr std::array TEST_CASES = { + TestCase{ + .name = "ConvFwdXdlBf16CompInstances2x_0", + .algorithm = + {.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 128, .k = 64}}, + .tuning_params{.ak1 = 16, .bk1 = 16, .m_xdl_per_wave = 2, .n_xdl_per_wave = 2}, + .block_transfer{ + .thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1}, + .thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1}, + .thread_cluster_lengths_c = + {.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8}, + }}, + .expected_type = + "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 128, 64, Default, 32, 32, " + "2, 2, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>", + }, + TestCase{ + .name = "GroupedConvFwdXdlBf16CompInstance0", + .algorithm = {.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 256, .k = 32}}, + .tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4}, + .block_transfer{ + .thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1}, + .thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1}, + .thread_cluster_lengths_c = {.m_block = 1, + .m_wave_per_xdl = 32, + .n_block = 1, + .n_wave_per_xdl = 8}, + }}, + .expected_type = + "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, " + "4, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>", + }}; + +static constexpr int NUM_TEST_CASES = std::tuple_size_v; + +// Helper to generate testing::Types, TestIndex<1>, ..., TestIndex>. +template +struct TestingIndices +{ + template + struct TestIndex + { + static constexpr int index = INDEX; + }; + template + static auto GenerateTypes(std::integer_sequence) + { + return ::testing::Types...>{}; + } + // testing::Types sequence of TestIndex types. + using Types = decltype(GenerateTypes(std::make_integer_sequence{})); +}; + +// A typed test suite so we can instantiate all the kernel builders. +template +class ConvBuilderInstancesTest : public ::testing::Test +{ + protected: + static constexpr int N = T::index; + static constexpr const std::string_view& NAME = TEST_CASES[N].name; + static constexpr auto& ALGORITHM = TEST_CASES[N].algorithm; + static constexpr const std::string_view& EXPECTED_TYPE = TEST_CASES[N].expected_type; +}; + +struct TestNameGenerator +{ + template + static std::string GetName(int index) + { + return std::to_string(index) + "." + std::string(TEST_CASES[index].name); + } +}; + +TYPED_TEST_SUITE(ConvBuilderInstancesTest, TestingIndices::Types, TestNameGenerator); + +// General test case, instantiated for each test case. +TYPED_TEST(ConvBuilderInstancesTest, KernelParamsConfigured) +{ + static constexpr const FwdConvAlgorithm& ALGORITHM = ConvBuilderInstancesTest::ALGORITHM; + using Builder = ckb::ConvBuilder; + EXPECT_EQ(Builder::Instance::TypeString(), ConvBuilderInstancesTest::EXPECTED_TYPE); + const auto& tp = ALGORITHM.tuning_params; + EXPECT_EQ(Builder::factory::TUNING.ak1, tp.ak1); + EXPECT_EQ(Builder::factory::TUNING.bk1, tp.bk1); + const auto& tcla = ALGORITHM.block_transfer.thread_cluster_lengths_a; + EXPECT_EQ(Builder::factory::A_BLOCK_TRANSFER.thread_cluster_lengths[0], tcla.k0); + EXPECT_EQ(Builder::factory::A_BLOCK_TRANSFER.thread_cluster_lengths[1], tcla.m); + EXPECT_EQ(Builder::factory::A_BLOCK_TRANSFER.thread_cluster_lengths[2], tcla.k1); + const auto& tclb = ALGORITHM.block_transfer.thread_cluster_lengths_b; + EXPECT_EQ(Builder::factory::B_BLOCK_TRANSFER.thread_cluster_lengths[0], tclb.k0); + EXPECT_EQ(Builder::factory::B_BLOCK_TRANSFER.thread_cluster_lengths[1], tclb.n); + EXPECT_EQ(Builder::factory::B_BLOCK_TRANSFER.thread_cluster_lengths[2], tclb.k1); + const auto& tclc = ALGORITHM.block_transfer.thread_cluster_lengths_c; + EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[0], tclc.m_block); + EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[1], tclc.m_wave_per_xdl); + EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[2], tclc.n_block); + EXPECT_EQ(Builder::factory::C_BLOCK_TRANSFER.thread_cluster_lengths[3], tclc.n_wave_per_xdl); +} + +} // namespace