diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index 6f25d1d7f4..ca8a2b1adb 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -21,6 +21,7 @@ add_ck_builder_test(test_conv_builder test_conv_grp_fwd_2d.cpp test_conv_grp_fwd_3d.cpp test_conv_grp_bwd_2d.cpp - test_conv_fwd_instances.cpp) + test_conv_fwd_instances.cpp + test_conv_bwd_instances.cpp) add_ck_builder_test(test_builder_utils test_builder_utils.cpp) \ No newline at end of file diff --git a/experimental/builder/test/test_conv_bwd_instances.cpp b/experimental/builder/test/test_conv_bwd_instances.cpp new file mode 100644 index 0000000000..244707abd9 --- /dev/null +++ b/experimental/builder/test/test_conv_bwd_instances.cpp @@ -0,0 +1,177 @@ +// This test is designed to verify that the ConvBuilder can instantiate the same +// kernel classes that are used in production code. Production code may have +// hundreds or thousands of kernel instances, so this test uses a GTest typed +// test suite to efficiently test a representative set of these kernel examples. +// Each test case defines a specific convolution algorithm configuration and the +// expected kernel type string that the builder should generate. +#include + +#include + +namespace { + +namespace ckb = ck_tile::builder; +using P = ckb::BlockGemmPipelineVersion; + +// Defines the signature of the convolution operation to be tested. +// This includes dimensionality, direction, data layout, and data type. +struct ConvSignature +{ + int spatial_dim = 2; + ckb::ConvDirection direction = ckb::ConvDirection::BACKWARD_DATA; + ckb::GroupConvLayout layout = ckb::GroupConvLayout::CHANNELS_LAST; + ckb::DataType data_type = ckb::DataType::FP16; +}; +static_assert(ckb::ConvSignatureDescriptor); + +// Defines the tunable algorithmic parameters for the convolution kernel. +// This includes thread block configuration, tuning parameters, data transfer +// settings, and the GEMM pipeline version. +struct FwdConvAlgorithm +{ + ckb::ThreadBlock thread_block; + ckb::ConvTuningParams tuning_params; + struct BlockTransfer + { + ckb::BlockATransferLengths thread_cluster_dims_a; + ckb::BlockBTransferLengths thread_cluster_dims_b; + ckb::BlockCTransferLengths thread_cluster_dims_c; + } block_transfer; +}; +static_assert(ckb::ConvAlgorithmDescriptor); +static_assert(ckb::SpecifiesThreadBlock); +static_assert(ckb::SpecifiesConvTuning); +static_assert(ckb::SpecifiesBlockATransfer); +static_assert(ckb::SpecifiesBlockBTransfer); +static_assert(ckb::SpecifiesBlockCTransfer); + +// A container for a single test case, bundling a descriptive name, the +// algorithm configuration, and the expected generated kernel type string. +struct TestCase +{ + std::string_view name; + FwdConvAlgorithm algorithm; + std::string_view expected_type; +}; + +// Helper function to set the sub_matrix size. +constexpr ckb::ThreadBlock set_submatrix(int m, int n, int k) +{ + return {.block_size = 256, .submatrix = {.m = m, .n = n, .k = k}}; +} + +// An array of test cases that drive the typed test suite. Each entry +// represents a unique kernel instance to be verified. +constexpr std::array TEST_CASES = { + TestCase{ + .name = "ConvFwdXdlBf16CompInstances2x_0", + .algorithm = + {.thread_block = set_submatrix(256, 128, 64), + .tuning_params = {.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 2}, + .block_transfer = {.thread_cluster_dims_a = {.k0 = 4, .m = 16, .k1 = 1}, + .thread_cluster_dims_b = {.k0 = 4, .n = 8, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 16, + .n_block = 1, + .n_wave_per_xdl = 4}}}, + .expected_type = + "DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1<256, 256, 128, 64, 8, 8, Default, " + "16, 16, 1, 4, 8, 8, 1, 1, TransposeTransferInScalarPerVectorAligned: 1, " + "TransposeTransferOutScalarPerVectorAligned: 1>", + }, +}; + +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 that will be instantiated for each type in TestingIndices::Types. +// This creates a separate test for each entry in the TEST_CASES array, allowing +// GTest to run and report on them individually. +template +class ConvBuilderBwdInstancesTest : 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; +}; + +// Custom test name generator to provide more descriptive names for each +// typed test instance, incorporating the index and the name from the TestCase. +struct TestNameGenerator +{ + template + static std::string GetName(int index) + { + return std::to_string(index) + "." + std::string(TEST_CASES[index].name); + } +}; + +TYPED_TEST_SUITE(ConvBuilderBwdInstancesTest, + TestingIndices::Types, + TestNameGenerator); + +// This is the body of the typed test. It will be executed for each TestCase. +// It verifies that the ConvBuilder, when configured with a specific algorithm, +// generates the correct kernel type string and correctly configures the +// underlying factory parameters. +TYPED_TEST(ConvBuilderBwdInstancesTest, KernelParamsConfigured) +{ + static constexpr const FwdConvAlgorithm& ALGORITHM = + ConvBuilderBwdInstancesTest::ALGORITHM; + static constexpr const ConvSignature SIGNATURE; + using Builder = ckb::ConvBuilder; + EXPECT_EQ(Builder::Instance::TypeString(), + ConvBuilderBwdInstancesTest::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& tcda = ALGORITHM.block_transfer.thread_cluster_dims_a; + EXPECT_EQ(Builder::Factory::A_BLOCK_TRANSFER.thread_cluster_dims[0], tcda.k0); + EXPECT_EQ(Builder::Factory::A_BLOCK_TRANSFER.thread_cluster_dims[1], tcda.m); + EXPECT_EQ(Builder::Factory::A_BLOCK_TRANSFER.thread_cluster_dims[2], tcda.k1); + const auto& tcdb = ALGORITHM.block_transfer.thread_cluster_dims_b; + EXPECT_EQ(Builder::Factory::B_BLOCK_TRANSFER.thread_cluster_dims[0], tcdb.k0); + EXPECT_EQ(Builder::Factory::B_BLOCK_TRANSFER.thread_cluster_dims[1], tcdb.n); + EXPECT_EQ(Builder::Factory::B_BLOCK_TRANSFER.thread_cluster_dims[2], tcdb.k1); + const auto& tcdc = ALGORITHM.block_transfer.thread_cluster_dims_c; + EXPECT_EQ(Builder::Factory::C_BLOCK_TRANSFER.thread_cluster_dims[0], tcdc.m_block); + EXPECT_EQ(Builder::Factory::C_BLOCK_TRANSFER.thread_cluster_dims[1], tcdc.m_wave_per_xdl); + EXPECT_EQ(Builder::Factory::C_BLOCK_TRANSFER.thread_cluster_dims[2], tcdc.n_block); + EXPECT_EQ(Builder::Factory::C_BLOCK_TRANSFER.thread_cluster_dims[3], tcdc.n_wave_per_xdl); +} + +// A standard GTest to ensure that all `expected_type` strings in the +// TEST_CASES array are unique. This helps prevent copy-paste errors and +// ensures that each test case is meaningful. +TEST(ConvBuilderBwdInstancesTest, TypeStringsAreUnique) +{ + std::set strings; + for(int i = 0; i < NUM_TEST_CASES; ++i) + { + const auto& [iter, inserted] = strings.insert(std::string(TEST_CASES[i].expected_type)); + EXPECT_TRUE(inserted) << "Duplicate expected_string " << *iter; + } + EXPECT_EQ(strings.size(), NUM_TEST_CASES) + << "Found fewer unique expected_strings than test cases"; +} + +} // namespace diff --git a/experimental/builder/test/test_conv_fwd_instances.cpp b/experimental/builder/test/test_conv_fwd_instances.cpp index 3a614ad297..0c337d155d 100644 --- a/experimental/builder/test/test_conv_fwd_instances.cpp +++ b/experimental/builder/test/test_conv_fwd_instances.cpp @@ -223,7 +223,7 @@ struct TestingIndices // This creates a separate test for each entry in the TEST_CASES array, allowing // GTest to run and report on them individually. template -class ConvBuilderInstancesTest : public ::testing::Test +class ConvBuilderFwdInstancesTest : public ::testing::Test { protected: static constexpr int N = T::index; @@ -243,7 +243,7 @@ struct TestNameGenerator } }; -TYPED_TEST_SUITE(ConvBuilderInstancesTest, +TYPED_TEST_SUITE(ConvBuilderFwdInstancesTest, TestingIndices::Types, TestNameGenerator); @@ -251,13 +251,13 @@ TYPED_TEST_SUITE(ConvBuilderInstancesTest, // It verifies that the ConvBuilder, when configured with a specific algorithm, // generates the correct kernel type string and correctly configures the // underlying factory parameters. -TYPED_TEST(ConvBuilderInstancesTest, KernelParamsConfigured) +TYPED_TEST(ConvBuilderFwdInstancesTest, KernelParamsConfigured) { static constexpr const FwdConvAlgorithm& ALGORITHM = - ConvBuilderInstancesTest::ALGORITHM; + ConvBuilderFwdInstancesTest::ALGORITHM; static constexpr const ConvSignature SIGNATURE; using Builder = ckb::ConvBuilder; - EXPECT_EQ(Builder::Instance::TypeString(), ConvBuilderInstancesTest::EXPECTED_TYPE); + EXPECT_EQ(Builder::Instance::TypeString(), ConvBuilderFwdInstancesTest::EXPECTED_TYPE); const auto& tp = ALGORITHM.tuning_params; EXPECT_EQ(Builder::Factory::TUNING.ak1, tp.ak1); EXPECT_EQ(Builder::Factory::TUNING.bk1, tp.bk1); @@ -279,7 +279,7 @@ TYPED_TEST(ConvBuilderInstancesTest, KernelParamsConfigured) // A standard GTest to ensure that all `expected_type` strings in the // TEST_CASES array are unique. This helps prevent copy-paste errors and // ensures that each test case is meaningful. -TEST(ConvBuilderInstancesTest, TypeStringsAreUnique) +TEST(ConvBuilderFwdInstancesTest, TypeStringsAreUnique) { std::set strings; for(int i = 0; i < NUM_TEST_CASES; ++i)