From 2ae886dbe5bb70aaf7eae052024ee3ddd4f1d802 Mon Sep 17 00:00:00 2001 From: Kevin Abraham Date: Fri, 13 Feb 2026 15:24:58 +0000 Subject: [PATCH] added device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3 to builder --- .../builder/conv_algorithm_concepts.hpp | 16 +- ..._data_multi_d_wmma_cshuffle_v3_factory.hpp | 85 +++-- .../builder/factory/conv_dispatcher.hpp | 11 +- ...v_bwd_data_multiple_d_wmma_cshuffle_v3.hpp | 354 ++++++++++++++++++ ..._conv_bwd_data_multiple_d_xdl_cshuffle.hpp | 6 +- experimental/builder/test/CMakeLists.txt | 1 + ...conv_bwd_data_multi_d_wmma_cshuffle_v3.cpp | 45 +++ .../test/impl/conv_algorithm_types.hpp | 31 +- .../test/utils/ckb_conv_test_configs.hpp | 7 + .../test/utils/conv_algorithm_type_utils.hpp | 27 ++ ...v_bwd_data_multiple_d_wmma_cshuffle_v3.hpp | 24 ++ 11 files changed, 568 insertions(+), 39 deletions(-) create mode 100644 experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp create mode 100644 experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_wmma_cshuffle_v3.cpp 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 bb87e50943..9cff75f049 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -47,11 +47,17 @@ concept BlockGemmPipelineDescriptor = requires(T t) { // Concept for parameters that describe a gridwise WMMA GEMM problem. template concept GridwiseWmmaGemmDescriptor = requires(T t) { - { t.k1 } -> SizeType; - { t.m_per_wmma } -> SizeType; - { t.n_per_wmma } -> SizeType; - { t.m_wmma_per_wave } -> SizeType; - { t.n_wmma_per_wave } -> SizeType; + ( + requires { { T::k1 } -> SizeType; } || + (requires { { T::ak1 } -> SizeType; } && + requires { { T::bk1 } -> SizeType; }) +) && +requires { + { T::m_per_wmma } -> SizeType; + { T::n_per_wmma } -> SizeType; + { T::m_wmma_per_wave } -> SizeType; + { T::n_wmma_per_wave } -> SizeType; +}; }; // Concept for vectorized data transfer for convolution input tensors. diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_wmma_cshuffle_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_wmma_cshuffle_v3_factory.hpp index fc7a8dea2c..c0d96eda4b 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_wmma_cshuffle_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_data_multi_d_wmma_cshuffle_v3_factory.hpp @@ -17,13 +17,13 @@ namespace ck_tile::builder::factory { -// Factory for DeviceGroupedConvBwdDataMultipleD_wmma_CShuffle instance +// Factory for DeviceGroupedConvBwdDataMultipleD_wmma_CShuffle_v3 instance // of a grouped bwd Data convolution kernel. template requires ConvDirectionIsBackwardData -struct ConvBwdDataMultiDWmmaFactory +struct ConvBwdDataMultiDWmmaV3Factory { static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; @@ -37,13 +37,12 @@ struct ConvBwdDataMultiDWmmaFactory static constexpr auto LOOP_SCHEDULER = internal::SetLoopScheduler(); 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 A_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto B_BLOCK_TRANSFER = internal::SetBwdConvBlockTransfer(); static constexpr auto C_BLOCK_TRANSFER = internal::SetCBlockTransfer(); + static constexpr auto BLOCK_GEMM = internal::SetBlockGemm(); // Check limits for the algorithm parameters. // TODO: Add more limits checks as needed. @@ -57,30 +56,60 @@ struct ConvBwdDataMultiDWmmaFactory // The backward convolution kernel class instance. using Instance = - ck::tensor_operation::device::DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle < SPATIAL_DIM, - typename Layouts::OutLayout, typename Layouts::WeiLayout, typename Layouts::DsLayout, - typename Layouts::InLayout, typename Types::OutDataType, typename Types::WeiDataType, - typename Types::AccDataType, typename Types::OutComputeType, typename Types::DsDataType, - typename Types::InDataType, typename Ops::OutElementwiseOp, - typename Ops::WeiElementwiseOp, typename Ops::InElementwiseOp, BWD_CONV_SPECIALIZATION, - bool DoPadGemmM, bool DoPadGemmN, BLOCK.block_size, BLOCK.per_block.m, BLOCK.per_block.n, - BLOCK.per_block.k, GRIDWISE_GEMM.ak1, GRIDWISE_GEMM.bk1, GRIDWISE_GEMM.m_per_wmma, - GRIDWISE_GEMM.n_per_wmma, GRIDWISE_GEMM.m_wmma_per_wave, GRIDWISE_GEMM.n_wmma_per_wave, - to_sequence_v, - to_sequence_v, - to_sequence_v, A_BLOCK_TRANSFER.src_vector_dim, - A_BLOCK_TRANSFER.src_scalar_per_vector, A_BLOCK_TRANSFER.lds_dst_scalar_per_vector, - A_BLOCK_TRANSFER.lds_padding, to_sequence_v, - to_sequence_v, - to_sequence_v, B_BLOCK_TRANSFER.src_vector_dim, - B_BLOCK_TRANSFER.src_scalar_per_vector, B_BLOCK_TRANSFER.lds_dst_scalar_per_vector, - B_BLOCK_TRANSFER.lds_padding, C_BLOCK_TRANSFER.m_xdl_per_wave_per_shuffle, - C_BLOCK_TRANSFER.n_xdl_per_wave_per_shuffle, - to_sequence_v, C_BLOCK_TRANSFER.scalar_per_vector, - ALGORITHM.num_gemm_k_prefetch_stages, LOOP_SCHEDULER, GRIDWISE_GEMM_PIPELINE_VERSION, - ALGORITHM.max_transpose_transfer_src_scalar_per_vector, - ALGORITHM.max_transpose_transfer_dst_scalar_per_vector >> - ; + ck::tensor_operation::device::DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3< + SPATIAL_DIM, + typename Layouts::OutLayout, + typename Layouts::WeiLayout, + typename Layouts::DsLayout, + typename Layouts::InLayout, + typename Types::OutDataType, + typename Types::WeiDataType, + typename Types::AccDataType, + typename Types::OutComputeType, + typename Types::DsDataType, + typename Types::InDataType, + typename Ops::OutElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::InElementwiseOp, + BWD_CONV_SPECIALIZATION, + ALGORITHM.DoPadGemmM, + ALGORITHM.DoPadGemmN, + BLOCK.block_size, + BLOCK.per_block.m, + BLOCK.per_block.n, + BLOCK.per_block.k, + GRIDWISE_GEMM.ak1, + GRIDWISE_GEMM.bk1, + GRIDWISE_GEMM.m_per_wmma, + GRIDWISE_GEMM.n_per_wmma, + GRIDWISE_GEMM.m_wmma_per_wave, + GRIDWISE_GEMM.n_wmma_per_wave, + to_sequence_v, + to_sequence_v, + to_sequence_v, + A_BLOCK_TRANSFER.src_vector_dim, + A_BLOCK_TRANSFER.src_scalar_per_vector, + A_BLOCK_TRANSFER.lds_dst_scalar_per_vector, + A_BLOCK_TRANSFER.lds_padding, + to_sequence_v, + to_sequence_v, + to_sequence_v, + B_BLOCK_TRANSFER.src_vector_dim, + B_BLOCK_TRANSFER.src_scalar_per_vector, + B_BLOCK_TRANSFER.lds_dst_scalar_per_vector, + B_BLOCK_TRANSFER.lds_padding, + C_BLOCK_TRANSFER.m_xdl_per_wave_per_shuffle, + C_BLOCK_TRANSFER.n_xdl_per_wave_per_shuffle, + to_sequence_v, + ck::Sequence, + BLOCK_GEMM.scheduler, + BLOCK_GEMM.pipeline_version, + typename Types::OutComputeType, + typename Types::InComputeType, + ALGORITHM.max_transpose_transfer_src_scalar_per_vector, + ALGORITHM.max_transpose_transfer_dst_scalar_per_vector>; }; } // namespace ck_tile::builder::factory 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 50ac4f1188..d06c61e266 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp @@ -79,6 +79,7 @@ #include "ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp" #include "ck_tile/builder/factory/conv_bwd_data_multi_d_xdl_factory.hpp" #include "ck_tile/builder/factory/conv_bwd_data_multi_d_wmma_factory.hpp" +#include "ck_tile/builder/factory/conv_bwd_data_multi_d_wmma_cshuffle_v3_factory.hpp" namespace ck_tile::builder::factory { @@ -157,6 +158,11 @@ constexpr auto make_conv_instance() { return typename ConvBwdDataMultiDXdlFactory::Instance{}; } + else if constexpr(BwdMultiDWmmaV3Algorithm) + { + return + typename ConvBwdDataMultiDWmmaV3Factory::Instance{}; + } else if constexpr(BwdMultiDWmmaAlgorithm) { return typename ConvBwdDataMultiDWmmaFactory::Instance{}; @@ -167,8 +173,9 @@ constexpr auto make_conv_instance() 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."); + "The ALGORITHM must satisfy requirements for one of: Reference, XDL multiple d, " + "Wmma multiple d, " + "or WMMA multiple d v3."); } } // Backward weight direction (will expand with more algorithms in the future) diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp new file mode 100644 index 0000000000..4c70af06f3 --- /dev/null +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp @@ -0,0 +1,354 @@ +// 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_Wmma_CShuffleV3; + +} // namespace ck::tensor_operation::device + +namespace ck_tile { +namespace reflect { +/// @brief Tag type for DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle_V3 device kernel +struct DeviceGroupedConvBwdData_multiple_d_Wmma_CShuffle_V3_Tag +{ +}; + +template +struct InstanceTraits< + ck::tensor_operation::device::DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3< + NDimSpatial, + OutLayout_, // output image + WeiLayout_, // weight + DsLayout_, // bias + InLayout_, // input image + OutDataType_, // output image + WeiDataType_, // weight + AccDataType_, + CShuffleDataType_, + DsDataType_, // bias + InDataType_, // input image + OutElementwiseOp_, // output image + WeiElementwiseOp_, // weight + InElementwiseOp_, // C, bias, and input image + ConvBackwardDataSpecialization, + DoPadGemmM, + DoPadGemmN, + BlockSize, + MPerBlock, + NPerBlock, + K0PerBlock, + AK1, + BK1, + MPerWMMA, + NPerWMMA, + MRepeat, + NRepeat, + ABlockTransferThreadClusterLengths_AK0_M_AK1_, + ABlockTransferThreadClusterArrangeOrder_, + ABlockTransferSrcAccessOrder_, + ABlockTransferSrcVectorDim, + ABlockTransferSrcScalarPerVector, + ABlockTransferDstScalarPerVector_AK1, + ABlockLdsExtraM, + BBlockTransferThreadClusterLengths_BK0_N_BK1_, + BBlockTransferThreadClusterArrangeOrder_, + BBlockTransferSrcAccessOrder_, + BBlockTransferSrcVectorDim, + BBlockTransferSrcScalarPerVector, + BBlockTransferDstScalarPerVector_BK1, + BBlockLdsExtraN, + CShuffleMRepeatPerShuffle, + CShuffleNRepeatPerShuffle, + CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_, + CDEShuffleBlockTransferScalarPerVector_NPerBlock_, + BlkGemmPipeSched, + BlkGemmPipelineVer, + ComputeTypeA_, + ComputeTypeB_, + max_transpose_transfer_src_scalar_per_vector, + max_transpose_transfer_dst_scalar_per_vector>> +{ + static constexpr auto kTensorOpName = "DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3"; + + /// @brief Tag type identifying this device kernel variant + using device_kernel_tag = DeviceGroupedConvBwdData_multiple_d_Wmma_CShuffle_V3_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 = InElementwiseOp_; + using WeiElementwiseOperation = WeiElementwiseOp_; + using OutElementwiseOperation = OutElementwiseOp_; + + 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 kMPerWmma = MPerWMMA; + static constexpr ck::index_t kNPerWmma = NPerWMMA; + static constexpr ck::index_t kMRepeat = MRepeat; + static constexpr ck::index_t kNRepeat = NRepeat; + static constexpr ck::index_t kCShuffleMRepeatPerShuffle = CShuffleMRepeatPerShuffle; + static constexpr ck::index_t kCShuffleNRepeatPerShuffle = CShuffleNRepeatPerShuffle; + 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 = DoPadGemmM; + static constexpr bool kDoPadGemmN = DoPadGemmN; + using CDEShuffleBlockTransferScalarPerVector_NPerBlock = + CDEShuffleBlockTransferScalarPerVector_NPerBlock_; + + static constexpr auto kCDEShuffleBlockTransferScalarPerVector_NPerBlock = + detail::SequenceToArray::value; + + static constexpr ck::BlockGemmPipelineScheduler kBlkGemmPipeSched = BlkGemmPipeSched; + static constexpr ck::BlockGemmPipelineVersion kBlkGemmPipelineVer = BlkGemmPipelineVer; + + using ABlockTransferThreadClusterLengths_AK0_M_AK1 = + ABlockTransferThreadClusterLengths_AK0_M_AK1_; + 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 = ABlockLdsExtraM; + + using BBlockTransferThreadClusterLengths_BK0_N_BK1 = + BBlockTransferThreadClusterLengths_BK0_N_BK1_; + using BBlockTransferThreadClusterArrangeOrder = BBlockTransferThreadClusterArrangeOrder_; + using BBlockTransferSrcAccessOrder = BBlockTransferSrcAccessOrder_; + // B block transfer thread cluster dimensions (converted to std::array) + // 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 = BBlockLdsExtraN; + + using CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock = + CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_; + + static constexpr auto kCThreadClusterLengths = detail::SequenceToArray< + CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock>::value; + + using ComputeTypeA = ComputeTypeA_; + using ComputeTypeB = ComputeTypeB_; + // Static member function to generate instance string + static std::string instance_string() + { + std::ostringstream oss; + + // Kernel type name + oss << "DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3"; + + // Template parameters in exact order + oss << "<" << kSpatialDim; // 1. NDimSpatial + oss << "," << detail::layout_name(); // 2. OutLayout + oss << "," << detail::layout_name(); // 3. WeiLayout + oss << "," << detail::tuple_name(); // 4. DsLayout + oss << "," << detail::layout_name(); // 5. InLayout + oss << "," << detail::type_name(); // 6. OutDataType + oss << "," << detail::type_name(); // 7. WeiDataType + oss << "," << detail::type_name(); // 8. AccDataType + oss << "," << detail::tuple_name(); // 9. DsDataType + oss << "," << detail::type_name(); // 10. InDataType + oss << "," + << detail::elementwise_op_name(); // 11. + // OutElementwiseOperation + oss << "," + << detail::elementwise_op_name(); // 12. + // WeiElementwiseOperation + oss << "," + << detail::elementwise_op_name(); // 13. InElementwiseOperation + oss << "," + << detail::conv_bwd_data_spec_name( + kConvBwdDataSpecialization); // 14. ConvBackwardDataSpecialization + oss << "," << kDoPadGemmM; + oss << "," << kDoPadGemmN; + oss << "," << kBlockSize; // 15. BlockSize + oss << "," << kMPerBlock; // 16. MPerBlock + oss << "," << kNPerBlock; // 17. NPerBlock + oss << "," << kK0PerBlock; // 18. K0PerBlock + oss << "," << kAK1; // 19. ABK1 + oss << "," << kBK1; // 19. ABK1 + oss << "," << kMPerWmma; // 20. MPerWmma + oss << "," << kNPerWmma; // 21. NPerWmma + oss << "," << kMRepeat; // 22. MRepeat + oss << "," << kNRepeat; // 23. NRepeat + 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 << "," << kCShuffleMRepeatPerShuffle; // 38. + oss << "," << kCShuffleNRepeatPerShuffle; // 39. + oss << "," + << detail::sequence_name< + CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock>(); // 40. + oss << "," << kCDEShuffleBlockTransferScalarPerVector_NPerBlock[0]; // 41. + oss << "," << detail::pipeline_scheduler_name(kBlkGemmPipeSched); // 43. + oss << "," << detail::pipeline_version_name(kBlkGemmPipelineVer); // 44. + oss << "," << detail::type_name(); // 45. + oss << "," << detail::type_name(); // 46. + oss << "," << kMaxTransposeTransferSrcScalarPerVector; // 47. + oss << "," << kMaxTransposeTransferDstScalarPerVector; // 48. + + oss << ">"; + + return oss.str(); + } +}; + +} // namespace reflect +} // namespace ck_tile 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 index 8d9c17cb21..08c6ebd1b9 100644 --- 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 @@ -220,8 +220,8 @@ struct InstanceTraits< 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 bool kDoPadGemmM = do_pad_gemm_m; + static constexpr bool kDoPadGemmN = do_pad_gemm_n; static constexpr int kNumGemmKPrefetchStage = num_gemm_k_prefetch_stages; using ABlockTransferThreadClusterLengths_K0_M_K1 = ABlockTransferThreadClusterLengths_K0_M_K1_; @@ -307,7 +307,7 @@ struct InstanceTraits< oss << "," << kNPerBlock; // 17. NPerBlock oss << "," << kK0PerBlock; // 18. K0PerBlock oss << "," << kAK1; // 19. AK1 - oss << "," << kBK1; // 19,5. BK1 + oss << "," << kBK1; // 19. BK1 oss << "," << kMPerXDL; // 20. MPerXDL oss << "," << kNPerXDL; // 21. NPerXDL oss << "," << kMXdlPerWave; // 22. MXdlPerWave diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index dea451c43f..476df4a776 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -196,6 +196,7 @@ 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 conv/ck/test_ckb_conv_bwd_data_multi_d_wmma_cshuffle.cpp + conv/ck/test_ckb_conv_bwd_data_multi_d_wmma_cshuffle_v3.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_wmma_cshuffle_v3.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_wmma_cshuffle_v3.cpp new file mode 100644 index 0000000000..eee9d86947 --- /dev/null +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_data_multi_d_wmma_cshuffle_v3.cpp @@ -0,0 +1,45 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT +#include "gmock/gmock.h" +#include "utils/ckb_conv_test_configs.hpp" +#include "utils/ckb_conv_test_utils.hpp" +#include "utils/conv_algorithm_type_utils.hpp" +#include "ck_tile/host/device_prop.hpp" + +namespace ckb = ck_tile::builder; +namespace ckt = ck_tile::builder::test; +namespace cku = ck_tile::builder::test_utils; + +constexpr auto SIGNATURE = + ckt::ConvSignature{.spatial_dim = 2, + .direction = ckb::ConvDirection::BACKWARD_DATA, + .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}}}; + +constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle_V3{} + .with_thread_block(cku::ThreadBlock_64_32x32x32) + .with_gemm_config(cku::GemmParamsABK1_Wmma_16x16_2x1_per_wave) + .with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3) + .with_bwd_data_specialization(ckb::ConvSpecialization::DEFAULT) + .with_prefetch_config(1, ckb::PipelineScheduler::DEFAULT) + .with_gemm_pad_params(0, 0) + .with_block_gemm(cku::BlockGemmDesc_v1_intrawave) + .with_transpose_params(2, 2); + +using Builder = ckb::ConvBuilder; +using Instance = Builder::Instance; + +TEST(BwdData_2DFp16_MultiD_Wmma_CShuffle_V3_GNHWC, Create) +{ + const auto expected_transfer_parameters = to_string(ALGORITHM); + std::cout << "Expected Transfer Parameters: " << expected_transfer_parameters << std::endl; + cku::run_test({"DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3", + expected_transfer_parameters, + "Default", + "GNHWK,GKYXC,EmptyTuple,GNHWC", + "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 b193348fc4..e403d8e255 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -72,6 +72,16 @@ struct GridwiseWmmaGemm size_t n_wmma_per_wave = 0; }; static_assert(ckb::GridwiseWmmaGemmDescriptor); +struct GridwiseWmmaGemmABK1 +{ + size_t ak1 = 0; + size_t bk1 = 0; + size_t m_per_wmma = 0; + size_t n_per_wmma = 0; + size_t m_wmma_per_wave = 0; + size_t n_wmma_per_wave = 0; +}; +static_assert(ckb::GridwiseWmmaGemmDescriptor); struct BlockGemmPipeline { @@ -227,6 +237,11 @@ struct WmmaGemm_ GridwiseWmmaGemm gridwise_gemm; }; +struct WmmaGemmABK1_ +{ + GridwiseWmmaGemmABK1 gridwise_gemm; +}; + template struct Transfer_ { @@ -426,6 +441,10 @@ struct ConvAlgorithmTemplate : Components... { result.gridwise_gemm = gemm; } + else if constexpr(std::is_base_of_v) + { + result.gridwise_gemm = gemm; + } else { static_assert(false, "Unrecognized GemmConfig type"); @@ -739,7 +758,6 @@ using ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle = TransposeParams_, GemmPad_>; -// Bwd Data algorithm types using ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle = ConvAlgorithmTemplate; +using ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle_V3 = + ConvAlgorithmTemplate, + ConvSpecializationBwdData_, + BlockGemm_, + MultipleDSpecialization_, + Prefetch_, + TransposeParams_, + GemmPad_>; + } // namespace ck_tile::builder::test diff --git a/experimental/builder/test/utils/ckb_conv_test_configs.hpp b/experimental/builder/test/utils/ckb_conv_test_configs.hpp index 4c7f0acdef..8b7d68f8db 100644 --- a/experimental/builder/test/utils/ckb_conv_test_configs.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_configs.hpp @@ -303,6 +303,13 @@ constexpr GridwiseWmmaGemm GemmParams_Wmma_2x1_per_wave{ constexpr GridwiseWmmaGemm GemmParams_Wmma_16x16_2x1_per_wave{ .k1 = 8, .m_per_wmma = 16, .n_per_wmma = 16, .m_wmma_per_wave = 2, .n_wmma_per_wave = 1}; +constexpr GridwiseWmmaGemmABK1 GemmParamsABK1_Wmma_16x16_2x1_per_wave{.ak1 = 8, + .bk1 = 8, + .m_per_wmma = 16, + .n_per_wmma = 16, + .m_wmma_per_wave = 2, + .n_wmma_per_wave = 1}; + constexpr ThreadBlock ThreadBlock_256_256x256x32{.block_size = 256, .tile_size = {.m = 256, .n = 256, .k = 32}}; diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp index 987dc73660..cc304c2fca 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -121,6 +121,15 @@ inline std::string to_string(GridwiseWmmaGemm t) return oss.str(); } +template <> +inline std::string to_string(GridwiseWmmaGemmABK1 t) +{ + std::ostringstream oss; + oss << t.ak1 << "," << t.bk1 << "," << t.m_per_wmma << "," << t.n_per_wmma << "," + << t.m_wmma_per_wave << "," << t.n_wmma_per_wave; + return oss.str(); +} + template <> inline std::string to_string(BlockGemmPipeline t) { @@ -304,6 +313,12 @@ inline std::string to_string(WmmaGemm_ t) return to_string(t.gridwise_gemm); } +template <> +inline std::string to_string(WmmaGemmABK1_ t) +{ + return to_string(t.gridwise_gemm); +} + template inline std::string to_string(Transfer_ t) { @@ -540,4 +555,16 @@ inline std::string to_string +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffle_V3 t) +{ + std::ostringstream oss; + oss << to_string(static_cast(t)) << "," + << to_string(static_cast(t)) << "," + << to_string(static_cast>(t)); + return oss.str(); + return oss.str(); +} + } // namespace ck_tile::builder::test diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp index dfdfd53725..bcce4ef9ca 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp @@ -28,6 +28,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_wmma_cshuffle_v3.hpp" +#endif + namespace ck { namespace tensor_operation { namespace device { @@ -1985,8 +1990,27 @@ struct DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3 "The argument pointer is not an object of " "DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3::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_wmma_cshuffle_v3.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 } // namespace tensor_operation } // namespace ck +