From 6ded2ab92a541956aa5895f3095ebc9a99da8423 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Mon, 3 Nov 2025 14:06:45 +0000 Subject: [PATCH 01/20] Remove code duplication in fwd conv builder tests. --- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 22 +- .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 22 +- .../test/conv/test_ckb_conv_fwd_1d_i8.cpp | 21 +- .../test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 36 +- .../test/conv/test_ckb_conv_fwd_2d_fp16.cpp | 22 +- .../test/conv/test_ckb_conv_fwd_2d_fp32.cpp | 22 +- .../test/conv/test_ckb_conv_fwd_3d_bf16.cpp | 21 +- .../test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 22 +- .../test/conv/test_ckb_conv_fwd_3d_fp32.cpp | 22 +- .../test/utils/ckb_conv_test_common.hpp | 332 +++++++----------- 10 files changed, 242 insertions(+), 300 deletions(-) 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 472c43438d..ab83f4d109 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,9 +1,9 @@ #include "utils/ckb_conv_test_common.hpp" +namespace { + using namespace ck_tile::builder::test_utils; -namespace ck_tile::builder::testing { - // 1D BF16 (channels-first) with Pipeline V2 and FILTER_1X1_STRIDE1_PAD0 specialization and SCALE // elementwise op TEST(FwdConvInstances, @@ -18,14 +18,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 256, .n = 256, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_256x256x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v2_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< - FwdConvSignature, - FwdThreadBlock, - BlockGemmPipelineVersion::V2, - ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0>(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 3f840ba2b0..8e6abdf6f5 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace { -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; // 1D FP16 (channels-last) with DEFAULT specialization TEST(FwdConvInstances, @@ -17,12 +17,18 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 64, - .tile_size = {.m = 64, .n = 32, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_64x32x32, + .gridwise_gemm = FwdGemmParams_Xdl_2x1_per_wave, + .block_transfer = FwdBlockTransfer_4x16x1, + .fwd_specialization = ConvFwdSpecialization::DEFAULT, + .gemm_specialization = GemmSpecialization::MNKPadding, + .num_gemm_k_prefetch_stages = 1, + .num_groups_to_merge = 2, + .loop_scheduler = LoopScheduler::DEFAULT}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 1819cca728..0bfadfd21f 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace { -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; // 1D I8 (channels-last) with and DEFAULT specialization TEST(FwdConvInstances, @@ -17,12 +17,17 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleD_Wmma_CShuffle}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 128, - .tile_size = {.m = 64, .n = 64, .k = 64}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_64x64x64, + .gridwise_gemm = FwdGemmParams_Wmma_2x1_per_wave, + .block_transfer = FwdBlockTransfer_4x32x1, + .fwd_specialization = ConvFwdSpecialization::DEFAULT, + .gemm_specialization = GemmSpecialization::MNKPadding, + .num_gemm_k_prefetch_stages = 1, + .loop_scheduler = LoopScheduler::DEFAULT}; - run_test_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 b9969f7e95..a4c3e06858 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace{ -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; // 2D BF16 NHWGC (channels-last) with Pipeline V1 and DEFAULT TEST(FwdConvInstances, @@ -17,13 +17,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 256, .n = 256, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_256x256x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::DEFAULT, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v1_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3(); + using Builder = ConvBuilder; + run_test(); } // 2D BF16 NHWGC (channels-last) with Pipeline V5 and FILTER_3x3 @@ -39,13 +42,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 256, .n = 256, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_256x256x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::FILTER_3x3, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v5_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 cd5186cc10..7acc0851cd 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace { -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; TEST(FwdConvInstances, Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_2D_FP16_GNHWC) @@ -16,14 +16,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 256, .n = 256, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_256x256x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::FILTER_1X1_PAD0, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v3_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< - FwdConvSignature, - FwdThreadBlock, - BlockGemmPipelineVersion::V3, - ConvFwdSpecialization::FILTER_1X1_PAD0>(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 584e0ab182..e4aa8775ee 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace { -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; TEST(FwdConvInstances, Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_2D_FP32_NGCHW_GKCYX) @@ -16,14 +16,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 128, .n = 128, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_128x128x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v4_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< - FwdConvSignature, - FwdThreadBlock, - BlockGemmPipelineVersion::V4, - ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0>(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 17caf98457..75dc313f34 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace { -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; // 3D BF16 GNDHWC (group-first, channels-last) with Pipeline V3 and DEFAULT TEST(FwdConvInstances, @@ -17,13 +17,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 256, .n = 256, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_256x256x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::DEFAULT, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v3_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 ec4649a6ff..c0cfccbd03 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,9 +1,9 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; - namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; + // 3D FP16 NDHWGC (channels-last) with Pipeline V4 and FILTER_1X1_PAD0 TEST(FwdConvInstances, Create_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3_Instance_3D_FP16_NDHWGC_ChannelsLast) @@ -17,14 +17,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 128, .n = 128, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_128x128x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::FILTER_1X1_PAD0, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v4_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< - FwdConvSignature, - FwdThreadBlock, - BlockGemmPipelineVersion::V4, - ConvFwdSpecialization::FILTER_1X1_PAD0>(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // 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 393ea9206d..09006948f0 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,8 +1,8 @@ #include "utils/ckb_conv_test_common.hpp" -using namespace ck_tile::builder::test_utils; +namespace { -namespace ck_tile::builder::testing { +using namespace ck_tile::builder::test_utils; // 3D FP32 NGCDHW (channels-first) with Pipeline V1 and FILTER_1X1_PAD0 TEST(FwdConvInstances, @@ -17,14 +17,16 @@ TEST(FwdConvInstances, .device_operation = FwdGroupConvDeviceOperation::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3}; - constexpr ThreadBlock FwdThreadBlock{.block_size = 256, - .tile_size = {.m = 128, .n = 128, .k = 32}}; + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ + .thread_block = FwdThreadBlock_256x256x32, + .gridwise_gemm = FwdGemmParams_Xdl_4x4_per_wave, + .block_transfer = FwdBlockTransfer_4x64_1, + .fwd_specialization = ConvFwdSpecialization::FILTER_1X1_PAD0, + .gemm_specialization = GemmSpecialization::MNKPadding, + .block_gemm = BlockGemmDesc_v1_intrawave}; - run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< - FwdConvSignature, - FwdThreadBlock, - BlockGemmPipelineVersion::V1, - ConvFwdSpecialization::FILTER_1X1_PAD0>(); + using Builder = ConvBuilder; + run_test(); } -} // namespace ck_tile::builder::testing +} // namespace diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_common.hpp index d18a008015..35d2a3ca6e 100644 --- a/experimental/builder/test/utils/ckb_conv_test_common.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_common.hpp @@ -13,224 +13,136 @@ namespace ck_tile::builder::test_utils { using namespace ck_tile::builder; using namespace test; +constexpr BlockTransferABC FwdBlockTransfer_4x64_1 {.block_transfer_a = {.k0 = 4, .m_n = 64, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 64, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 32, + .n_block = 1, + .n_wave_per_xdl = 8}, + .lds_transfer_a = {.src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = false}, + .lds_transfer_b = {.src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = false}, + .epilogue_c = {.m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, + .block_transfer_access_order_a = {1, 0, 2}, + .block_transfer_access_order_b = {1, 0, 2}, + .src_access_order_a = {1, 0, 2}, + .src_access_order_b = {1, 0, 2}}; + +constexpr BlockTransferABC FwdBlockTransfer_4x16x1 {.block_transfer_a = {.k0 = 4, .m_n = 16, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 16, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 16, + .n_block = 1, + .n_wave_per_xdl = 4}, + .lds_transfer_a = {.src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, + .lds_transfer_b = {.src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, + .epilogue_c = {.m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, + .block_transfer_access_order_a = {1, 0, 2}, + .block_transfer_access_order_b = {1, 0, 2}, + .src_access_order_a = {1, 0, 2}, + .src_access_order_b = {1, 0, 2}}; + +constexpr BlockTransferABC FwdBlockTransfer_4x32x1 {.block_transfer_a = {.k0 = 4, .m_n = 32, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 32, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 32, + .n_block = 1, + .n_wave_per_xdl = 4}, + .lds_transfer_a = {.src_vector_dim = 2, + .src_scalar_per_vector = 16, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, + .lds_transfer_b = {.src_vector_dim = 2, + .src_scalar_per_vector = 16, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, + .epilogue_c = {.m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, + .block_transfer_access_order_a = {1, 0, 2}, + .block_transfer_access_order_b = {1, 0, 2}, + .src_access_order_a = {1, 0, 2}, + .src_access_order_b = {1, 0, 2}}; + +constexpr GridwiseXdlGemm FwdGemmParams_Xdl_4x4_per_wave {.ak1 = 8, + .bk1 = 8, + .m_per_xdl = 32, + .n_per_xdl = 32, + .m_xdl_per_wave = 4, + .n_xdl_per_wave = 4}; + +constexpr GridwiseXdlGemm FwdGemmParams_Xdl_2x1_per_wave {.ak1 = 8, + .bk1 = 8, + .m_per_xdl = 32, + .n_per_xdl = 32, + .m_xdl_per_wave = 2, + .n_xdl_per_wave = 1}; + +constexpr GridwiseWmmaGemm FwdGemmParams_Wmma_2x1_per_wave {.k1 = 8, + .m_per_wmma = 32, + .n_per_wmma = 32, + .m_wmma_per_wave = 2, + .n_wmma_per_wave = 1, + .pipeline_version = GridwiseGemmPipelineVersion::V1}; + +constexpr ThreadBlock FwdThreadBlock_256x256x32 {.block_size = 256, + .tile_size = {.m = 256, .n = 256, .k = 32}}; + +constexpr ThreadBlock FwdThreadBlock_128x128x32 {.block_size = 256, + .tile_size = {.m = 128, .n = 128, .k = 32}}; + +constexpr ThreadBlock FwdThreadBlock_64x32x32 {.block_size = 64, + .tile_size = {.m = 64, .n = 32, .k = 32}}; + +constexpr ThreadBlock FwdThreadBlock_64x64x64 {.block_size = 128, + .tile_size = {.m = 64, .n = 64, .k = 64}}; + +constexpr BlockGemm BlockGemmDesc_v1_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V1, + .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + +constexpr BlockGemm BlockGemmDesc_v2_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V2, + .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + +constexpr BlockGemm BlockGemmDesc_v3_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V3, + .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + +constexpr BlockGemm BlockGemmDesc_v4_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V4, + .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + +constexpr BlockGemm BlockGemmDesc_v5_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V5, + .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + // Common test implementation -template -constexpr void run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3() +template +constexpr void run_test() { - constexpr GridwiseXdlGemm FwdGemmParams{.ak1 = 8, - .bk1 = 8, - .m_per_xdl = 32, - .n_per_xdl = 32, - .m_xdl_per_wave = 4, - .n_xdl_per_wave = 4}; - - constexpr BlockTransferABC FwdBlockTransfer{.block_transfer_a = {.k0 = 4, .m_n = 64, .k1 = 1}, - .block_transfer_b = {.k0 = 4, .m_n = 64, .k1 = 1}, - .thread_cluster_dims_c = {.m_block = 1, - .m_wave_per_xdl = 32, - .n_block = 1, - .n_wave_per_xdl = 8}, - .lds_transfer_a = {.src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = false}, - .lds_transfer_b = {.src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = false}, - .epilogue_c = {.m_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, - .block_transfer_access_order_a = {1, 0, 2}, - .block_transfer_access_order_b = {1, 0, 2}, - .src_access_order_a = {1, 0, 2}, - .src_access_order_b = {1, 0, 2}}; - - constexpr BlockGemm BlockGemmDesc = {.pipeline_version = FwdPipelineVersion, - .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; - - constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 FwdConvAlgorithm{ - .thread_block = FwdThreadBlock, - .gridwise_gemm = FwdGemmParams, - .block_transfer = FwdBlockTransfer, - .fwd_specialization = FwdConvSpecialization, - .gemm_specialization = GemmSpecialization::MNKPadding, - .block_gemm = BlockGemmDesc}; - - using Builder = ConvBuilder; - 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); - EXPECT_TRUE(kernel_string.starts_with("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3")); - - // Verify pipeline version is correct - if(FwdPipelineVersion == BlockGemmPipelineVersion::V1) - EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v1") != std::string::npos); - else if(FwdPipelineVersion == BlockGemmPipelineVersion::V3) - EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v3") != std::string::npos); - else if(FwdPipelineVersion == BlockGemmPipelineVersion::V4) - EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v4") != std::string::npos); - else if(FwdPipelineVersion == BlockGemmPipelineVersion::V5) - EXPECT_TRUE(kernel_string.find("BlkGemmPipelineVersion: v5") != std::string::npos); - - // Verify specialization is correct - if(FwdConvSpecialization == ConvFwdSpecialization::DEFAULT) - EXPECT_TRUE(kernel_string.find("Default") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_PAD0) - EXPECT_TRUE(kernel_string.find("Filter1x1Pad0") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0) - EXPECT_TRUE(kernel_string.find("Filter1x1Stride1Pad0") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_3x3) - EXPECT_TRUE(kernel_string.find("Filter3x3") != std::string::npos); - - const auto invoker_ptr = instance.MakeInvokerPointer(); - EXPECT_NE(invoker_ptr, nullptr); -} - -template -constexpr void run_test_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle() -{ - constexpr GridwiseXdlGemm FwdGemmParams{.ak1 = 8, - .bk1 = 8, - .m_per_xdl = 32, - .n_per_xdl = 32, - .m_xdl_per_wave = 2, - .n_xdl_per_wave = 1}; - - constexpr BlockTransferABC FwdBlockTransfer{.block_transfer_a = {.k0 = 4, .m_n = 16, .k1 = 1}, - .block_transfer_b = {.k0 = 4, .m_n = 16, .k1 = 1}, - .thread_cluster_dims_c = {.m_block = 1, - .m_wave_per_xdl = 16, - .n_block = 1, - .n_wave_per_xdl = 4}, - .lds_transfer_a = {.src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, - .lds_transfer_b = {.src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, - .epilogue_c = {.m_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, - .block_transfer_access_order_a = {1, 0, 2}, - .block_transfer_access_order_b = {1, 0, 2}, - .src_access_order_a = {1, 0, 2}, - .src_access_order_b = {1, 0, 2}}; - - constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle FwdConvAlgorithm{ - .thread_block = FwdThreadBlock, - .gridwise_gemm = FwdGemmParams, - .block_transfer = FwdBlockTransfer, - .fwd_specialization = FwdConvSpecialization, - .gemm_specialization = GemmSpecialization::MNKPadding, - .num_gemm_k_prefetch_stages = 1, - .num_groups_to_merge = 2, - .loop_scheduler = LoopScheduler::DEFAULT}; - - using Builder = ConvBuilder; - - 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); - - EXPECT_TRUE(kernel_string.starts_with("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle")); - - // Verify specialization is correct - if(FwdConvSpecialization == ConvFwdSpecialization::DEFAULT) - EXPECT_TRUE(kernel_string.find("Default") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_PAD0) - EXPECT_TRUE(kernel_string.find("Filter1x1Pad0") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0) - EXPECT_TRUE(kernel_string.find("Filter1x1Stride1Pad0") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_3x3) - EXPECT_TRUE(kernel_string.find("Filter3x3") != std::string::npos); - - const auto invoker_ptr = instance.MakeInvokerPointer(); - EXPECT_NE(invoker_ptr, nullptr); -} - -template -constexpr void run_test_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle() -{ - constexpr GridwiseWmmaGemm FwdGemmParams{.k1 = 8, - .m_per_wmma = 32, - .n_per_wmma = 32, - .m_wmma_per_wave = 2, - .n_wmma_per_wave = 1, - .pipeline_version = GridwiseGemmPipelineVersion::V1}; - - constexpr BlockTransferABC FwdBlockTransfer{.block_transfer_a = {.k0 = 4, .m_n = 32, .k1 = 1}, - .block_transfer_b = {.k0 = 4, .m_n = 32, .k1 = 1}, - .thread_cluster_dims_c = {.m_block = 1, - .m_wave_per_xdl = 32, - .n_block = 1, - .n_wave_per_xdl = 4}, - .lds_transfer_a = {.src_vector_dim = 2, - .src_scalar_per_vector = 16, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, - .lds_transfer_b = {.src_vector_dim = 2, - .src_scalar_per_vector = 16, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, - .epilogue_c = {.m_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, - .block_transfer_access_order_a = {1, 0, 2}, - .block_transfer_access_order_b = {1, 0, 2}, - .src_access_order_a = {1, 0, 2}, - .src_access_order_b = {1, 0, 2}}; - - constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle FwdConvAlgorithm{ - .thread_block = FwdThreadBlock, - .gridwise_gemm = FwdGemmParams, - .block_transfer = FwdBlockTransfer, - .fwd_specialization = FwdConvSpecialization, - .gemm_specialization = GemmSpecialization::MNKPadding, - .num_gemm_k_prefetch_stages = 1, - .loop_scheduler = LoopScheduler::DEFAULT}; - - using Builder = ConvBuilder; - - 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); - - EXPECT_TRUE(kernel_string.starts_with("DeviceGroupedConvFwdMultipleD_Wmma_CShuffle")); - - // Verify specialization is correct - if(FwdConvSpecialization == ConvFwdSpecialization::DEFAULT) - EXPECT_TRUE(kernel_string.find("Default") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_PAD0) - EXPECT_TRUE(kernel_string.find("Filter1x1Pad0") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0) - EXPECT_TRUE(kernel_string.find("Filter1x1Stride1Pad0") != std::string::npos); - else if(FwdConvSpecialization == ConvFwdSpecialization::FILTER_3x3) - EXPECT_TRUE(kernel_string.find("Filter3x3") != std::string::npos); - const auto invoker_ptr = instance.MakeInvokerPointer(); EXPECT_NE(invoker_ptr, nullptr); } From 9fc3d64e16332b226374e2e5de31339cc3d29e1c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Mon, 3 Nov 2025 14:33:56 +0000 Subject: [PATCH 02/20] Rename test config file. --- .../builder/test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 2 +- experimental/builder/test/conv/test_ckb_conv_fwd_1d_i8.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_2d_fp16.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_2d_fp32.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_3d_bf16.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 2 +- .../builder/test/conv/test_ckb_conv_fwd_3d_fp32.cpp | 2 +- ...{ckb_conv_test_common.hpp => ckb_conv_test_configs.hpp} | 7 ++++++- 10 files changed, 15 insertions(+), 10 deletions(-) rename experimental/builder/test/utils/{ckb_conv_test_common.hpp => ckb_conv_test_configs.hpp} (98%) 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 ab83f4d109..36c9a7de64 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 8e6abdf6f5..d6060d3b88 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 0bfadfd21f..411659d863 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 a4c3e06858..05f219a49b 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 7acc0851cd..10c293c258 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 e4aa8775ee..9ae507d3e7 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 75dc313f34..d100e6f1d6 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" 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 c0cfccbd03..399c85d2cc 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" namespace ck_tile::builder::testing { 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 09006948f0..63948559a7 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,4 +1,4 @@ -#include "utils/ckb_conv_test_common.hpp" +#include "utils/ckb_conv_test_configs.hpp" namespace { diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp similarity index 98% rename from experimental/builder/test/utils/ckb_conv_test_common.hpp rename to experimental/builder/test/utils/ckb_conv_test_configs.hpp index 35d2a3ca6e..cd34d6c288 100644 --- a/experimental/builder/test/utils/ckb_conv_test_common.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp @@ -135,7 +135,7 @@ constexpr BlockGemm BlockGemmDesc_v5_intrawave = {.pipeline_version = BlockGemmP // Common test implementation template -constexpr void run_test() +constexpr void run_test(std::vector> checks = {}) { auto instance = typename Builder::Instance{}; @@ -145,6 +145,11 @@ constexpr void run_test() const auto invoker_ptr = instance.MakeInvokerPointer(); EXPECT_NE(invoker_ptr, nullptr); + + for (const auto& check : checks) + { + check(kernel_string); + } } } // namespace ck_tile::builder::test_utils From 0d68298d43984941aaee55d296537937dc34c1b0 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Mon, 3 Nov 2025 14:39:54 +0000 Subject: [PATCH 03/20] Separate test configs from other testing assets. --- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 1 + .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 1 + .../test/conv/test_ckb_conv_fwd_1d_i8.cpp | 1 + .../test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 1 + .../test/conv/test_ckb_conv_fwd_2d_fp16.cpp | 1 + .../test/conv/test_ckb_conv_fwd_2d_fp32.cpp | 1 + .../test/conv/test_ckb_conv_fwd_3d_bf16.cpp | 1 + .../test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 1 + .../test/conv/test_ckb_conv_fwd_3d_fp32.cpp | 1 + .../test/utils/ckb_conv_test_common.hpp | 35 +++++++++++++++++++ .../test/utils/ckb_conv_test_configs.hpp | 20 ----------- 11 files changed, 44 insertions(+), 20 deletions(-) create mode 100644 experimental/builder/test/utils/ckb_conv_test_common.hpp 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 36c9a7de64..b7e7edd463 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 d6060d3b88..d442d9f273 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 411659d863..b97e9fbd64 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 05f219a49b..5094978219 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 10c293c258..1cf861cd02 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 9ae507d3e7..86e8b8e453 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 d100e6f1d6..bef4cbcea1 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" 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 399c85d2cc..860fdefd8f 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" namespace ck_tile::builder::testing { 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 63948559a7..6713509b2d 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,4 +1,5 @@ #include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_common.hpp" namespace { diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_common.hpp new file mode 100644 index 0000000000..29eced2344 --- /dev/null +++ b/experimental/builder/test/utils/ckb_conv_test_common.hpp @@ -0,0 +1,35 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#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(std::vector> checks = {}) +{ + 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& check : checks) + { + check(kernel_string); + } +} + +} // namespace ck_tile::builder::test_utils diff --git a/experimental/builder/test/utils/ckb_conv_test_configs.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp index cd34d6c288..0db8af2dc0 100644 --- a/experimental/builder/test/utils/ckb_conv_test_configs.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp @@ -3,7 +3,6 @@ #pragma once -#include #include "impl/conv_algorithm_types.hpp" #include "impl/conv_signature_types.hpp" #include "ck_tile/builder/conv_builder.hpp" @@ -133,23 +132,4 @@ constexpr BlockGemm BlockGemmDesc_v4_intrawave = {.pipeline_version = BlockGemmP constexpr BlockGemm BlockGemmDesc_v5_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V5, .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; -// Common test implementation -template -constexpr void run_test(std::vector> checks = {}) -{ - 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& check : checks) - { - check(kernel_string); - } -} - } // namespace ck_tile::builder::test_utils From 15f010138b46e94b5bbab530039b4435415df981 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Mon, 3 Nov 2025 16:06:14 +0000 Subject: [PATCH 04/20] Fix namespace. --- experimental/builder/test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 860fdefd8f..63bef01864 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,7 +1,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_common.hpp" -namespace ck_tile::builder::testing { +namespace { using namespace ck_tile::builder::test_utils; From 1893079fdc14129e62d1892c4d26db1011b2884b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Mon, 3 Nov 2025 17:06:01 +0000 Subject: [PATCH 05/20] WIP: Builder for expected test results. --- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 8 +++- .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 7 +++- .../test/utils/ckb_conv_test_common.hpp | 39 ++++++++++++++++--- 3 files changed, 47 insertions(+), 7 deletions(-) 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 b7e7edd463..0cc1e83810 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 @@ -28,7 +28,13 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v2_intrawave}; using Builder = ConvBuilder; - run_test(); + + const auto& asserts = InstanceNameAsserts{} + .StartsWith("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3") + .Contains("Filter1x1Stride1Pad0") + .Contains("BlkGemmPipelineVersion: v2"); + + run_test(asserts); } } // 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 d442d9f273..382e40e5a8 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 @@ -29,7 +29,12 @@ TEST(FwdConvInstances, .loop_scheduler = LoopScheduler::DEFAULT}; using Builder = ConvBuilder; - run_test(); + + const auto& asserts = InstanceNameAsserts{} + .StartsWith("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle") + .Contains("Default"); + + run_test(asserts); } } // namespace diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_common.hpp index 29eced2344..7fc9b6aeee 100644 --- a/experimental/builder/test/utils/ckb_conv_test_common.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_common.hpp @@ -4,6 +4,7 @@ #pragma once #include +#include #include "impl/conv_algorithm_types.hpp" #include "impl/conv_signature_types.hpp" #include "ck_tile/builder/conv_builder.hpp" @@ -13,9 +14,40 @@ 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(std::vector> checks = {}) +constexpr void run_test(const InstanceNameAsserts& asserts) { auto instance = typename Builder::Instance{}; @@ -26,10 +58,7 @@ constexpr void run_test(std::vector> che const auto invoker_ptr = instance.MakeInvokerPointer(); EXPECT_NE(invoker_ptr, nullptr); - for (const auto& check : checks) - { - check(kernel_string); - } + asserts.Check(kernel_string); } } // namespace ck_tile::builder::test_utils 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 06/20] 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 From adf0a80290d4671b3bb3c067a3591fb4a22d880b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Tue, 4 Nov 2025 08:59:57 +0000 Subject: [PATCH 07/20] clang-format --- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 13 ++++++------- .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 6 ++---- .../test/conv/test_ckb_conv_fwd_1d_i8.cpp | 6 ++---- .../test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 17 ++++++++--------- .../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 ++++---- 9 files changed, 38 insertions(+), 44 deletions(-) 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 8792882314..ca358d0d61 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 @@ -2,7 +2,7 @@ #include "utils/ckb_conv_test_utils.hpp" namespace { - + using namespace ck_tile::builder::test_utils; // 1D BF16 (channels-first) with Pipeline V2 and FILTER_1X1_STRIDE1_PAD0 specialization and SCALE @@ -28,12 +28,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v2_intrawave}; using Builder = ConvBuilder; - run_test({ - "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 256, 256, 32", - "Filter1x1Stride1Pad0", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: v2"}); + 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 651e2b78a2..5c6d563f46 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 @@ -29,10 +29,8 @@ TEST(FwdConvInstances, .loop_scheduler = LoopScheduler::DEFAULT}; using Builder = ConvBuilder; - run_test({ - "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle", - "64, 64, 32, 32", - "Default"}); + 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 978532dfd4..f5a5123f5c 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 @@ -28,10 +28,8 @@ TEST(FwdConvInstances, .loop_scheduler = LoopScheduler::DEFAULT}; using Builder = ConvBuilder; - run_test({ - "DeviceGroupedConvFwdMultipleD_Wmma_CShuffle", - "128, 64, 64, 64", - "Default"}); + 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 e04d5a72c2..c79620db20 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,7 +1,7 @@ #include "utils/ckb_conv_test_configs.hpp" #include "utils/ckb_conv_test_utils.hpp" -namespace{ +namespace { using namespace ck_tile::builder::test_utils; @@ -27,12 +27,11 @@ TEST(FwdConvInstances, .block_gemm = BlockGemmDesc_v1_intrawave}; using Builder = ConvBuilder; - run_test({ - "DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 256, 256, 32", - "Default", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: v1"}); + 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 @@ -58,8 +57,8 @@ TEST(FwdConvInstances, using Builder = ConvBuilder; run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "Filter3x3", - "BlkGemmPipelineVersion: v5"}); + "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 d5953e7315..dc059e2914 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 @@ -27,10 +27,10 @@ TEST(FwdConvInstances, using Builder = ConvBuilder; run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 256, 256, 32", - "Filter1x1Pad0", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: 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 f1ea331bc5..e8474881fc 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 @@ -27,10 +27,10 @@ TEST(FwdConvInstances, using Builder = ConvBuilder; run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 128, 128, 32", - "Filter1x1Stride1Pad0", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: v4"}); + "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 86e557fd92..be3bb31717 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 @@ -28,10 +28,10 @@ TEST(FwdConvInstances, using Builder = ConvBuilder; run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 256, 256, 32", - "Default", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: 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 8590b7e93f..b7d611ef5f 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 @@ -28,10 +28,10 @@ TEST(FwdConvInstances, using Builder = ConvBuilder; run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 128, 128, 32", - "Filter1x1Pad0", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: v4"}); + "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 030e000144..48b8f8c215 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 @@ -28,10 +28,10 @@ TEST(FwdConvInstances, using Builder = ConvBuilder; run_test({"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3", - "256, 256, 256, 32", - "Filter1x1Pad0", - "BlkGemmPipelineScheduler: Intrawave", - "BlkGemmPipelineVersion: v1"}); + "256, 256, 256, 32", + "Filter1x1Pad0", + "BlkGemmPipelineScheduler: Intrawave", + "BlkGemmPipelineVersion: v1"}); } } // namespace From 69a93a57f094596a8d203112276ca73356e573cd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Tue, 4 Nov 2025 10:57:50 +0000 Subject: [PATCH 08/20] Change if-else statements into switch in conv factory. --- .../include/ck_tile/builder/conv_factory.hpp | 286 +++++++----------- 1 file changed, 112 insertions(+), 174 deletions(-) diff --git a/experimental/builder/include/ck_tile/builder/conv_factory.hpp b/experimental/builder/include/ck_tile/builder/conv_factory.hpp index 8ea3e18d65..35bc0cf5eb 100644 --- a/experimental/builder/include/ck_tile/builder/conv_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_factory.hpp @@ -290,49 +290,44 @@ struct BlockGemmSpec }; template -constexpr BlockGemmSpec SetBlockGemm() +consteval BlockGemmSpec SetBlockGemm() { constexpr auto& BG = ALGORITHM.block_gemm; ck::BlockGemmPipelineScheduler scheduler; ck::BlockGemmPipelineVersion version; - if constexpr(BG.scheduler == BlockGemmPipelineScheduler::INTRAWAVE) + switch(BG.scheduler) { - scheduler = ck::BlockGemmPipelineScheduler::Intrawave; - } - else if constexpr(BG.scheduler == BlockGemmPipelineScheduler::INTERWAVE) - { - scheduler = ck::BlockGemmPipelineScheduler::Interwave; - } - else - { - static_assert(false, "Unknown BlockGemmPipelineScheduler"); + case BlockGemmPipelineScheduler::INTRAWAVE: + scheduler = ck::BlockGemmPipelineScheduler::Intrawave; + break; + case BlockGemmPipelineScheduler::INTERWAVE: + scheduler = ck::BlockGemmPipelineScheduler::Interwave; + break; + default: + throw "Unknown BlockGemmPipelineScheduler"; } - if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V1) + switch(BG.pipeline_version) { - version = ck::BlockGemmPipelineVersion::v1; - } - else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V2) - { - version = ck::BlockGemmPipelineVersion::v2; - } - else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V3) - { - version = ck::BlockGemmPipelineVersion::v3; - } - else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V4) - { - version = ck::BlockGemmPipelineVersion::v4; - } - else if constexpr(BG.pipeline_version == BlockGemmPipelineVersion::V5) - { - version = ck::BlockGemmPipelineVersion::v5; - } - else - { - static_assert(false, "Unknown BlockGemmPipelineVersion"); + case BlockGemmPipelineVersion::V1: + version = ck::BlockGemmPipelineVersion::v1; + break; + case BlockGemmPipelineVersion::V2: + version = ck::BlockGemmPipelineVersion::v2; + break; + case BlockGemmPipelineVersion::V3: + version = ck::BlockGemmPipelineVersion::v3; + break; + case BlockGemmPipelineVersion::V4: + version = ck::BlockGemmPipelineVersion::v4; + break; + case BlockGemmPipelineVersion::V5: + version = ck::BlockGemmPipelineVersion::v5; + break; + default: + throw "Unknown BlockGemmPipelineVersion"; } return BlockGemmSpec{.pipeline_version = version, .scheduler = scheduler}; @@ -441,18 +436,15 @@ template consteval ck::LoopScheduler SetLoopScheduler() { constexpr auto loop_scheduler = ALGORITHM.loop_scheduler; - - if constexpr(loop_scheduler == LoopScheduler::DEFAULT) + using ck_loop_sched = ck::LoopScheduler; + switch(loop_scheduler) { - return ck::LoopScheduler::Default; - } - else if constexpr(loop_scheduler == LoopScheduler::INTERWAVE) - { - return ck::LoopScheduler::Interwave; - } - else - { - static_assert(false, "Unknown LoopScheduler"); + case LoopScheduler::DEFAULT: + return ck_loop_sched::Default; + case LoopScheduler::INTERWAVE: + return ck_loop_sched::Interwave; + default: + throw "Unknown LoopScheduler"; } } @@ -460,29 +452,21 @@ template consteval ck::PipelineVersion SetGridwiseGemmPipelineVersion() { constexpr auto pipeline_version = ALGORITHM.gridwise_gemm.pipeline_version; - if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V1) + using ck_pipeline = ck::PipelineVersion; + switch(pipeline_version) { - return ck::PipelineVersion::v1; - } - else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V2) - { - return ck::PipelineVersion::v2; - } - else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V3) - { - static_assert(false, "V3 is used only for stream-K."); - } - else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::V4) - { - return ck::PipelineVersion::v4; - } - else if constexpr(pipeline_version == GridwiseGemmPipelineVersion::WEIGHT_ONLY) - { - return ck::PipelineVersion::weight_only; - } - else - { - static_assert(false, "Unknown GridwiseGemmPipelineVersion"); + case GridwiseGemmPipelineVersion::V1: + return ck_pipeline::v1; + case GridwiseGemmPipelineVersion::V2: + return ck_pipeline::v2; + case GridwiseGemmPipelineVersion::V4: + return ck_pipeline::v4; + case GridwiseGemmPipelineVersion::WEIGHT_ONLY: + return ck_pipeline::weight_only; + case GridwiseGemmPipelineVersion::V3: + throw "GridwiseGemmPipelineVersion::V3 is used only for stream-K."; + default: + throw "Unknown GridwiseGemmPipelineVersion"; } } @@ -490,74 +474,44 @@ template consteval ck::tensor_operation::device::GemmSpecialization SetGemmSpecialization() { constexpr auto gemm_spec = ALGORITHM.gemm_specialization; + using ck_gemm_spec = ck::tensor_operation::device::GemmSpecialization; - if constexpr(gemm_spec == GemmSpecialization::Default) + switch(gemm_spec) { - return ck::tensor_operation::device::GemmSpecialization::Default; - } - else if constexpr(gemm_spec == GemmSpecialization::MPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::NPadding) - { - return ck::tensor_operation::device::GemmSpecialization::NPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::KPadding) - { - return ck::tensor_operation::device::GemmSpecialization::KPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MNPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MNPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MKPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MKPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::NKPadding) - { - return ck::tensor_operation::device::GemmSpecialization::NKPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MNKPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MNKPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::OPadding) - { - return ck::tensor_operation::device::GemmSpecialization::OPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MOPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::NOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::NOPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::KOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::KOPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MNOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MNOPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MKOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MKOPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::NKOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::NKOPadding; - } - else if constexpr(gemm_spec == GemmSpecialization::MNKOPadding) - { - return ck::tensor_operation::device::GemmSpecialization::MNKOPadding; - } - else - { - static_assert(false, "Unknown GemmSpecialization"); + case GemmSpecialization::Default: + return ck_gemm_spec::Default; + case GemmSpecialization::MPadding: + return ck_gemm_spec::MPadding; + case GemmSpecialization::NPadding: + return ck_gemm_spec::NPadding; + case GemmSpecialization::KPadding: + return ck_gemm_spec::KPadding; + case GemmSpecialization::MNPadding: + return ck_gemm_spec::MNPadding; + case GemmSpecialization::MKPadding: + return ck_gemm_spec::MKPadding; + case GemmSpecialization::NKPadding: + return ck_gemm_spec::NKPadding; + case GemmSpecialization::MNKPadding: + return ck_gemm_spec::MNKPadding; + case GemmSpecialization::OPadding: + return ck_gemm_spec::OPadding; + case GemmSpecialization::MOPadding: + return ck_gemm_spec::MOPadding; + case GemmSpecialization::NOPadding: + return ck_gemm_spec::NOPadding; + case GemmSpecialization::KOPadding: + return ck_gemm_spec::KOPadding; + case GemmSpecialization::MNOPadding: + return ck_gemm_spec::MNOPadding; + case GemmSpecialization::MKOPadding: + return ck_gemm_spec::MKOPadding; + case GemmSpecialization::NKOPadding: + return ck_gemm_spec::NKOPadding; + case GemmSpecialization::MNKOPadding: + return ck_gemm_spec::MNKOPadding; + default: + throw "Unknown GemmSpecialization"; } } @@ -565,30 +519,21 @@ template consteval ck::BlockGemmPipelineVersion SetBlockGemmPipelineVersion() { constexpr auto version = ALGORITHM.pipeline_version; - - if constexpr(version == BlockGemmPipelineVersion::V1) + using ck_block_gemm = ck::BlockGemmPipelineVersion; + switch(version) { - return ck::BlockGemmPipelineVersion::v1; - } - else if constexpr(version == BlockGemmPipelineVersion::V2) - { - return ck::BlockGemmPipelineVersion::v2; - } - else if constexpr(version == BlockGemmPipelineVersion::V3) - { - return ck::BlockGemmPipelineVersion::v3; - } - else if constexpr(version == BlockGemmPipelineVersion::V4) - { - return ck::BlockGemmPipelineVersion::v4; - } - else if constexpr(version == BlockGemmPipelineVersion::V5) - { - return ck::BlockGemmPipelineVersion::v5; - } - else - { - static_assert(false, "Unknown BlockGemmPipelineVersion"); + case BlockGemmPipelineVersion::V1: + return ck_block_gemm::v1; + case BlockGemmPipelineVersion::V2: + return ck_block_gemm::v2; + case BlockGemmPipelineVersion::V3: + return ck_block_gemm::v3; + case BlockGemmPipelineVersion::V4: + return ck_block_gemm::v4; + case BlockGemmPipelineVersion::V5: + return ck_block_gemm::v5; + default: + throw "Unknown BlockGemmPipelineVersion"; } } @@ -596,26 +541,19 @@ template consteval ck::tensor_operation::device::ConvolutionForwardSpecialization SetFwdConvSpecialization() { constexpr auto specialization = ALGORITHM.fwd_specialization; - - if constexpr(specialization == ConvFwdSpecialization::DEFAULT) + using ck_conv_spec = ck::tensor_operation::device::ConvolutionForwardSpecialization; + switch(specialization) { - return ck::tensor_operation::device::ConvolutionForwardSpecialization::Default; - } - else if constexpr(specialization == ConvFwdSpecialization::FILTER_1X1_PAD0) - { - return ck::tensor_operation::device::ConvolutionForwardSpecialization::Filter1x1Pad0; - } - else if constexpr(specialization == ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0) - { - return ck::tensor_operation::device::ConvolutionForwardSpecialization::Filter1x1Stride1Pad0; - } - else if constexpr(specialization == ConvFwdSpecialization::FILTER_3x3) - { - return ck::tensor_operation::device::ConvolutionForwardSpecialization::Filter3x3; - } - else - { - static_assert(false, "Unknown ConvFwdSpecialization"); + case ConvFwdSpecialization::DEFAULT: + return ck_conv_spec::Default; + case ConvFwdSpecialization::FILTER_1X1_PAD0: + return ck_conv_spec::Filter1x1Pad0; + case ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0: + return ck_conv_spec::Filter1x1Stride1Pad0; + case ConvFwdSpecialization::FILTER_3x3: + return ck_conv_spec::Filter3x3; + default: + throw "Unknown ConvFwdSpecialization"; } } From c1db7497af8655818edeb90ddcc8b7b86e27df3c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Tue, 4 Nov 2025 12:04:09 +0000 Subject: [PATCH 09/20] Fix clang-formatting. --- .../include/ck_tile/builder/conv_factory.hpp | 160 +++++--------- .../test/utils/ckb_conv_test_configs.hpp | 201 +++++++++--------- .../test/utils/ckb_conv_test_utils.hpp | 3 +- 3 files changed, 158 insertions(+), 206 deletions(-) diff --git a/experimental/builder/include/ck_tile/builder/conv_factory.hpp b/experimental/builder/include/ck_tile/builder/conv_factory.hpp index 35bc0cf5eb..336e668099 100644 --- a/experimental/builder/include/ck_tile/builder/conv_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_factory.hpp @@ -299,35 +299,23 @@ consteval BlockGemmSpec SetBlockGemm() switch(BG.scheduler) { - case BlockGemmPipelineScheduler::INTRAWAVE: - scheduler = ck::BlockGemmPipelineScheduler::Intrawave; - break; - case BlockGemmPipelineScheduler::INTERWAVE: - scheduler = ck::BlockGemmPipelineScheduler::Interwave; - break; - default: - throw "Unknown BlockGemmPipelineScheduler"; + case BlockGemmPipelineScheduler::INTRAWAVE: + scheduler = ck::BlockGemmPipelineScheduler::Intrawave; + break; + case BlockGemmPipelineScheduler::INTERWAVE: + scheduler = ck::BlockGemmPipelineScheduler::Interwave; + break; + default: throw "Unknown BlockGemmPipelineScheduler"; } switch(BG.pipeline_version) { - case BlockGemmPipelineVersion::V1: - version = ck::BlockGemmPipelineVersion::v1; - break; - case BlockGemmPipelineVersion::V2: - version = ck::BlockGemmPipelineVersion::v2; - break; - case BlockGemmPipelineVersion::V3: - version = ck::BlockGemmPipelineVersion::v3; - break; - case BlockGemmPipelineVersion::V4: - version = ck::BlockGemmPipelineVersion::v4; - break; - case BlockGemmPipelineVersion::V5: - version = ck::BlockGemmPipelineVersion::v5; - break; - default: - throw "Unknown BlockGemmPipelineVersion"; + case BlockGemmPipelineVersion::V1: version = ck::BlockGemmPipelineVersion::v1; break; + case BlockGemmPipelineVersion::V2: version = ck::BlockGemmPipelineVersion::v2; break; + case BlockGemmPipelineVersion::V3: version = ck::BlockGemmPipelineVersion::v3; break; + case BlockGemmPipelineVersion::V4: version = ck::BlockGemmPipelineVersion::v4; break; + case BlockGemmPipelineVersion::V5: version = ck::BlockGemmPipelineVersion::v5; break; + default: throw "Unknown BlockGemmPipelineVersion"; } return BlockGemmSpec{.pipeline_version = version, .scheduler = scheduler}; @@ -436,15 +424,12 @@ template consteval ck::LoopScheduler SetLoopScheduler() { constexpr auto loop_scheduler = ALGORITHM.loop_scheduler; - using ck_loop_sched = ck::LoopScheduler; + using ck_loop_sched = ck::LoopScheduler; switch(loop_scheduler) { - case LoopScheduler::DEFAULT: - return ck_loop_sched::Default; - case LoopScheduler::INTERWAVE: - return ck_loop_sched::Interwave; - default: - throw "Unknown LoopScheduler"; + case LoopScheduler::DEFAULT: return ck_loop_sched::Default; + case LoopScheduler::INTERWAVE: return ck_loop_sched::Interwave; + default: throw "Unknown LoopScheduler"; } } @@ -452,21 +437,16 @@ template consteval ck::PipelineVersion SetGridwiseGemmPipelineVersion() { constexpr auto pipeline_version = ALGORITHM.gridwise_gemm.pipeline_version; - using ck_pipeline = ck::PipelineVersion; + using ck_pipeline = ck::PipelineVersion; switch(pipeline_version) { - case GridwiseGemmPipelineVersion::V1: - return ck_pipeline::v1; - case GridwiseGemmPipelineVersion::V2: - return ck_pipeline::v2; - case GridwiseGemmPipelineVersion::V4: - return ck_pipeline::v4; - case GridwiseGemmPipelineVersion::WEIGHT_ONLY: - return ck_pipeline::weight_only; - case GridwiseGemmPipelineVersion::V3: - throw "GridwiseGemmPipelineVersion::V3 is used only for stream-K."; - default: - throw "Unknown GridwiseGemmPipelineVersion"; + case GridwiseGemmPipelineVersion::V1: return ck_pipeline::v1; + case GridwiseGemmPipelineVersion::V2: return ck_pipeline::v2; + case GridwiseGemmPipelineVersion::V4: return ck_pipeline::v4; + case GridwiseGemmPipelineVersion::WEIGHT_ONLY: return ck_pipeline::weight_only; + case GridwiseGemmPipelineVersion::V3: + throw "GridwiseGemmPipelineVersion::V3 is used only for stream-K."; + default: throw "Unknown GridwiseGemmPipelineVersion"; } } @@ -474,44 +454,27 @@ template consteval ck::tensor_operation::device::GemmSpecialization SetGemmSpecialization() { constexpr auto gemm_spec = ALGORITHM.gemm_specialization; - using ck_gemm_spec = ck::tensor_operation::device::GemmSpecialization; + using ck_gemm_spec = ck::tensor_operation::device::GemmSpecialization; switch(gemm_spec) { - case GemmSpecialization::Default: - return ck_gemm_spec::Default; - case GemmSpecialization::MPadding: - return ck_gemm_spec::MPadding; - case GemmSpecialization::NPadding: - return ck_gemm_spec::NPadding; - case GemmSpecialization::KPadding: - return ck_gemm_spec::KPadding; - case GemmSpecialization::MNPadding: - return ck_gemm_spec::MNPadding; - case GemmSpecialization::MKPadding: - return ck_gemm_spec::MKPadding; - case GemmSpecialization::NKPadding: - return ck_gemm_spec::NKPadding; - case GemmSpecialization::MNKPadding: - return ck_gemm_spec::MNKPadding; - case GemmSpecialization::OPadding: - return ck_gemm_spec::OPadding; - case GemmSpecialization::MOPadding: - return ck_gemm_spec::MOPadding; - case GemmSpecialization::NOPadding: - return ck_gemm_spec::NOPadding; - case GemmSpecialization::KOPadding: - return ck_gemm_spec::KOPadding; - case GemmSpecialization::MNOPadding: - return ck_gemm_spec::MNOPadding; - case GemmSpecialization::MKOPadding: - return ck_gemm_spec::MKOPadding; - case GemmSpecialization::NKOPadding: - return ck_gemm_spec::NKOPadding; - case GemmSpecialization::MNKOPadding: - return ck_gemm_spec::MNKOPadding; - default: - throw "Unknown GemmSpecialization"; + case GemmSpecialization::Default: return ck_gemm_spec::Default; + case GemmSpecialization::MPadding: return ck_gemm_spec::MPadding; + case GemmSpecialization::NPadding: return ck_gemm_spec::NPadding; + case GemmSpecialization::KPadding: return ck_gemm_spec::KPadding; + case GemmSpecialization::MNPadding: return ck_gemm_spec::MNPadding; + case GemmSpecialization::MKPadding: return ck_gemm_spec::MKPadding; + case GemmSpecialization::NKPadding: return ck_gemm_spec::NKPadding; + case GemmSpecialization::MNKPadding: return ck_gemm_spec::MNKPadding; + case GemmSpecialization::OPadding: return ck_gemm_spec::OPadding; + case GemmSpecialization::MOPadding: return ck_gemm_spec::MOPadding; + case GemmSpecialization::NOPadding: return ck_gemm_spec::NOPadding; + case GemmSpecialization::KOPadding: return ck_gemm_spec::KOPadding; + case GemmSpecialization::MNOPadding: return ck_gemm_spec::MNOPadding; + case GemmSpecialization::MKOPadding: return ck_gemm_spec::MKOPadding; + case GemmSpecialization::NKOPadding: return ck_gemm_spec::NKOPadding; + case GemmSpecialization::MNKOPadding: return ck_gemm_spec::MNKOPadding; + default: throw "Unknown GemmSpecialization"; } } @@ -519,21 +482,15 @@ template consteval ck::BlockGemmPipelineVersion SetBlockGemmPipelineVersion() { constexpr auto version = ALGORITHM.pipeline_version; - using ck_block_gemm = ck::BlockGemmPipelineVersion; + using ck_block_gemm = ck::BlockGemmPipelineVersion; switch(version) { - case BlockGemmPipelineVersion::V1: - return ck_block_gemm::v1; - case BlockGemmPipelineVersion::V2: - return ck_block_gemm::v2; - case BlockGemmPipelineVersion::V3: - return ck_block_gemm::v3; - case BlockGemmPipelineVersion::V4: - return ck_block_gemm::v4; - case BlockGemmPipelineVersion::V5: - return ck_block_gemm::v5; - default: - throw "Unknown BlockGemmPipelineVersion"; + case BlockGemmPipelineVersion::V1: return ck_block_gemm::v1; + case BlockGemmPipelineVersion::V2: return ck_block_gemm::v2; + case BlockGemmPipelineVersion::V3: return ck_block_gemm::v3; + case BlockGemmPipelineVersion::V4: return ck_block_gemm::v4; + case BlockGemmPipelineVersion::V5: return ck_block_gemm::v5; + default: throw "Unknown BlockGemmPipelineVersion"; } } @@ -541,19 +498,14 @@ template consteval ck::tensor_operation::device::ConvolutionForwardSpecialization SetFwdConvSpecialization() { constexpr auto specialization = ALGORITHM.fwd_specialization; - using ck_conv_spec = ck::tensor_operation::device::ConvolutionForwardSpecialization; + using ck_conv_spec = ck::tensor_operation::device::ConvolutionForwardSpecialization; switch(specialization) { - case ConvFwdSpecialization::DEFAULT: - return ck_conv_spec::Default; - case ConvFwdSpecialization::FILTER_1X1_PAD0: - return ck_conv_spec::Filter1x1Pad0; - case ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0: - return ck_conv_spec::Filter1x1Stride1Pad0; - case ConvFwdSpecialization::FILTER_3x3: - return ck_conv_spec::Filter3x3; - default: - throw "Unknown ConvFwdSpecialization"; + case ConvFwdSpecialization::DEFAULT: return ck_conv_spec::Default; + case ConvFwdSpecialization::FILTER_1X1_PAD0: return ck_conv_spec::Filter1x1Pad0; + case ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0: return ck_conv_spec::Filter1x1Stride1Pad0; + case ConvFwdSpecialization::FILTER_3x3: return ck_conv_spec::Filter3x3; + default: throw "Unknown ConvFwdSpecialization"; } } diff --git a/experimental/builder/test/utils/ckb_conv_test_configs.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp index 0db8af2dc0..02c3dfec9b 100644 --- a/experimental/builder/test/utils/ckb_conv_test_configs.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp @@ -12,124 +12,125 @@ namespace ck_tile::builder::test_utils { using namespace ck_tile::builder; using namespace test; -constexpr BlockTransferABC FwdBlockTransfer_4x64_1 {.block_transfer_a = {.k0 = 4, .m_n = 64, .k1 = 1}, - .block_transfer_b = {.k0 = 4, .m_n = 64, .k1 = 1}, - .thread_cluster_dims_c = {.m_block = 1, - .m_wave_per_xdl = 32, - .n_block = 1, - .n_wave_per_xdl = 8}, - .lds_transfer_a = {.src_vector_dim = 2, - .src_scalar_per_vector = 2, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = false}, - .lds_transfer_b = {.src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = false}, - .epilogue_c = {.m_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, - .block_transfer_access_order_a = {1, 0, 2}, - .block_transfer_access_order_b = {1, 0, 2}, - .src_access_order_a = {1, 0, 2}, - .src_access_order_b = {1, 0, 2}}; +constexpr BlockTransferABC FwdBlockTransfer_4x64_1{ + .block_transfer_a = {.k0 = 4, .m_n = 64, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 64, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 32, + .n_block = 1, + .n_wave_per_xdl = 8}, + .lds_transfer_a = {.src_vector_dim = 2, + .src_scalar_per_vector = 2, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = false}, + .lds_transfer_b = {.src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = false}, + .epilogue_c = {.m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, + .block_transfer_access_order_a = {1, 0, 2}, + .block_transfer_access_order_b = {1, 0, 2}, + .src_access_order_a = {1, 0, 2}, + .src_access_order_b = {1, 0, 2}}; -constexpr BlockTransferABC FwdBlockTransfer_4x16x1 {.block_transfer_a = {.k0 = 4, .m_n = 16, .k1 = 1}, - .block_transfer_b = {.k0 = 4, .m_n = 16, .k1 = 1}, - .thread_cluster_dims_c = {.m_block = 1, - .m_wave_per_xdl = 16, - .n_block = 1, - .n_wave_per_xdl = 4}, - .lds_transfer_a = {.src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, - .lds_transfer_b = {.src_vector_dim = 2, - .src_scalar_per_vector = 8, - .lds_dst_scalar_per_vector = 8, - .is_direct_load = false, - .lds_padding = true}, - .epilogue_c = {.m_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, - .block_transfer_access_order_a = {1, 0, 2}, - .block_transfer_access_order_b = {1, 0, 2}, - .src_access_order_a = {1, 0, 2}, - .src_access_order_b = {1, 0, 2}}; +constexpr BlockTransferABC FwdBlockTransfer_4x16x1{ + .block_transfer_a = {.k0 = 4, .m_n = 16, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 16, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 16, + .n_block = 1, + .n_wave_per_xdl = 4}, + .lds_transfer_a = {.src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, + .lds_transfer_b = {.src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true}, + .epilogue_c = {.m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, + .block_transfer_access_order_a = {1, 0, 2}, + .block_transfer_access_order_b = {1, 0, 2}, + .src_access_order_a = {1, 0, 2}, + .src_access_order_b = {1, 0, 2}}; -constexpr BlockTransferABC FwdBlockTransfer_4x32x1 {.block_transfer_a = {.k0 = 4, .m_n = 32, .k1 = 1}, - .block_transfer_b = {.k0 = 4, .m_n = 32, .k1 = 1}, - .thread_cluster_dims_c = {.m_block = 1, - .m_wave_per_xdl = 32, - .n_block = 1, - .n_wave_per_xdl = 4}, - .lds_transfer_a = {.src_vector_dim = 2, - .src_scalar_per_vector = 16, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, - .lds_transfer_b = {.src_vector_dim = 2, - .src_scalar_per_vector = 16, - .lds_dst_scalar_per_vector = 16, - .is_direct_load = false, - .lds_padding = true}, - .epilogue_c = {.m_per_wave_per_shuffle = 1, - .n_per_wave_per_shuffle = 1, - .scalar_per_vector = 8}, - .block_transfer_access_order_a = {1, 0, 2}, - .block_transfer_access_order_b = {1, 0, 2}, - .src_access_order_a = {1, 0, 2}, - .src_access_order_b = {1, 0, 2}}; +constexpr BlockTransferABC FwdBlockTransfer_4x32x1{ + .block_transfer_a = {.k0 = 4, .m_n = 32, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 32, .k1 = 1}, + .thread_cluster_dims_c = {.m_block = 1, + .m_wave_per_xdl = 32, + .n_block = 1, + .n_wave_per_xdl = 4}, + .lds_transfer_a = {.src_vector_dim = 2, + .src_scalar_per_vector = 16, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, + .lds_transfer_b = {.src_vector_dim = 2, + .src_scalar_per_vector = 16, + .lds_dst_scalar_per_vector = 16, + .is_direct_load = false, + .lds_padding = true}, + .epilogue_c = {.m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8}, + .block_transfer_access_order_a = {1, 0, 2}, + .block_transfer_access_order_b = {1, 0, 2}, + .src_access_order_a = {1, 0, 2}, + .src_access_order_b = {1, 0, 2}}; -constexpr GridwiseXdlGemm FwdGemmParams_Xdl_4x4_per_wave {.ak1 = 8, - .bk1 = 8, - .m_per_xdl = 32, - .n_per_xdl = 32, - .m_xdl_per_wave = 4, - .n_xdl_per_wave = 4}; +constexpr GridwiseXdlGemm FwdGemmParams_Xdl_4x4_per_wave{ + .ak1 = 8, .bk1 = 8, .m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4}; -constexpr GridwiseXdlGemm FwdGemmParams_Xdl_2x1_per_wave {.ak1 = 8, - .bk1 = 8, - .m_per_xdl = 32, - .n_per_xdl = 32, - .m_xdl_per_wave = 2, - .n_xdl_per_wave = 1}; +constexpr GridwiseXdlGemm FwdGemmParams_Xdl_2x1_per_wave{ + .ak1 = 8, .bk1 = 8, .m_per_xdl = 32, .n_per_xdl = 32, .m_xdl_per_wave = 2, .n_xdl_per_wave = 1}; -constexpr GridwiseWmmaGemm FwdGemmParams_Wmma_2x1_per_wave {.k1 = 8, - .m_per_wmma = 32, - .n_per_wmma = 32, - .m_wmma_per_wave = 2, - .n_wmma_per_wave = 1, - .pipeline_version = GridwiseGemmPipelineVersion::V1}; +constexpr GridwiseWmmaGemm FwdGemmParams_Wmma_2x1_per_wave{.k1 = 8, + .m_per_wmma = 32, + .n_per_wmma = 32, + .m_wmma_per_wave = 2, + .n_wmma_per_wave = 1, + .pipeline_version = + GridwiseGemmPipelineVersion::V1}; -constexpr ThreadBlock FwdThreadBlock_256x256x32 {.block_size = 256, - .tile_size = {.m = 256, .n = 256, .k = 32}}; +constexpr ThreadBlock FwdThreadBlock_256x256x32{.block_size = 256, + .tile_size = {.m = 256, .n = 256, .k = 32}}; -constexpr ThreadBlock FwdThreadBlock_128x128x32 {.block_size = 256, - .tile_size = {.m = 128, .n = 128, .k = 32}}; +constexpr ThreadBlock FwdThreadBlock_128x128x32{.block_size = 256, + .tile_size = {.m = 128, .n = 128, .k = 32}}; -constexpr ThreadBlock FwdThreadBlock_64x32x32 {.block_size = 64, - .tile_size = {.m = 64, .n = 32, .k = 32}}; +constexpr ThreadBlock FwdThreadBlock_64x32x32{.block_size = 64, + .tile_size = {.m = 64, .n = 32, .k = 32}}; -constexpr ThreadBlock FwdThreadBlock_64x64x64 {.block_size = 128, - .tile_size = {.m = 64, .n = 64, .k = 64}}; +constexpr ThreadBlock FwdThreadBlock_64x64x64{.block_size = 128, + .tile_size = {.m = 64, .n = 64, .k = 64}}; constexpr BlockGemm BlockGemmDesc_v1_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V1, - .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + .scheduler = + BlockGemmPipelineScheduler::INTRAWAVE}; constexpr BlockGemm BlockGemmDesc_v2_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V2, - .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + .scheduler = + BlockGemmPipelineScheduler::INTRAWAVE}; constexpr BlockGemm BlockGemmDesc_v3_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V3, - .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + .scheduler = + BlockGemmPipelineScheduler::INTRAWAVE}; constexpr BlockGemm BlockGemmDesc_v4_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V4, - .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + .scheduler = + BlockGemmPipelineScheduler::INTRAWAVE}; constexpr BlockGemm BlockGemmDesc_v5_intrawave = {.pipeline_version = BlockGemmPipelineVersion::V5, - .scheduler = BlockGemmPipelineScheduler::INTRAWAVE}; + .scheduler = + BlockGemmPipelineScheduler::INTRAWAVE}; } // 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 index c04117e8d6..c2c6ae63cf 100644 --- a/experimental/builder/test/utils/ckb_conv_test_utils.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_utils.hpp @@ -14,7 +14,6 @@ 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) @@ -28,7 +27,7 @@ constexpr void run_test(const std::vector& kernel_instance_componen const auto invoker_ptr = instance.MakeInvokerPointer(); EXPECT_NE(invoker_ptr, nullptr); - for (const auto& component : kernel_instance_components) + for(const auto& component : kernel_instance_components) { EXPECT_THAT(kernel_string, ::testing::HasSubstr(component)); } From 5d6427a0fd69514b4a321bde2005c635d47399ac Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Tue, 4 Nov 2025 12:16:09 +0000 Subject: [PATCH 10/20] Removed unnecessary includes. --- experimental/builder/test/utils/ckb_conv_test_utils.hpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/experimental/builder/test/utils/ckb_conv_test_utils.hpp b/experimental/builder/test/utils/ckb_conv_test_utils.hpp index c2c6ae63cf..09f9523c3a 100644 --- a/experimental/builder/test/utils/ckb_conv_test_utils.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_utils.hpp @@ -5,13 +5,8 @@ #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 From 8681ced9629f6e952afa5b77c5f3549d60920efa Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Bart=C5=82omiej=20Kocot?= Date: Tue, 4 Nov 2025 15:04:53 +0100 Subject: [PATCH 11/20] [CK TILE] Refactor Conv configs and Conv Elementwise (#3151) * [CK TILE] Refactor Conv configs and Conv Elementwise * fix --- .../20_grouped_convolution/conv_configs.hpp | 8 +- .../grouped_convolution_backward_data.cpp | 6 +- ...uped_convolution_backward_data_invoker.hpp | 53 ++-- ...ed_convolution_backward_weight_invoker.hpp | 9 +- ...tion_backward_weight_two_stage_invoker.hpp | 9 +- .../grouped_convolution_forward.cpp | 6 +- ...grouped_convolution_forward_bias_clamp.cpp | 6 +- .../grouped_convolution_forward_invoker.hpp | 57 ++-- ...ouped_convolution_forward_large_tensor.cpp | 6 +- ...nvolution_forward_large_tensor_invoker.hpp | 248 ++++++++++-------- ...n_grouped_convolution_bwd_data_example.inc | 16 +- .../run_grouped_convolution_fwd_example.inc | 16 +- .../grouped_convolution_forward_kernel.hpp | 7 +- .../utils/grouped_convolution_utils.hpp | 2 - 14 files changed, 230 insertions(+), 219 deletions(-) diff --git a/example/ck_tile/20_grouped_convolution/conv_configs.hpp b/example/ck_tile/20_grouped_convolution/conv_configs.hpp index c688215280..b9edf247cc 100644 --- a/example/ck_tile/20_grouped_convolution/conv_configs.hpp +++ b/example/ck_tile/20_grouped_convolution/conv_configs.hpp @@ -18,11 +18,7 @@ struct ConvConfigBase static constexpr bool kPadN = true; static constexpr bool kPadK = true; - static constexpr bool PermuteA = false; - static constexpr bool PermuteB = false; - - static constexpr bool TransposeC = false; - static constexpr bool UseStructuredSparsity = false; + static constexpr bool TransposeC = false; static constexpr ck_tile::index_t VectorSizeA = 4; static constexpr ck_tile::index_t VectorSizeB = 8; @@ -34,8 +30,6 @@ struct ConvConfigBase static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave; static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V3; static constexpr ck_tile::index_t NumWaveGroups = 1; - static constexpr bool Preshuffle = false; - static constexpr bool TiledMMAPermuteN = false; static constexpr ck_tile::index_t NumGroupsToMerge = 1; }; diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp index ad593b1418..8ea892a215 100644 --- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp +++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data.cpp @@ -14,7 +14,7 @@ #include "grouped_convolution_backward_data_invoker.hpp" #include "run_grouped_convolution_bwd_data_example.inc" -template