From 0ac48abe611acffeea291d0f014138f77c07eb06 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Tue, 4 Nov 2025 08:58:25 +0000 Subject: [PATCH] Improve ckb fwd conv instance tests. --- experimental/builder/test/CMakeLists.txt | 3 +- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 15 ++--- .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 12 ++-- .../test/conv/test_ckb_conv_fwd_1d_i8.cpp | 7 +- .../test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 13 +++- .../test/conv/test_ckb_conv_fwd_2d_fp16.cpp | 8 ++- .../test/conv/test_ckb_conv_fwd_2d_fp32.cpp | 8 ++- .../test/conv/test_ckb_conv_fwd_3d_bf16.cpp | 8 ++- .../test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 8 ++- .../test/conv/test_ckb_conv_fwd_3d_fp32.cpp | 8 ++- .../test/utils/ckb_conv_test_common.hpp | 64 ------------------- .../test/utils/ckb_conv_test_utils.hpp | 37 +++++++++++ 12 files changed, 96 insertions(+), 95 deletions(-) delete mode 100644 experimental/builder/test/utils/ckb_conv_test_common.hpp create mode 100644 experimental/builder/test/utils/ckb_conv_test_utils.hpp diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index 26a666a805..846225883c 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -43,7 +43,8 @@ add_ck_builder_test(test_ckb_build_fwd_instances conv/test_ckb_conv_fwd_2d_fp32.cpp conv/test_ckb_conv_fwd_3d_bf16.cpp conv/test_ckb_conv_fwd_3d_fp16.cpp - conv/test_ckb_conv_fwd_3d_fp32.cpp) + conv/test_ckb_conv_fwd_3d_fp32.cpp + ) function(add_ck_factory_test test_name) add_ck_builder_test(${test_name} ${ARGN}) diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp index 0cc1e83810..8792882314 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -28,13 +28,12 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v2_intrawave}; using Builder = ConvBuilder; - - const auto& asserts = InstanceNameAsserts{} - .StartsWith("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3") - .Contains("Filter1x1Stride1Pad0") - .Contains("BlkGemmPipelineVersion: v2"); - - run_test(asserts); + run_test({ + "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 256, 256, 32", + "Filter1x1Stride1Pad0", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v2"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp index 382e40e5a8..651e2b78a2 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -29,12 +29,10 @@ TEST(FwdConvInstances, .loop_scheduler = LoopScheduler::DEFAULT}; using Builder = ConvBuilder; - - const auto& asserts = InstanceNameAsserts{} - .StartsWith("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle") - .Contains("Default"); - - run_test(asserts); + run_test({ + "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle", + "64, 64, 32, 32", + "Default"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp index b97e9fbd64..978532dfd4 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -28,7 +28,10 @@ TEST(FwdConvInstances, .loop_scheduler = LoopScheduler::DEFAULT}; using Builder = ConvBuilder; - run_test(); + run_test({ + "DeviceGroupedConvFwdMultipleD_Wmma_CShuffle", + "128, 64, 64, 64", + "Default"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp index 5094978219..e04d5a72c2 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace{ @@ -27,7 +27,12 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v1_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({ + "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 256, 256, 32", + "Default", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v1"}); } // 2D BF16 NHWGC (channels-last) with Pipeline V5 and FILTER_3x3 @@ -52,7 +57,9 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v5_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "Filter3x3", + "BlkGemmPipelineVersion: v5"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp index 1cf861cd02..d5953e7315 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -26,7 +26,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v3_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 256, 256, 32", + "Filter1x1Pad0", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v3"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp index 86e8b8e453..f1ea331bc5 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -26,7 +26,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v4_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 128, 128, 32", + "Filter1x1Stride1Pad0", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v4"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp index bef4cbcea1..86e557fd92 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -27,7 +27,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v3_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 256, 256, 32", + "Default", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v3"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp index 63bef01864..8590b7e93f 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -27,7 +27,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v4_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 128, 128, 32", + "Filter1x1Pad0", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v4"}); } } // namespace diff --git a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp index 6713509b2d..030e000144 100644 --- a/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp +++ b/experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp @@ -1,5 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_utils.hpp" namespace { @@ -27,7 +27,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v1_intrawave}; using Builder = ConvBuilder; - run_test(); + run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", + "256, 256, 256, 32", + "Filter1x1Pad0", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v1"}); } } // namespace diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_common.hpp deleted file mode 100644 index 7fc9b6aeee..0000000000 --- a/experimental/builder/test/utils/ckb_conv_test_common.hpp +++ /dev/null @@ -1,64 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. - -#pragma once - -#include -#include -#include "impl/conv_algorithm_types.hpp" -#include "impl/conv_signature_types.hpp" -#include "ck_tile/builder/conv_builder.hpp" - -namespace ck_tile::builder::test_utils { - -using namespace ck_tile::builder; -using namespace test; - -class InstanceNameAsserts -{ -public: - InstanceNameAsserts& StartsWith(const char* prefix) - { - prefixes_.push_back(std::string(prefix)); - return *this; - } - - InstanceNameAsserts& Contains(const char* substring) - { - substrings_.push_back(std::string(substring)); - return *this; - } - - void Check(const std::string& kernel_string) const - { - for (const auto& prefix : prefixes_) - { - EXPECT_THAT(kernel_string, ::testing::StartsWith(prefix)); - } - for (const auto& substr : substrings_) - { - EXPECT_THAT(kernel_string, ::testing::HasSubstr(substr)); - } - } -private: - std::vector prefixes_; - std::vector substrings_; -}; - -// Common test implementation -template -constexpr void run_test(const InstanceNameAsserts& asserts) -{ - auto instance = typename Builder::Instance{}; - - const auto kernel_string = instance.GetTypeString(); - std::cout << "Generated kernel: " << kernel_string << std::endl; - EXPECT_GT(kernel_string.size(), 0); - - const auto invoker_ptr = instance.MakeInvokerPointer(); - EXPECT_NE(invoker_ptr, nullptr); - - asserts.Check(kernel_string); -} - -} // namespace ck_tile::builder::test_utils diff --git a/experimental/builder/test/utils/ckb_conv_test_utils.hpp b/experimental/builder/test/utils/ckb_conv_test_utils.hpp new file mode 100644 index 0000000000..c04117e8d6 --- /dev/null +++ b/experimental/builder/test/utils/ckb_conv_test_utils.hpp @@ -0,0 +1,37 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include +#include +#include "impl/conv_algorithm_types.hpp" +#include "impl/conv_signature_types.hpp" +#include "ck_tile/builder/conv_builder.hpp" + +namespace ck_tile::builder::test_utils { + +using namespace ck_tile::builder; +using namespace test; + + +// Common test implementation +template +constexpr void run_test(const std::vector& kernel_instance_components) +{ + auto instance = typename Builder::Instance{}; + + const auto kernel_string = instance.GetTypeString(); + std::cout << "Generated kernel: " << kernel_string << std::endl; + EXPECT_GT(kernel_string.size(), 0); + + const auto invoker_ptr = instance.MakeInvokerPointer(); + EXPECT_NE(invoker_ptr, nullptr); + + for (const auto& component : kernel_instance_components) + { + EXPECT_THAT(kernel_string, ::testing::HasSubstr(component)); + } +} + +} // namespace ck_tile::builder::test_utils