From 32a26d371b537c14f38023350b7db7f7c5d5d5b9 Mon Sep 17 00:00:00 2001 From: "assistant-librarian[bot]" Date: Tue, 4 Nov 2025 15:13:12 +0000 Subject: [PATCH] Merge commit '0be0288f58879123c228373525c4b438d354694f' into develop --- .../include/ck_tile/builder/builder_utils.hpp | 2 +- .../builder/conv_algorithm_concepts.hpp | 2 +- .../ck_tile/builder/conv_algorithm_limits.hpp | 2 +- .../include/ck_tile/builder/conv_builder.hpp | 2 +- .../include/ck_tile/builder/conv_factory.hpp | 2 +- .../builder/conv_signature_concepts.hpp | 2 +- .../builder/conv_signature_predicates.hpp | 2 +- .../builder/reflect/instance_traits.hpp | 2 +- ...e_grouped_conv_bwd_weight_xdl_cshuffle.hpp | 285 ++++++++++++++++++ ..._conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp | 2 +- ...ped_conv_fwd_multiple_abd_xdl_cshuffle.hpp | 2 +- ..._conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp | 2 +- ...uped_conv_fwd_multiple_d_wmma_cshuffle.hpp | 2 +- ...d_multiple_d_xdl_large_tensor_cshuffle.hpp | 2 +- .../builder/reflect/instance_traits_util.hpp | 17 +- .../builder/include/ck_tile/builder/types.hpp | 2 +- .../include/ck_tile/builder/versions.hpp | 3 + experimental/builder/test/CMakeLists.txt | 4 +- .../test/conv/test_ckb_conv_fwd_1d_bf16.cpp | 3 + .../test/conv/test_ckb_conv_fwd_1d_fp16.cpp | 3 + .../test/conv/test_ckb_conv_fwd_1d_i8.cpp | 3 + .../test/conv/test_ckb_conv_fwd_2d_bf16.cpp | 3 + .../test/conv/test_ckb_conv_fwd_2d_fp16.cpp | 3 + .../test/conv/test_ckb_conv_fwd_2d_fp32.cpp | 3 + .../test/conv/test_ckb_conv_fwd_3d_bf16.cpp | 3 + .../test/conv/test_ckb_conv_fwd_3d_fp16.cpp | 3 + .../test/conv/test_ckb_conv_fwd_3d_fp32.cpp | 3 + .../test/impl/conv_algorithm_types.hpp | 2 +- .../test/impl/conv_signature_types.hpp | 2 +- .../test/test_bwd_weight_instance_traits.cpp | 112 +++++++ ...ck_factory_grouped_convolution_forward.cpp | 2 +- ...d_convolution_forward_bias_bnorm_clamp.cpp | 2 +- ...grouped_convolution_forward_bias_clamp.cpp | 2 +- ...y_grouped_convolution_forward_bilinear.cpp | 2 +- ...tory_grouped_convolution_forward_clamp.cpp | 2 +- ..._grouped_convolution_forward_convscale.cpp | 2 +- ...grouped_convolution_forward_dynamic_op.cpp | 2 +- ...tory_grouped_convolution_forward_scale.cpp | 2 +- ...rouped_convolution_forward_scaleadd_ab.cpp | 2 +- ...olution_forward_scaleadd_scaleadd_relu.cpp | 2 +- .../builder/test/test_conv_builder.cpp | 3 + .../builder/test/test_fwd_instance_traits.cpp | 2 +- ...nstance_string_bwd_weight_grp_conv_xdl.cpp | 86 ++++++ .../test_get_instance_string_fwd_grp_conv.cpp | 2 +- ...st_get_instance_string_fwd_grp_conv_dl.cpp | 2 +- ...tance_string_fwd_grp_conv_large_tensor.cpp | 2 +- ...st_get_instance_string_fwd_grp_conv_v3.cpp | 2 +- ..._get_instance_string_fwd_grp_conv_wmma.cpp | 2 +- .../builder/test/test_inline_diff.cpp | 2 +- .../test/test_instance_traits_util.cpp | 2 +- .../builder/test/test_testing_utils.cpp | 2 +- experimental/builder/test/testing_utils.cpp | 2 +- experimental/builder/test/testing_utils.hpp | 2 +- .../test/utils/ckb_conv_test_common.hpp | 2 +- ...e_grouped_conv_bwd_weight_xdl_cshuffle.hpp | 19 +- 55 files changed, 591 insertions(+), 41 deletions(-) create mode 100644 experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp create mode 100644 experimental/builder/test/test_bwd_weight_instance_traits.cpp create mode 100644 experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp diff --git a/experimental/builder/include/ck_tile/builder/builder_utils.hpp b/experimental/builder/include/ck_tile/builder/builder_utils.hpp index a69471c9ed..5b4981c630 100644 --- a/experimental/builder/include/ck_tile/builder/builder_utils.hpp +++ b/experimental/builder/include/ck_tile/builder/builder_utils.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once 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 586a119c75..365835684e 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp index 68d5ec5a83..b24d0f47e9 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_limits.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/include/ck_tile/builder/conv_builder.hpp b/experimental/builder/include/ck_tile/builder/conv_builder.hpp index d74948709b..bfb7386ea1 100644 --- a/experimental/builder/include/ck_tile/builder/conv_builder.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_builder.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/include/ck_tile/builder/conv_factory.hpp b/experimental/builder/include/ck_tile/builder/conv_factory.hpp index 8ea3e18d65..a3932f524c 100644 --- a/experimental/builder/include/ck_tile/builder/conv_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_factory.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // A factory for instantiating CK convolution kernels. // diff --git a/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp index 370e7b6521..742dfbb89c 100644 --- a/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // This file defines the compile-time "signature" for grouped convolution operations. // A signature is a collection of properties that fully describe a convolution kernel's diff --git a/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp b/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp index f947c7e329..f016a342d3 100644 --- a/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_signature_predicates.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp index a47ad0ef57..c9b45691cc 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // Compile-time reflection for CK device kernel instances. // diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp new file mode 100644 index 0000000000..a0def3e5d9 --- /dev/null +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp @@ -0,0 +1,285 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "instance_traits.hpp" +#include "ck/tensor_operation/gpu/device/convolution_backward_weight_specialization.hpp" + +// Forward declaration to avoid circular dependency +namespace ck::tensor_operation::device { + +template +struct DeviceGroupedConvBwdWeight_Xdl_CShuffle; + +} // namespace ck::tensor_operation::device + +namespace ck_tile { +namespace reflect { + +template +struct InstanceTraits> +{ + static constexpr auto kTensorOpName = "DeviceGroupedConvBwdWeight_Xdl_CShuffle"; + + static constexpr ck::index_t kNDimSpatial = NDimSpatial; + + using InLayout = InLayout_; + using WeiLayout = WeiLayout_; + using OutLayout = OutLayout_; + + using InDataType = InDataType_; + using WeiDataType = WeiDataType_; + using OutDataType = OutDataType_; + using AccDataType = AccDataType_; + + using InElementwiseOperation = InElementwiseOperation_; + using WeiElementwiseOperation = WeiElementwiseOperation_; + using OutElementwiseOperation = OutElementwiseOperation_; + + static constexpr auto kConvBackwardWeightSpecialization = ConvBackwardWeightSpecialization; + + 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 kK1 = K1; + 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; + + using ABlockTransferThreadClusterLengths_K0_M_K1 = ABlockTransferThreadClusterLengths_K0_M_K1_; + using ABlockTransferThreadClusterArrangeOrder = ABlockTransferThreadClusterArrangeOrder_; + using ABlockTransferSrcAccessOrder = ABlockTransferSrcAccessOrder_; + static constexpr ck::index_t kABlockTransferSrcVectorDim = ABlockTransferSrcVectorDim; + static constexpr ck::index_t kABlockTransferSrcScalarPerVector = + ABlockTransferSrcScalarPerVector; + static constexpr ck::index_t kABlockTransferDstScalarPerVector_K1 = + ABlockTransferDstScalarPerVector_K1; + static constexpr bool kABlockLdsAddExtraM = ABlockLdsAddExtraM; + + using BBlockTransferThreadClusterLengths_K0_N_K1 = BBlockTransferThreadClusterLengths_K0_N_K1_; + using BBlockTransferThreadClusterArrangeOrder = BBlockTransferThreadClusterArrangeOrder_; + using BBlockTransferSrcAccessOrder = BBlockTransferSrcAccessOrder_; + static constexpr ck::index_t kBBlockTransferSrcVectorDim = BBlockTransferSrcVectorDim; + static constexpr ck::index_t kBBlockTransferSrcScalarPerVector = + BBlockTransferSrcScalarPerVector; + static constexpr ck::index_t kBBlockTransferDstScalarPerVector_K1 = + BBlockTransferDstScalarPerVector_K1; + static constexpr bool kBBlockLdsAddExtraN = BBlockLdsAddExtraN; + + static constexpr ck::index_t kCShuffleMXdlPerWavePerShuffle = CShuffleMXdlPerWavePerShuffle; + static constexpr ck::index_t kCShuffleNXdlPerWavePerShuffle = CShuffleNXdlPerWavePerShuffle; + + using CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock = + CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock_; + static constexpr ck::index_t kCBlockTransferScalarPerVector_NWaveNPerXdl = + CBlockTransferScalarPerVector_NWaveNPerXdl; + + using ComputeTypeA = ComputeTypeA_; + using ComputeTypeB = ComputeTypeB_; + + static constexpr ck::index_t kMaxTransposeTransferSrcScalarPerVector = + MaxTransposeTransferSrcScalarPerVector; + static constexpr ck::index_t kMaxTransposeTransferDstScalarPerVector = + MaxTransposeTransferDstScalarPerVector; + + // Static member function to generate instance string + static std::string instance_string() + { + std::ostringstream oss; + + // Kernel type name + oss << "DeviceGroupedConvBwdWeight_Xdl_CShuffle"; + + // Template parameters in exact order + oss << "<" << kNDimSpatial; // 1. NDimSpatial + oss << "," << detail::layout_name(); // 2. InLayout + oss << "," << detail::layout_name(); // 3. WeiLayout + oss << "," << detail::layout_name(); // 4. OutLayout + oss << "," << detail::type_name(); // 5. InDataType + oss << "," << detail::type_name(); // 6. WeiDataType + oss << "," << detail::type_name(); // 7. OutDataType + oss << "," << detail::type_name(); // 8. AccDataType + oss << "," + << detail::elementwise_op_name(); // 9. InElementwiseOperation + oss << "," + << detail::elementwise_op_name(); // 10. + // WeiElementwiseOperation + oss << "," + << detail::elementwise_op_name(); // 11. + // OutElementwiseOperation + oss << "," + << detail::conv_bwd_weight_spec_name( + kConvBackwardWeightSpecialization); // 12. ConvBackwardWeightSpecialization + oss << "," << kBlockSize; // 13. BlockSize + oss << "," << kMPerBlock; // 14. MPerBlock + oss << "," << kNPerBlock; // 15. NPerBlock + oss << "," << kK0PerBlock; // 16. K0PerBlock + oss << "," << kK1; // 17. K1 + oss << "," << kMPerXDL; // 18. MPerXDL + oss << "," << kNPerXDL; // 19. NPerXDL + oss << "," << kMXdlPerWave; // 20. MXdlPerWave + oss << "," << kNXdlPerWave; // 21. NXdlPerWave + oss << "," << detail::sequence_name(); // 22. + oss << "," << detail::sequence_name(); // 23. + oss << "," << detail::sequence_name(); // 24. + oss << "," << kABlockTransferSrcVectorDim; // 25. + oss << "," << kABlockTransferSrcScalarPerVector; // 26. + oss << "," << kABlockTransferDstScalarPerVector_K1; // 27. + oss << "," << (kABlockLdsAddExtraM ? "true" : "false"); // 28. + oss << "," << detail::sequence_name(); // 29. + oss << "," << detail::sequence_name(); // 30. + oss << "," << detail::sequence_name(); // 31. + oss << "," << kBBlockTransferSrcVectorDim; // 32. + oss << "," << kBBlockTransferSrcScalarPerVector; // 33. + oss << "," << kBBlockTransferDstScalarPerVector_K1; // 34. + oss << "," << (kBBlockLdsAddExtraN ? "true" : "false"); // 35. + oss << "," << kCShuffleMXdlPerWavePerShuffle; // 36. + oss << "," << kCShuffleNXdlPerWavePerShuffle; // 37. + oss << "," + << detail::sequence_name< + CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock>(); // 38. + oss << "," << kCBlockTransferScalarPerVector_NWaveNPerXdl; // 39. + oss << "," << detail::type_name(); // 40. + oss << "," << detail::type_name(); // 41. + oss << "," << kMaxTransposeTransferSrcScalarPerVector; // 42. + oss << "," << kMaxTransposeTransferDstScalarPerVector; // 43. + oss << ">"; + + return oss.str(); + } +}; + +} // namespace reflect +} // namespace ck_tile diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp index 54ba224f7f..8756825c3f 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // InstanceTraits specialization for DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK // diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp index 346b731c33..572b04e75b 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // InstanceTraits specialization for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle // diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp index 5784938fc6..7eb4c883b0 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // InstanceTraits specialization for DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 // diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp index 9f64c7d6e8..60f991b1fc 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // InstanceTraits specialization for DeviceGroupedConvFwdMultipleD_Wmma_CShuffle // diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp index b41aed784c..f60d5e7b68 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // InstanceTraits specialization for DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor // 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 b13675a7b9..95d1c94de4 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 @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // Utility functions and helpers for instance_traits.hpp // Contains helper functions to convert types, enums, and sequences to string representations. @@ -21,6 +21,7 @@ #include #include #include +#include #include namespace ck_tile::reflect::detail { @@ -112,6 +113,20 @@ conv_fwd_spec_name(ck::tensor_operation::device::ConvolutionForwardSpecializatio } } +// Convert ConvolutionBackwardWeightSpecialization enum to string +constexpr std::string_view conv_bwd_weight_spec_name( + ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization spec) +{ + using enum ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization; + switch(spec) + { + case Default: return "Default"; + case Filter1x1Stride1Pad0: return "Filter1x1Stride1Pad0"; + case Filter1x1Pad0: return "Filter1x1Pad0"; + case OddC: return "OddC"; + } +} + // Convert GemmSpecialization enum to string constexpr std::string_view gemm_spec_name(ck::tensor_operation::device::GemmSpecialization spec) { diff --git a/experimental/builder/include/ck_tile/builder/types.hpp b/experimental/builder/include/ck_tile/builder/types.hpp index a2ef89da2e..2650f0de16 100644 --- a/experimental/builder/include/ck_tile/builder/types.hpp +++ b/experimental/builder/include/ck_tile/builder/types.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/include/ck_tile/builder/versions.hpp b/experimental/builder/include/ck_tile/builder/versions.hpp index e8fb2fe4de..4ee45b730b 100644 --- a/experimental/builder/include/ck_tile/builder/versions.hpp +++ b/experimental/builder/include/ck_tile/builder/versions.hpp @@ -1,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once #include diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index 26a666a805..8b5c4519a9 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -20,6 +20,7 @@ endfunction() add_ck_builder_test(test_ckb_conv_builder test_conv_builder.cpp test_fwd_instance_traits.cpp + test_bwd_weight_instance_traits.cpp test_instance_traits_util.cpp) add_ck_builder_test(test_ckb_inline_diff test_inline_diff.cpp) @@ -30,7 +31,8 @@ add_ck_builder_test(test_ckb_get_instance_string test_get_instance_string_fwd_grp_conv.cpp test_get_instance_string_fwd_grp_conv_large_tensor.cpp test_get_instance_string_fwd_grp_conv_wmma.cpp - test_get_instance_string_fwd_grp_conv_dl.cpp) + test_get_instance_string_fwd_grp_conv_dl.cpp + test_get_instance_string_bwd_weight_grp_conv_xdl.cpp) # Testing the fwd convolution builder requires kernel compilation. # To enable parallel compilation, the individual tests are split into separate files. 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..b58de836de 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..a83ca84297 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..3ceac2a047 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..b8dbf2ca97 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..aba2f29ffd 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..4d01323600 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..a30158aa8e 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..c0b2e613a2 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; 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..0fea260eac 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,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include "utils/ckb_conv_test_common.hpp" using namespace ck_tile::builder::test_utils; diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp index 9c5ca9b97b..1a78028862 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/test/impl/conv_signature_types.hpp b/experimental/builder/test/impl/conv_signature_types.hpp index cc5490c711..5e6684c4cd 100644 --- a/experimental/builder/test/impl/conv_signature_types.hpp +++ b/experimental/builder/test/impl/conv_signature_types.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/experimental/builder/test/test_bwd_weight_instance_traits.cpp b/experimental/builder/test/test_bwd_weight_instance_traits.cpp new file mode 100644 index 0000000000..24c28c2b9d --- /dev/null +++ b/experimental/builder/test/test_bwd_weight_instance_traits.cpp @@ -0,0 +1,112 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +namespace { + +TEST(InstanceTraits, BwdWeightXdlCShuffleInstanceStringReturnsCorrectFormat) +{ + using DeviceInstance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle< + 2, // NDimSpatial + ck::tensor_layout::convolution::GNHWC, // InLayout + ck::tensor_layout::convolution::GKYXC, // WeiLayout + ck::tensor_layout::convolution::GNHWK, // OutLayout + ck::half_t, // InDataType + ck::half_t, // WeiDataType + ck::half_t, // OutDataType + float, // AccDataType + ck::tensor_operation::element_wise::PassThrough, // InElementwiseOperation + ck::tensor_operation::element_wise::PassThrough, // WeiElementwiseOperation + ck::tensor_operation::element_wise::PassThrough, // OutElementwiseOperation + ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization:: + Default, // ConvBackwardWeightSpecialization + 256, // BlockSize + 128, // MPerBlock + 128, // NPerBlock + 4, // K0PerBlock + 8, // K1 + 32, // MPerXDL + 32, // NPerXDL + 2, // MXdlPerWave + 2, // NXdlPerWave + ck::Sequence<4, 64, 1>, // ABlockTransferThreadClusterLengths_K0_M_K1 + ck::Sequence<1, 0, 2>, // ABlockTransferThreadClusterArrangeOrder + ck::Sequence<1, 0, 2>, // ABlockTransferSrcAccessOrder + 2, // ABlockTransferSrcVectorDim + 8, // ABlockTransferSrcScalarPerVector + 8, // ABlockTransferDstScalarPerVector_K1 + false, // ABlockLdsAddExtraM + ck::Sequence<4, 64, 1>, // BBlockTransferThreadClusterLengths_K0_N_K1 + ck::Sequence<1, 0, 2>, // BBlockTransferThreadClusterArrangeOrder + ck::Sequence<1, 0, 2>, // BBlockTransferSrcAccessOrder + 2, // BBlockTransferSrcVectorDim + 8, // BBlockTransferSrcScalarPerVector + 8, // BBlockTransferDstScalarPerVector_K1 + false, // BBlockLdsAddExtraN + 1, // CShuffleMXdlPerWavePerShuffle + 1, // CShuffleNXdlPerWavePerShuffle + ck::Sequence<1, + 32, + 1, + 8>, // CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock + 8, // CBlockTransferScalarPerVector_NWaveNPerXdl + ck::half_t, // ComputeTypeA + ck::half_t, // ComputeTypeB + 1, // MaxTransposeTransferSrcScalarPerVector + 1>; // MaxTransposeTransferDstScalarPerVector + + std::string instance_str = ck_tile::reflect::instance_string(); + + std::string expected_str = "DeviceGroupedConvBwdWeight_Xdl_CShuffle" + "<2" // NDimSpatial + ",GNHWC" // InLayout + ",GKYXC" // WeiLayout + ",GNHWK" // OutLayout + ",fp16" // InDataType + ",fp16" // WeiDataType + ",fp16" // OutDataType + ",fp32" // AccDataType + ",PassThrough" // InElementwiseOperation + ",PassThrough" // WeiElementwiseOperation + ",PassThrough" // OutElementwiseOperation + ",Default" // ConvBackwardWeightSpecialization + ",256" // BlockSize + ",128" // MPerBlock + ",128" // NPerBlock + ",4" // K0PerBlock + ",8" // K1 + ",32" // MPerXDL + ",32" // NPerXDL + ",2" // MXdlPerWave + ",2" // NXdlPerWave + ",Seq(4,64,1)" // ABlockTransferThreadClusterLengths_K0_M_K1 + ",Seq(1,0,2)" // ABlockTransferThreadClusterArrangeOrder + ",Seq(1,0,2)" // ABlockTransferSrcAccessOrder + ",2" // ABlockTransferSrcVectorDim + ",8" // ABlockTransferSrcScalarPerVector + ",8" // ABlockTransferDstScalarPerVector_K1 + ",false" // ABlockLdsAddExtraM + ",Seq(4,64,1)" // BBlockTransferThreadClusterLengths_K0_N_K1 + ",Seq(1,0,2)" // BBlockTransferThreadClusterArrangeOrder + ",Seq(1,0,2)" // BBlockTransferSrcAccessOrder + ",2" // BBlockTransferSrcVectorDim + ",8" // BBlockTransferSrcScalarPerVector + ",8" // BBlockTransferDstScalarPerVector_K1 + ",false" // BBlockLdsAddExtraN + ",1" // CShuffleMXdlPerWavePerShuffle + ",1" // CShuffleNXdlPerWavePerShuffle + ",Seq(1,32,1,8)" // CBlockTransferClusterLengths + ",8" // CBlockTransferScalarPerVector_NWaveNPerXdl + ",fp16" // ComputeTypeA + ",fp16" // ComputeTypeB + ",1" // MaxTransposeTransferSrcScalarPerVector + ",1>"; // MaxTransposeTransferDstScalarPerVector + + EXPECT_EQ(instance_str, expected_str); +} + +} // anonymous namespace diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp index 53199f1d77..73a731a69c 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp index c68fdb4b24..75cb58018e 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp index a9e5e39e8f..b301bab966 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp index cc9cbd3dda..8b77b11e0c 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include "ck/utility/data_type.hpp" diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp index a8d1d1763f..e678fa1258 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include "ck/utility/data_type.hpp" diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp index 0da300e3d8..4882c6fde2 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp index e918785fd7..7437385e2a 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include "ck/utility/data_type.hpp" diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp index 428d1c81f3..8144d7bedd 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp index 774c30c05e..beebc3f853 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp index ba8726a643..79feb298cd 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_conv_builder.cpp b/experimental/builder/test/test_conv_builder.cpp index 4ec189daa4..27b30946b2 100644 --- a/experimental/builder/test/test_conv_builder.cpp +++ b/experimental/builder/test/test_conv_builder.cpp @@ -1,3 +1,6 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include class ConvBuilderTest : public ::testing::Test diff --git a/experimental/builder/test/test_fwd_instance_traits.cpp b/experimental/builder/test/test_fwd_instance_traits.cpp index 2e3b6ac264..b57b20eb7d 100644 --- a/experimental/builder/test/test_fwd_instance_traits.cpp +++ b/experimental/builder/test/test_fwd_instance_traits.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp b/experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp new file mode 100644 index 0000000000..68b43c6a99 --- /dev/null +++ b/experimental/builder/test/test_get_instance_string_bwd_weight_grp_conv_xdl.cpp @@ -0,0 +1,86 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +// Test GetInstanceString through base class pointer for backward weight XDL variant +TEST(GetInstanceString, ReturnsStringForBwdWeightGrpConvXdlInstance) +{ + // Use the template helper to get a working instance configuration + using InstanceTuple = ck::tensor_operation::device::instance:: + device_grouped_conv_bwd_weight_xdl_c_shuffle_f16_instances< + 2, // NDimSpatial + ck::tensor_operation::device::instance::GNHWC, // InLayout + ck::tensor_operation::device::instance::GKYXC, // WeiLayout + ck::tensor_operation::device::instance::GNHWK, // OutLayout + ck::tensor_operation::device::instance:: + ConvBwdWeightDefault>; // ConvBwdWeightSpecialization + + // Get the first instance from the tuple + using DeviceInstance = typename std::tuple_element<0, InstanceTuple>::type; + + // Define the base class type using the most general operator base + using BaseClass = ck::tensor_operation::device::BaseOperator; + + // Create an instance of the derived class + DeviceInstance device_instance; + + // Get a pointer to the base class + BaseClass* base_ptr = &device_instance; + + // Call GetInstanceString through the base class pointer + std::string instance_str = base_ptr->GetInstanceString(); + + // Expected complete instance string based on the first instance from + // device_grouped_conv_bwd_weight_xdl_c_shuffle_f16_instances + // This corresponds to the configuration with BlockSize=64, MPerBlock=64, NPerBlock=64, etc. + std::string expected_str = "DeviceGroupedConvBwdWeight_Xdl_CShuffle" + "<2" // NDimSpatial + ",GNHWC" // InLayout + ",GKYXC" // WeiLayout + ",GNHWK" // OutLayout + ",fp16" // InDataType + ",fp16" // WeiDataType + ",fp16" // OutDataType + ",fp32" // AccDataType + ",PassThrough" // InElementwiseOperation + ",PassThrough" // WeiElementwiseOperation + ",PassThrough" // OutElementwiseOperation + ",Default" // ConvBackwardWeightSpecialization + ",64" // BlockSize + ",64" // MPerBlock + ",64" // NPerBlock + ",4" // K0PerBlock + ",8" // K1 + ",32" // MPerXDL + ",32" // NPerXDL + ",2" // MXdlPerWave + ",2" // NXdlPerWave + ",Seq(1,4,8,2)" // ABlockTransferThreadClusterLengths_K0_M_K1 + ",Seq(0,3,1,2)" // ABlockTransferThreadClusterArrangeOrder + ",Seq(0,2,1,3)" // ABlockTransferSrcAccessOrder + ",2" // ABlockTransferSrcVectorDim + ",2" // ABlockTransferSrcScalarPerVector + ",4" // ABlockTransferDstScalarPerVector_K1 + ",true" // ABlockLdsAddExtraM + ",Seq(1,4,8,2)" // BBlockTransferThreadClusterLengths_K0_N_K1 + ",Seq(0,3,1,2)" // BBlockTransferThreadClusterArrangeOrder + ",Seq(0,2,1,3)" // BBlockTransferSrcAccessOrder + ",2" // BBlockTransferSrcVectorDim + ",2" // BBlockTransferSrcScalarPerVector + ",4" // BBlockTransferDstScalarPerVector_K1 + ",true" // BBlockLdsAddExtraN + ",1" // CShuffleMXdlPerWavePerShuffle + ",1" // CShuffleNXdlPerWavePerShuffle + ",Seq(1,16,1,4)" // CBlockTransferClusterLengths + ",2" // CBlockTransferScalarPerVector_NWaveNPerXdl + ",fp16" // ComputeTypeA + ",fp16" // ComputeTypeB + ",1" // MaxTransposeTransferSrcScalarPerVector + ",1>"; // MaxTransposeTransferDstScalarPerVector + + EXPECT_EQ(instance_str, expected_str); +} diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp index 7c43c428ce..ca683abedc 100644 --- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp +++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp index 54e026b308..abf55322b6 100644 --- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp +++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_dl.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp index fd75fd9e6b..d6900d041f 100644 --- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp +++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_large_tensor.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp index 45ab5db322..67e55173de 100644 --- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp +++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_v3.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp index e1505659f5..ccce080f9d 100644 --- a/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp +++ b/experimental/builder/test/test_get_instance_string_fwd_grp_conv_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_inline_diff.cpp b/experimental/builder/test/test_inline_diff.cpp index 41692fb40e..7a1b0ab8c3 100644 --- a/experimental/builder/test/test_inline_diff.cpp +++ b/experimental/builder/test/test_inline_diff.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/test_instance_traits_util.cpp b/experimental/builder/test/test_instance_traits_util.cpp index 04d3cb978b..5623028699 100644 --- a/experimental/builder/test/test_instance_traits_util.cpp +++ b/experimental/builder/test/test_instance_traits_util.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/test_testing_utils.cpp b/experimental/builder/test/test_testing_utils.cpp index 24a1c9bc81..89d8c06e59 100644 --- a/experimental/builder/test/test_testing_utils.cpp +++ b/experimental/builder/test/test_testing_utils.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/experimental/builder/test/testing_utils.cpp b/experimental/builder/test/testing_utils.cpp index 34793b601e..02663317f1 100644 --- a/experimental/builder/test/testing_utils.cpp +++ b/experimental/builder/test/testing_utils.cpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #include "testing_utils.hpp" #include diff --git a/experimental/builder/test/testing_utils.hpp b/experimental/builder/test/testing_utils.hpp index 3ff2eb32de..ae5811bd4b 100644 --- a/experimental/builder/test/testing_utils.hpp +++ b/experimental/builder/test/testing_utils.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/experimental/builder/test/utils/ckb_conv_test_common.hpp b/experimental/builder/test/utils/ckb_conv_test_common.hpp index d18a008015..d145bdfd6c 100644 --- a/experimental/builder/test/utils/ckb_conv_test_common.hpp +++ b/experimental/builder/test/utils/ckb_conv_test_common.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp index c7ee3e9ecf..650c6f11d3 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp @@ -1,5 +1,5 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -24,6 +24,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#ifdef CK_EXPERIMENTAL_BUILDER +#include "ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp" +#endif + namespace ck { namespace tensor_operation { namespace device { @@ -1225,6 +1229,19 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle return str.str(); } +#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_weight_xdl_cshuffle.hpp " + "for the given template parameters."); + return ck_tile::reflect::instance_string(); + } +#endif + size_t GetWorkSpaceSize(const BaseArgument* p_arg) const override { auto arg = dynamic_cast(p_arg);