From cee90b800ee83eb290a3b7bfd26fe08c7aa96905 Mon Sep 17 00:00:00 2001 From: John Shumway Date: Thu, 28 Aug 2025 02:01:46 +0000 Subject: [PATCH] Fix test files for convolution builder. --- experimental/builder/test/CMakeLists.txt | 7 +- experimental/builder/test/builder.cpp | 88 ------------------- .../builder/test/test_conv_builder.cpp | 44 ++++++++++ 3 files changed, 45 insertions(+), 94 deletions(-) delete mode 100644 experimental/builder/test/builder.cpp create mode 100644 experimental/builder/test/test_conv_builder.cpp diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index 5215c66996..f9c0b0cb84 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -1,13 +1,8 @@ set(CMAKE_CXX_STANDARD 20) -add_executable(gemm_example gemm_example.cpp) -target_include_directories(gemm_example PRIVATE - "${PROJECT_SOURCE_DIR}/include" -) - include(gtest) -add_executable(test_conv_builder builder.cpp) +add_executable(test_conv_builder test_conv_builder.cpp) target_include_directories(test_conv_builder PRIVATE "${PROJECT_SOURCE_DIR}/experimental/builder/include" "${PROJECT_SOURCE_DIR}/include" diff --git a/experimental/builder/test/builder.cpp b/experimental/builder/test/builder.cpp deleted file mode 100644 index 4d4c96698d..0000000000 --- a/experimental/builder/test/builder.cpp +++ /dev/null @@ -1,88 +0,0 @@ -#include - -#include - -namespace { - -namespace ckb = ck_tile::builder; - -// Example of kernel description for Forward Conv with default settings -struct GroupedConvFwdXdlImplicitGemm : public GroupedConvBaseXdlV1 -{ - static constexpr ConvolutionDirection ConvolutionDirection_ = ConvolutionDirection::Forward; - static constexpr ElementwiseOperation ElementwiseOperation_ = ElementwiseOperation::Bias; -}; - -// Example of kernel description for Backward Weight Conv with default settings and Split K Two -// Stage -struct GroupedConvBwdWeightXdlImplicitGemmTwoStage : public GroupedConvBaseXdlV1 -{ - [[maybe_unused]] static constexpr ConvolutionDirection ConvolutionDirection_ = - ConvolutionDirection::BackwardWeight; - [[maybe_unused]] static constexpr SplitKSupport SplitKSupport_ = - SplitKSupport::SupportedTwoStage; -}; - -struct Implementation16x16 : ImplementationDefaultV1 -{ - static constexpr ck::index_t BlockSize_ = 64; - static constexpr auto TileSizes_ = std::make_tuple(16, 16, 32); - static constexpr ck::index_t K1_ = 8; - static constexpr MFMAInstructionSize MFMAInstructionSize_ = MFMAInstructionSize::M16N16; - static constexpr auto XdlPerWave_ = std::make_tuple(16, 16); - static constexpr auto GlobalTransferVectorSize_ = std::make_tuple(1, 1, 1); - static constexpr auto LDSStoreVectorSize_ = std::make_tuple(4, 4); -}; - -struct ProblemBF16NHWGC : public BF16ProblemBaseV1, public NHWGCProblemBaseV1 -{ -}; - -TEST(ConvBuilderTest, TestBuilderV0_0_0) -{ - ConvolutionBuilder - builder_fwd; - - EXPECT_EQ(builder_fwd.GetInstanceName(), - "GroupedConvFwdMultipleABD_Xdl_CShuffle<64, 16, 16, 32, Default, 8, 16x16, 16, 16, " - "1, 4, 1, 4, 1, Intrawave, v1, 1>"); - // It would be nice if this worked, but it fails. - // [[maybe_unused]] auto instance = builder_fwd.GetInstance(); -} - -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; -}; - -TEST(ConvBuilderTest, TestSignature) -{ - static_assert(ckb::ConvSignature); - SUCCEED(); -} - -struct FwdConvAlgorithm -{ - // -}; - -TEST(ConvBuilderTest, TestAlgorithm) -{ - static_assert(ckb::ConvAlgorithm); - SUCCEED(); -} - -static constexpr char API_VERSION[] = "0.1.0"; -using FwdConvBuilder = ckb::ConvBuilder; - -TEST(ConvBuilderTest, TestKernel) -{ - EXPECT_EQ( - FwdConvBuilder::Instance::TypeString(), - "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, " - "8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>"); -} -} // namespace diff --git a/experimental/builder/test/test_conv_builder.cpp b/experimental/builder/test/test_conv_builder.cpp new file mode 100644 index 0000000000..d761b277ba --- /dev/null +++ b/experimental/builder/test/test_conv_builder.cpp @@ -0,0 +1,44 @@ +#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; +}; + +TEST(ConvBuilderTest, TestSignature) +{ + static_assert(ckb::ConvSignature); + SUCCEED(); +} + +struct FwdConvAlgorithm +{ + // TODO: Add algorithm info. +}; + +TEST(ConvBuilderTest, TestAlgorithm) +{ + static_assert(ckb::ConvAlgorithm); + SUCCEED(); +} + +static constexpr char API_VERSION[] = "0.1.0"; +using FwdConvBuilder = ckb::ConvBuilder; + +TEST(ConvBuilderTest, TestInstance) +{ + EXPECT_EQ( + FwdConvBuilder::Instance::TypeString(), + "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, " + "8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>"); +} +} // namespace