From 1dcea1825f162561365699abedce22b77b4bf7b3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Mon, 5 Jan 2026 03:09:22 -0500 Subject: [PATCH] Fix DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3 factory and compute types for input and output tensor in bwd weigth convs. --- .../ck_tile/builder/conv_algorithm_concepts.hpp | 2 +- .../builder/conv_algorithm_diagnostics.hpp | 2 +- .../ck_tile/builder/factory/conv_algorithms.hpp | 12 ++++++------ .../conv_bwd_weight_multi_d_wmma_v3_factory.hpp | 16 +++++++--------- .../conv_bwd_weight_multi_d_xdl_factory.hpp | 4 ++-- ...conv_bwd_weight_two_stage_wmma_v3_factory.hpp | 2 +- .../conv_bwd_weight_two_stage_xdl_factory.hpp | 2 +- .../factory/conv_bwd_weight_wmma_v3_factory.hpp | 2 +- .../factory/conv_bwd_weight_xdl_factory.hpp | 2 +- .../factory/conv_bwd_weight_xdl_v3_factory.hpp | 4 ++-- experimental/builder/test/CMakeLists.txt | 2 +- ...conv_bwd_weight_multi_d_wmma_cshuffle_v3.cpp} | 16 ++++++++-------- .../builder/test/impl/conv_algorithm_types.hpp | 2 +- .../test/utils/conv_algorithm_type_utils.hpp | 4 ++-- 14 files changed, 35 insertions(+), 37 deletions(-) rename experimental/builder/test/conv/ck/{test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle.cpp => test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle_v3.cpp} (85%) diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp index 326350eb27..511c0c2b2d 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -255,7 +255,7 @@ concept SpecifiesBlockGemm = requires { }; template -concept SpecifiedGridwiseGemmPipeline = requires +concept SpecifiesGridwiseGemmPipeline = requires { { T::pipeline_version } -> std::convertible_to; }; diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp index df81b20f8e..047eb7e4bd 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_diagnostics.hpp @@ -1193,7 +1193,7 @@ consteval auto detailed_diagnostic_SpecifiesLdsTransfer() -> std::string { } template -consteval auto detailed_diagnostic_SpecifiedGridwiseGemmPipeline() -> std::string { +consteval auto detailed_diagnostic_SpecifiesGridwiseGemmPipeline() -> std::string { if constexpr (requires { T::pipeline_version; }) { using PipelineType = decltype(T::pipeline_version); constexpr bool convertible = std::convertible_to; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp index 800403b16d..891ea8f730 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -153,7 +153,7 @@ struct FwdWmmaAlgorithm { CHECK_CONCEPT(T, SpecifiesGemmSpecialization) CHECK_CONCEPT(T, SpecifiesNumPrefetchStages) CHECK_CONCEPT(T, SpecifiesLoopScheduler) - CHECK_CONCEPT(T, SpecifiedGridwiseGemmPipeline) + CHECK_CONCEPT(T, SpecifiesGridwiseGemmPipeline) static constexpr bool c1 = c_ConvAlgorithmDescriptor; static constexpr bool c2 = c_SpecifiesThreadBlock; @@ -166,7 +166,7 @@ struct FwdWmmaAlgorithm { static constexpr bool c9 = c_SpecifiesGemmSpecialization; static constexpr bool c10 = c_SpecifiesNumPrefetchStages; static constexpr bool c11 = c_SpecifiesLoopScheduler; - static constexpr bool c12 = c_SpecifiedGridwiseGemmPipeline; + static constexpr bool c12 = c_SpecifiesGridwiseGemmPipeline; static consteval bool is_valid() { return c1 && c2 && c3 && c4 && c5 && c6 && c7 && c8 && c9 && c10 && c11 && c12; @@ -186,7 +186,7 @@ struct FwdWmmaAlgorithm { DIAGNOSTIC_LINE(SpecifiesGemmSpecialization) + DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + DIAGNOSTIC_LINE(SpecifiesLoopScheduler) + - DIAGNOSTIC_LINE(SpecifiedGridwiseGemmPipeline); + DIAGNOSTIC_LINE(SpecifiesGridwiseGemmPipeline); } }; @@ -479,12 +479,12 @@ template struct BwdWmmaAlgorithm : public BwdWmmaAlgorithmBase { CHECK_CONCEPT(T, SpecifiesNumPrefetchStages) CHECK_CONCEPT(T, SpecifiesLoopScheduler) - CHECK_CONCEPT(T, SpecifiedGridwiseGemmPipeline) + CHECK_CONCEPT(T, SpecifiesGridwiseGemmPipeline) CHECK_CONCEPT(T, SpecifiesGenericInstance) static constexpr bool c9 = c_SpecifiesNumPrefetchStages; static constexpr bool c10 = c_SpecifiesLoopScheduler; - static constexpr bool c11 = c_SpecifiedGridwiseGemmPipeline; + static constexpr bool c11 = c_SpecifiesGridwiseGemmPipeline; static constexpr bool c12 = c_SpecifiesGenericInstance; static consteval bool is_valid() { @@ -497,7 +497,7 @@ struct BwdWmmaAlgorithm : public BwdWmmaAlgorithmBase { BwdWmmaAlgorithmBase::message() + DIAGNOSTIC_LINE(SpecifiesNumPrefetchStages) + DIAGNOSTIC_LINE(SpecifiesLoopScheduler) + - DIAGNOSTIC_LINE(SpecifiedGridwiseGemmPipeline) + + DIAGNOSTIC_LINE(SpecifiesGridwiseGemmPipeline) + DIAGNOSTIC_LINE(SpecifiesGenericInstance); } }; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp index 2ba1300537..9d94404a88 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp @@ -35,10 +35,6 @@ struct ConvBwdWeightMultiDWmmaV3Factory static constexpr auto BLOCK = internal::SetThreadBlockInfo(); static constexpr auto GRIDWISE_GEMM = ALGORITHM.gridwise_gemm; - static constexpr auto GRIDWISE_GEMM_PIPELINE_VERSION = - internal::SetGridwiseGemmPipelineVersion(); - static constexpr auto LOOP_SCHEDULER = internal::SetLoopScheduler(); - static constexpr auto A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = @@ -51,10 +47,10 @@ struct ConvBwdWeightMultiDWmmaV3Factory static_assert(InputVectorTransferLimits, "Invalid A block transfer config"); static_assert(InputVectorTransferLimits, "Invalid B block transfer config"); static_assert(OutputVectorTransferLimits, "Invalid C block transfer config"); - static_assert(AccessOrderLimits4D, "Invalid A thread cluster access order"); - static_assert(AccessOrderLimits4D, "Invalid B thread cluster access order"); - static_assert(AccessOrderLimits4D, "Invalid A source access order"); - static_assert(AccessOrderLimits4D, "Invalid B source access order"); + static_assert(AccessOrderLimits3D, "Invalid A thread cluster access order"); + static_assert(AccessOrderLimits3D, "Invalid B thread cluster access order"); + static_assert(AccessOrderLimits3D, "Invalid A source access order"); + static_assert(AccessOrderLimits3D, "Invalid B source access order"); // The forward convolution kernel class instance. using Instance = ck::tensor_operation::device::DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3< @@ -100,7 +96,9 @@ struct ConvBwdWeightMultiDWmmaV3Factory to_sequence_v, C_BLOCK_TRANSFER.scalar_per_vector, BLOCK_GEMM.scheduler, - BLOCK_GEMM.pipeline_version>; + BLOCK_GEMM.pipeline_version, + typename Types::OutComputeType, + typename Types::InComputeType>; }; } // namespace ck_tile::builder::factory diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp index 2c11f3b436..dae333d99e 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp @@ -95,8 +95,8 @@ struct ConvBwdWeightMultiDXdlFactory C_BLOCK_TRANSFER.n_xdl_per_wave_per_shuffle, to_sequence_v, C_BLOCK_TRANSFER.scalar_per_vector, - typename Types::InComputeType, - typename Types::WeiComputeType>; + typename Types::OutComputeType, + typename Types::InComputeType>; }; } // namespace ck_tile::builder::factory diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp index 635b822293..8288874f0e 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp @@ -96,8 +96,8 @@ struct ConvBwdWeightTwoStageWmmaV3Factory BLOCK_GEMM.scheduler, BLOCK_GEMM.pipeline_version, ALGORITHM.num_conv_groups_to_merge, + typename Types::OutComputeType, typename Types::InComputeType, - typename Types::WeiComputeType, ALGORITHM.max_transpose_transfer_src_scalar_per_vector, ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; }; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp index fae8ad7b87..df3e4c01b2 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp @@ -97,8 +97,8 @@ struct ConvBwdWeightTwoStageXdlFactory BLOCK_GEMM.scheduler, BLOCK_GEMM.pipeline_version, ALGORITHM.num_conv_groups_to_merge, + typename Types::OutComputeType, typename Types::InComputeType, - typename Types::WeiComputeType, ALGORITHM.max_transpose_transfer_src_scalar_per_vector, ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; }; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp index ab75f5b072..d1cfbe9e8d 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp @@ -95,8 +95,8 @@ struct ConvBwdWeightWmmaV3Factory C_BLOCK_TRANSFER.scalar_per_vector, BLOCK_GEMM.scheduler, BLOCK_GEMM.pipeline_version, + typename Types::OutComputeType, typename Types::InComputeType, - typename Types::WeiComputeType, ALGORITHM.max_transpose_transfer_src_scalar_per_vector, ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; }; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp index 8ffbb495ec..e89d227f82 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp @@ -93,8 +93,8 @@ struct ConvBwdWeightXdlFactory C_BLOCK_TRANSFER.n_xdl_per_wave_per_shuffle, to_sequence_v, C_BLOCK_TRANSFER.scalar_per_vector, + typename Types::OutComputeType, typename Types::InComputeType, - typename Types::WeiComputeType, ALGORITHM.max_transpose_transfer_src_scalar_per_vector, ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; }; diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp index 20aac55f31..e93b1456cd 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp @@ -96,8 +96,8 @@ struct ConvBwdWeightXdlV3Factory C_BLOCK_TRANSFER.scalar_per_vector, BLOCK_GEMM.scheduler, BLOCK_GEMM.pipeline_version, - typename Types::InComputeType, - typename Types::WeiComputeType>; + typename Types::OutComputeType, + typename Types::InComputeType>; }; } // namespace ck_tile::builder::factory diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index b2f5970d2e..667490151f 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -178,7 +178,7 @@ if (CK_USE_WMMA) conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle_v3.cpp conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp conv/ck/test_ckb_conv_bwd_weight_two_stage_wmma_cshuffle_v3.cpp - conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle.cpp + conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle_v3.cpp ) endif() diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle_v3.cpp similarity index 85% rename from experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle.cpp rename to experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle_v3.cpp index e2bcd4a926..404d1dbacd 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_multi_d_wmma_cshuffle_v3.cpp @@ -11,15 +11,15 @@ namespace ckt = ck_tile::builder::test; namespace cku = ck_tile::builder::test_utils; constexpr auto SIGNATURE = - ckt::ConvSignature{.spatial_dim = 2, + ckt::ConvSignature{.spatial_dim = 3, .direction = ckb::ConvDirection::BACKWARD_WEIGHT, .data_type = ckb::DataType::FP16, .accumulation_data_type = ckb::DataType::FP32, - .input = {.config = {.layout = ckb::TensorLayout::GNHWC}}, - .weight = {.config = {.layout = ckb::TensorLayout::GKYXC}}, - .output = {.config = {.layout = ckb::TensorLayout::GNHWK}}}; + .input = {.config = {.layout = ckb::TensorLayout::GNDHWC}}, + .weight = {.config = {.layout = ckb::TensorLayout::GKZYXC}}, + .output = {.config = {.layout = ckb::TensorLayout::GNDHWK}}}; -constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle{} +constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3{} .with_thread_block(cku::ThreadBlock_64_32x32x32) .with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave) .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3) @@ -29,14 +29,14 @@ constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeightMultiple using Builder = ckb::ConvBuilder; using Instance = Builder::Instance; -TEST(BwdWeight_2DFp16_MultiD_Wmma_Shuffle_GNHWC, Create) +TEST(BwdWeight_3DFp16_MultiD_Wmma_ShuffleV3_GNHWC, Create) { const auto expected_transfer_parameters = to_string(ALGORITHM); std::cout << "Expected Transfer Parameters: " << expected_transfer_parameters << std::endl; - cku::run_test({"DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle", + cku::run_test({"DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3", expected_transfer_parameters, "Default", - "GNHWC,GKYXC,GNHWK", + "GNDHWC,GKZYXC,GNDHWK", "PassThrough,PassThrough,PassThrough", "fp16,fp16>"}); // check compute types } diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp index eb03ecfab2..4cce9475cb 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -609,7 +609,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 = using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, GridGemm_, Prefetch_>; -using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle = +using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, BlockGemm_, MultipleDSpecialization_>; diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp index ee4bd1f597..6ad6179780 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -444,8 +444,8 @@ inline std::string to_string -inline std::string to_string( - ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle t) +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 t) { std::ostringstream oss; oss << to_string(static_cast(t)) << "," << to_string(static_cast(t))