diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp index 8c59dd21b1..3eca32c897 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_bf16.cpp @@ -33,7 +33,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x256x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0, + .with_fwd_specializations(ConvSpecialization::FILTER_1X1_STRIDE1_PAD0, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v2_intrawave); diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp index 7ab3ac605b..e543ce6fa0 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_64_64x32x32) .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) .with_transfer(Transfer_4x16x1) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_prefetch_config(1, 2, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp index d8fbd3827e..c87ffbd066 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp @@ -32,7 +32,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_128_64x64x64) .with_gemm_config(FwdGemmParams_Wmma_2x1_per_wave) .with_transfer(Transfer_4x32x1) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_prefetch_config(1, 0, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp index 5f3bdfe414..0ce530c1f8 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x256x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v1_intrawave); using Builder = ConvBuilder; @@ -67,7 +67,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x256x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvFwdSpecialization::FILTER_3x3, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::FILTER_3x3, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v5_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp index f6403b312c..56efb70f24 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp @@ -35,7 +35,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_64_64x32x32) .with_gemm_config(FwdGemmParams_Xdl_2x2_per_wave) .with_transfer(Transfer_4x16x1) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp index a49f55e6d6..b740ac8704 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_dl_fp16.cpp @@ -26,7 +26,7 @@ TEST(FwdConvInstances, Create_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK_Ins constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK{} .with_thread_block(ThreadBlock_256_128x128x16) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_dl_thread_config(DlThreadConfig_16x2x4x4x1) .with_dl_thread_cluster(DlThreadCluster_8x2) .with_dl_transfer(DlFwdTransfer); @@ -60,7 +60,7 @@ TEST(FwdConvInstances, constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK{} .with_thread_block(ThreadBlock_256_128x128x16) - .with_fwd_specializations(ConvFwdSpecialization::FILTER_1X1_PAD0, + .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, GemmSpecialization::MNKPadding) .with_dl_thread_config(DlThreadConfig_16x2x4x4x1) .with_dl_thread_cluster(DlThreadCluster_8x2) diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp index 7b543e0e3e..0f09ce6e8e 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp @@ -24,7 +24,7 @@ constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xd .with_thread_block(cku::ThreadBlock_256_256x256x32) .with_gemm_config(cku::FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(cku::Transfer_4x64x1) - .with_fwd_specializations(ckb::ConvFwdSpecialization::DEFAULT, + .with_fwd_specializations(ckb::ConvSpecialization::DEFAULT, ckb::GemmSpecialization::MNKPadding) .with_block_gemm(cku::BlockGemmDesc_v3_intrawave); diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp index a3f493c1d7..a8f97c417e 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp32.cpp @@ -29,7 +29,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_128x128x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0, + .with_fwd_specializations(ConvSpecialization::FILTER_1X1_STRIDE1_PAD0, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v4_intrawave); diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp index 279c942ba9..c95d48d712 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x128x32) .with_gemm_config(FwdGemmParams_Xdl_4x2_per_wave) .with_transfer(Transfer_4x64x1_fp8) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp index cd9655f0b4..4ad3283883 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x128x32) .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) .with_transfer(Transfer_4x16x1) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT)}; @@ -67,7 +67,7 @@ TEST( .with_thread_block(ThreadBlock_128_128x128x32) .with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave) .with_transfer(Transfer_4x16x1) - .with_specializations(ConvFwdSpecialization::FILTER_1X1_PAD0, + .with_specializations(ConvSpecialization::FILTER_1X1_PAD0, GemmSpecialization::MNKPadding) .with_prefetch_config(1, 1, PipelineScheduler::DEFAULT)}; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_bf16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_bf16.cpp index b29b4471c3..9e6ca00e58 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_bf16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_bf16.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x256x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvFwdSpecialization::DEFAULT, GemmSpecialization::MNKPadding) + .with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v3_intrawave); using Builder = ConvBuilder; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp index 9c4b9d4ec0..5e4e84543f 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp16.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_128x128x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd_specializations(ConvFwdSpecialization::FILTER_1X1_PAD0, + .with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v4_intrawave); diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp index ed0ec0e3c1..2c2e80b1a7 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_3d_fp32.cpp @@ -30,7 +30,7 @@ TEST(FwdConvInstances, .with_thread_block(ThreadBlock_256_256x256x32) .with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave) .with_transfer(Transfer_4x64x1) - .with_fwd-specializations(ConvFwdSpecialization::FILTER_1X1_PAD0, + .with_fwd-specializations(ConvSpecialization::FILTER_1X1_PAD0, GemmSpecialization::MNKPadding) .with_block_gemm(BlockGemmDesc_v1_intrawave); diff --git a/experimental/builder/test/conv/ck/test_conv_traits.cpp b/experimental/builder/test/conv/ck/test_conv_traits.cpp index d5661ad67b..b3a76e4e11 100644 --- a/experimental/builder/test/conv/ck/test_conv_traits.cpp +++ b/experimental/builder/test/conv/ck/test_conv_traits.cpp @@ -101,7 +101,7 @@ TEST_F(ConvTraitsTest, ConvFwdTraitsExtraction) // Verify specializations EXPECT_EQ(Traits::gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvSpecialization::DEFAULT); // Verify algorithm information EXPECT_EQ(Traits::thread_block_size, 256); @@ -229,7 +229,7 @@ TEST_F(ConvTraitsTest, ConvFwdBaseTraitsExtraction) // Verify specializations EXPECT_EQ(Traits::gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvSpecialization::DEFAULT); // Verify algorithm information EXPECT_EQ(Traits::thread_block_size, 256); @@ -313,7 +313,7 @@ TEST_F(ConvTraitsTest, ConvFwdLargeTensorTraitsExtraction) // Verify specializations EXPECT_EQ(Traits::gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(Traits::conv_specialization, ck_tile::builder::ConvSpecialization::DEFAULT); // Verify algorithm information EXPECT_EQ(Traits::thread_block_size, 256); diff --git a/experimental/builder/test/test_conv_description.cpp b/experimental/builder/test/test_conv_description.cpp index 158cb2668f..98c8ebbab1 100644 --- a/experimental/builder/test/test_conv_description.cpp +++ b/experimental/builder/test/test_conv_description.cpp @@ -148,7 +148,7 @@ struct DefaultAlgorithm }, }; - ckb::ConvFwdSpecialization fwd_specialization = ckb::ConvFwdSpecialization::DEFAULT; + ckb::ConvFwdSpecialization fwd_specialization = ckb::ConvSpecialization::DEFAULT; ckb::GemmSpecialization gemm_specialization = ckb::GemmSpecialization::Default; ckb::test::BlockGemm block_gemm{.pipeline_version = ckb::PipelineVersion::V4, .scheduler = ckb::PipelineScheduler::INTRAWAVE}; diff --git a/experimental/builder/test/unit_conv_tuning_params.cpp b/experimental/builder/test/unit_conv_tuning_params.cpp index b35a1ced55..2ea181ffb0 100644 --- a/experimental/builder/test/unit_conv_tuning_params.cpp +++ b/experimental/builder/test/unit_conv_tuning_params.cpp @@ -79,7 +79,7 @@ TEST(ConvTuningParams, AssignsFwdConvSpecialization) constexpr struct Algorithm { ckb::ConvFwdSpecialization fwd_specialization = - ckb::ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0; + ckb::ConvSpecialization::FILTER_1X1_STRIDE1_PAD0; } kAlgorithm; constexpr auto conv_spec = SetFwdConvSpecialization();