From eefadf7afda419b1c2f644a60fd3a77e487f4fc2 Mon Sep 17 00:00:00 2001 From: Kevin Abraham Date: Tue, 10 Feb 2026 21:59:08 +0000 Subject: [PATCH] first working implementation for bwd_data_multiple_d_xdl_v1 --- .../conv_bwd_data_multi_d_xdl_factory.hpp | 8 +- .../builder/factory/conv_dispatcher.hpp | 28 +- ..._conv_bwd_data_multiple_d_xdl_cshuffle.hpp | 343 ++++++++++++++++++ .../builder/reflect/instance_traits_util.hpp | 12 + experimental/builder/test/CMakeLists.txt | 2 +- ...ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp | 8 +- .../test/impl/conv_algorithm_types.hpp | 2 +- .../test/utils/conv_algorithm_type_utils.hpp | 2 +- ...nv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp | 23 ++ 9 files changed, 403 insertions(+), 25 deletions(-) create mode 100644 experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle.hpp diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_xdl_factory.hpp index 1b21b92341..ea19bc673f 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_xdl_factory.hpp @@ -49,10 +49,10 @@ struct ConvBwdDataMultiDXdlFactory static_assert(InputVectorTransferLimits); static_assert(InputVectorTransferLimits); static_assert(OutputVectorTransferLimits); - static_assert(AccessOrderLimits4D); - static_assert(AccessOrderLimits4D); - static_assert(AccessOrderLimits4D); - static_assert(AccessOrderLimits4D); + static_assert(AccessOrderLimits3D); + static_assert(AccessOrderLimits3D); + static_assert(AccessOrderLimits3D); + static_assert(AccessOrderLimits3D); // The backward convolution kernel class instance. using Instance = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp index de289f6d91..d82bebde01 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp @@ -146,25 +146,25 @@ constexpr auto make_conv_instance() false, "No suitable forward convolution kernel factory found for the provided ALGORITHM. " "The ALGORITHM must satisfy requirements for one of: Reference, Tile, XDL V3, XDL, " - "WMMA, DL (NHWC layout), or Large Tensor variant."); + "WMMA, DL (NHWC layout)"); } } // Backward data direction (will expand with more algorithms in the future) else if constexpr(ConvDirectionIsBackwardData) { - if constexpr(BwdMultiDXdlAlgorithm) - { - return typename ConvBwdDataMultiDXdlFactory::Instance{}; - } - else - { - static_assert( - false, - "No suitable backward data convolution kernel factory found for the provided " - "ALGORITHM. " - "The ALGORITHM must satisfy requirements for one of: Reference, Tile, XDL V3, XDL, " - "WMMA, DL (NHWC layout), or Large Tensor variant."); - } + // if constexpr(BwdMultiDXdlAlgorithm) + // { + return typename ConvBwdDataMultiDXdlFactory::Instance{}; + // } + // else + // { + // static_assert( + // false, + // "No suitable backward data convolution kernel factory found for the provided " + // "ALGORITHM. " + // "The ALGORITHM must satisfy requirements for one of: Reference, Tile, XDL V3, + // XDL, " "WMMA, DL (NHWC layout), or Large Tensor variant."); + // } } // Backward weight direction (will expand with more algorithms in the future) else if constexpr(ConvDirectionIsBackwardWeight) diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle.hpp new file mode 100644 index 0000000000..8774667ae1 --- /dev/null +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle.hpp @@ -0,0 +1,343 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "instance_traits.hpp" +#include "instance_traits_util.hpp" +#include "ck/tensor_operation/gpu/device/convolution_backward_data_specialization.hpp" + +namespace ck::tensor_operation::device { + +template +struct DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1; + +} // namespace ck::tensor_operation::device + +namespace ck_tile { +namespace reflect { +/// @brief Tag type for DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle device kernel +struct DeviceGroupedConvBwdData_multiple_d_Xdl_CShuffle_Tag +{ +}; + +template +struct InstanceTraits< + ck::tensor_operation::device::DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1< + NDimSpatial, + OutLayout_, + WeiLayout_, + DsLayout_, + InLayout_, + OutDataType_, + WeiDataType_, + AccDataType_, + OutComputeType_, + DsDataType_, + InDataType_, + OutElementwiseOperation_, + WeiElementwiseOperation_, + InElementwiseOperation_, + ConvBackwardDataSpecialization, + do_pad_gemm_m, + do_pad_gemm_n, + num_gemm_k_prefetch_stages, + BlockSize, + MPerBlock, + NPerBlock, + K0PerBlock, + AK1, + BK1, + MPerXDL, + NPerXDL, + MXdlPerWave, + NXdlPerWave, + ABlockTransferThreadClusterLengths_K0_M_K1_, + ABlockTransferThreadClusterArrangeOrder_, + ABlockTransferSrcAccessOrder_, + ABlockTransferSrcVectorDim, + ABlockTransferSrcScalarPerVector, + ABlockTransferDstScalarPerVector_AK1, + ABlockLdsAddExtraM, + BBlockTransferThreadClusterLengths_K0_N_K1_, + BBlockTransferThreadClusterArrangeOrder_, + BBlockTransferSrcAccessOrder_, + BBlockTransferSrcVectorDim, + BBlockTransferSrcScalarPerVector, + BBlockTransferDstScalarPerVector_BK1, + BBlockLdsAddExtraN, + CShuffleMXdlPerWavePerShuffle, + CShuffleNXdlPerWavePerShuffle, + CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_, + CBlockTransferScalarPerVector_NWaveNPerXdl, + LoopSched, + ComputeTypeA_, + ComputeTypeB_, + max_transpose_transfer_src_scalar_per_vector, + max_transpose_transfer_dst_scalar_per_vector>> +{ + static constexpr auto kTensorOpName = "DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle"; + + /// @brief Tag type identifying this device kernel variant + using device_kernel_tag = DeviceGroupedConvBwdData_multiple_d_Xdl_CShuffle_Tag; + + static constexpr ck::index_t kSpatialDim = NDimSpatial; + + using InLayout = InLayout_; + using WeiLayout = WeiLayout_; + using OutLayout = OutLayout_; + using DsLayout = DsLayout_; + + using InDataType = InDataType_; + using WeiDataType = WeiDataType_; + using OutDataType = OutDataType_; + using AccDataType = AccDataType_; + using DsDataType = DsDataType_; + + using InElementwiseOperation = InElementwiseOperation_; + using WeiElementwiseOperation = WeiElementwiseOperation_; + using OutElementwiseOperation = OutElementwiseOperation_; + + static constexpr auto kConvBwdDataSpecialization = ConvBackwardDataSpecialization; + + static constexpr ck::index_t kBlockSize = BlockSize; + static constexpr ck::index_t kMPerBlock = MPerBlock; + static constexpr ck::index_t kNPerBlock = NPerBlock; + static constexpr ck::index_t kK0PerBlock = K0PerBlock; + static constexpr ck::index_t kAK1 = AK1; + static constexpr ck::index_t kBK1 = BK1; + static constexpr ck::index_t kMPerXDL = MPerXDL; + static constexpr ck::index_t kNPerXDL = NPerXDL; + static constexpr ck::index_t kMXdlPerWave = MXdlPerWave; + static constexpr ck::index_t kNXdlPerWave = NXdlPerWave; + static constexpr ck::index_t kCShuffleMXdlPerWavePerShuffle = CShuffleMXdlPerWavePerShuffle; + static constexpr ck::index_t kCShuffleNXdlPerWavePerShuffle = CShuffleNXdlPerWavePerShuffle; + static constexpr ck::index_t kCBlockTransferScalarPerVector_NWaveNPerXdl = + CBlockTransferScalarPerVector_NWaveNPerXdl; + static constexpr ck::index_t kMaxTransposeTransferSrcScalarPerVector = + max_transpose_transfer_src_scalar_per_vector; + static constexpr ck::index_t kMaxTransposeTransferDstScalarPerVector = + max_transpose_transfer_dst_scalar_per_vector; + static constexpr bool kDoPadGemmM = do_pad_gemm_n; + static constexpr bool kDoPadGemmN = do_pad_gemm_m; + static constexpr int kNumGemmKPrefetchStage = num_gemm_k_prefetch_stages; + + using ABlockTransferThreadClusterLengths_K0_M_K1 = ABlockTransferThreadClusterLengths_K0_M_K1_; + using ABlockTransferThreadClusterArrangeOrder = ABlockTransferThreadClusterArrangeOrder_; + using ABlockTransferSrcAccessOrder = ABlockTransferSrcAccessOrder_; + // A block transfer thread cluster dimensions (converted to std::array) + static constexpr auto kAThreadClusterLengths = + detail::SequenceToArray::value; + static constexpr auto kAThreadClusterArrangeOrder = + detail::SequenceToArray::value; + static constexpr auto kABlockTransferSrcAccessOrder = + detail::SequenceToArray::value; + static constexpr ck::index_t kABlockTransferSrcVectorDim = ABlockTransferSrcVectorDim; + static constexpr ck::index_t kABlockTransferSrcScalarPerVector = + ABlockTransferSrcScalarPerVector; + static constexpr ck::index_t kABlockTransferDstScalarPerVectorK1 = + ABlockTransferDstScalarPerVector_AK1; + static constexpr bool kABlockLdsExtraM = ABlockLdsAddExtraM; + + using BBlockTransferThreadClusterLengths_K0_N_K1 = BBlockTransferThreadClusterLengths_K0_N_K1_; + using BBlockTransferThreadClusterArrangeOrder = BBlockTransferThreadClusterArrangeOrder_; + using BBlockTransferSrcAccessOrder = BBlockTransferSrcAccessOrder_; + // B block transfer thread cluster dimensions (converted to std::array) + static constexpr auto kBThreadClusterLengths = + detail::SequenceToArray::value; + static constexpr auto kBThreadClusterArrangeOrder = + detail::SequenceToArray::value; + static constexpr auto kBBlockTransferSrcAccessOrder = + detail::SequenceToArray::value; + static constexpr ck::index_t kBBlockTransferSrcVectorDim = BBlockTransferSrcVectorDim; + static constexpr ck::index_t kBBlockTransferSrcScalarPerVector = + BBlockTransferSrcScalarPerVector; + static constexpr ck::index_t kBBlockTransferDstScalarPerVectorK1 = + BBlockTransferDstScalarPerVector_BK1; + static constexpr bool kBBlockLdsExtraN = BBlockLdsAddExtraN; + + using CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock = + CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_; + + using ComputeTypeA = ComputeTypeA_; + using ComputeTypeB = ComputeTypeB_; + + static constexpr ck::LoopScheduler kLoopScheduler = LoopSched; + + static constexpr auto kCThreadClusterLengths = detail::SequenceToArray< + CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock>::value; + + // Static member function to generate instance string + static std::string instance_string() + { + std::ostringstream oss; + + // Kernel type name + oss << "DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle"; + + // Template parameters in exact order + oss << "<" << kSpatialDim; // 1. NDimSpatial + oss << "," << detail::layout_name(); // 2. InLayout + oss << "," << detail::layout_name(); // 3. WeiLayout + oss << "," << detail::layout_name(); // 4. OutLayout + oss << "," << detail::tuple_name(); // 5. DsLayout + oss << "," << detail::type_name(); // 6. InDataType + oss << "," << detail::type_name(); // 7. WeiDataType + oss << "," << detail::type_name(); // 8. OutDataType + oss << "," << detail::type_name(); // 9. AccDataType + oss << "," << detail::tuple_name(); // 10. DsDataType + oss << "," + << detail::elementwise_op_name(); // 11. InElementwiseOperation + oss << "," + << detail::elementwise_op_name(); // 12. + // WeiElementwiseOperation + oss << "," + << detail::elementwise_op_name(); // 13. + // OutElementwiseOperation + oss << "," + << detail::conv_bwd_data_spec_name( + kConvBwdDataSpecialization); // 14. ConvBackwardDataSpecialization + oss << "," << kBlockSize; // 15. BlockSize + oss << "," << kMPerBlock; // 16. MPerBlock + oss << "," << kNPerBlock; // 17. NPerBlock + oss << "," << kK0PerBlock; // 18. K0PerBlock + oss << "," << kAK1; // 19. AK1 + oss << "," << kBK1; // 19,5. BK1 + oss << "," << kMPerXDL; // 20. MPerXDL + oss << "," << kNPerXDL; // 21. NPerXDL + oss << "," << kMXdlPerWave; // 22. MXdlPerWave + oss << "," << kNXdlPerWave; // 23. NXdlPerWave + oss << "," << detail::sequence_name(); // 24. + oss << "," << detail::sequence_name(); // 25. + oss << "," << detail::sequence_name(); // 26. + oss << "," << kABlockTransferSrcVectorDim; // 27. + oss << "," << kABlockTransferSrcScalarPerVector; // 28. + oss << "," << kABlockTransferDstScalarPerVectorK1; // 29. + oss << "," << (kABlockLdsExtraM ? "true" : "false"); // 30. + oss << "," << detail::sequence_name(); // 31. + oss << "," << detail::sequence_name(); // 32. + oss << "," << detail::sequence_name(); // 33. + oss << "," << kBBlockTransferSrcVectorDim; // 34. + oss << "," << kBBlockTransferSrcScalarPerVector; // 35. + oss << "," << kBBlockTransferDstScalarPerVectorK1; // 36. + oss << "," << (kBBlockLdsExtraN ? "true" : "false"); // 37. + oss << "," << kCShuffleMXdlPerWavePerShuffle; // 38. + oss << "," << kCShuffleNXdlPerWavePerShuffle; // 39. + oss << "," + << detail::sequence_name< + CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock>(); // 40. + oss << "," << kCBlockTransferScalarPerVector_NWaveNPerXdl; // 41. + oss << "," << detail::type_name(); // 42. + oss << "," << detail::type_name(); // 43. + oss << "," << detail::loop_scheduler_name(kLoopScheduler); // 44. LoopSched + + oss << ">"; + + return oss.str(); + } +}; + +} // namespace reflect +} // namespace ck_tile diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp index 1055cbc038..bd2cd1afd0 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp @@ -138,6 +138,18 @@ constexpr std::string_view conv_bwd_weight_spec_name( } } +// Convert ConvolutionBackwardDataSpecialization enum to string +constexpr std::string_view +conv_bwd_data_spec_name(ck::tensor_operation::device::ConvolutionBackwardDataSpecialization spec) +{ + using enum ck::tensor_operation::device::ConvolutionBackwardDataSpecialization; + switch(spec) + { + case Default: return "Default"; + case Filter1x1Stride1Pad0: return "Filter1x1Stride1Pad0"; + } +} + // Convert GemmSpecialization enum to string constexpr std::string_view gemm_spec_name(ck::tensor_operation::device::GemmSpecialization spec) { diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index 6b63d221d3..4ddfe27d93 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -178,7 +178,6 @@ set(BWD_WEIGHT_TESTS conv/ck/test_ckb_conv_bwd_weight_xdl_cshuffle_v3.cpp conv/ck/test_ckb_conv_bwd_weight_dl.cpp conv/ck_tile/test_ckb_conv_bwd_weight_2d_fp16_v3.cpp - conv/ck/test_ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp ) if (CK_USE_WMMA) @@ -195,6 +194,7 @@ target_link_libraries(test_ckb_build_bwd_weight_instances PRIVATE utility) add_ck_builder_test(test_ckb_build_bwd_data_instances conv/ck_tile/test_ckb_conv_bwd_data_2d_fp16_v3.cpp + conv/ck/test_ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp ) target_link_libraries(test_ckb_build_bwd_data_instances PRIVATE utility) diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp index 62682f828c..5559c7ec2c 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_xdl_cshuffle.cpp @@ -20,9 +20,9 @@ constexpr auto SIGNATURE = .output = {.config = {.layout = ckb::TensorLayout::GNHWK}}}; constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle{} - .with_thread_block(cku::ThreadBlock_256_128x128x8) + .with_thread_block(cku::ThreadBlock_256_256x128x32) .with_gemm_config(cku::BwdDataGemmParams_Xdl_4x4_per_wave) - .with_transfer(cku::BwdTransfer_4x64x1) + .with_transfer(cku::Transfer_4x64x1) .with_prefetch_config(1, ckb::PipelineScheduler::DEFAULT) .with_bwd_data_specialization(ckb::ConvSpecialization::DEFAULT) .with_gemm_pad_params(0, 0) @@ -35,10 +35,10 @@ TEST(BwdData_2DFp16_MultiD_CShuffle_GNHWC, Create) { const auto expected_transfer_parameters = to_string(ALGORITHM); std::cout << "Expected Transfer Parameters: " << expected_transfer_parameters << std::endl; - cku::run_test({"DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1", + cku::run_test({"DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle", expected_transfer_parameters, "Default", "GNHWC,GKYXC,GNHWK", "PassThrough,PassThrough,PassThrough", - "fp16,fp16>"}); // check compute types + "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 02fbdee164..7a00a3c7e6 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -732,7 +732,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 = using ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle = ConvAlgorithmTemplate, + Transfer_<>, ConvSpecializationBwdData_, MultipleDSpecialization_, Prefetch_, diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp index 7e40940295..8d6f526345 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -525,7 +525,7 @@ inline std::string to_string(t)) << "," << to_string(static_cast(t)) << "," - << to_string(static_cast>(t)); + << to_string(static_cast>(t)); return oss.str(); } diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp index 16b12cf386..8b71a4fa40 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp @@ -27,6 +27,11 @@ #include "ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp" +#ifdef CK_EXPERIMENTAL_BUILDER +#include "ck_tile/builder/reflect/description.hpp" +#include "ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle.hpp" +#endif + namespace ck { namespace tensor_operation { namespace device { @@ -1942,6 +1947,24 @@ struct DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 "The argument pointer is not an object of " "DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1::Argument structure!"); } + +#ifdef CK_EXPERIMENTAL_BUILDER + std::string GetInstanceString() const override + { + static_assert(ck_tile::reflect::HasInstanceTraits, + "Specialization of instance_traits not found. Please check that a " + "specialization exists in file " + "ck_tile/builder/reflect/" + "instance_traits_device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle.hpp " + "for the given template parameters."); + return ck_tile::reflect::instance_string(); + } + + std::unique_ptr describe() const override + { + return std::make_unique(GetInstanceString()); + } +#endif }; } // namespace device