From 92d1b386b2706bbf6f99225cd1c0062047ed4790 Mon Sep 17 00:00:00 2001 From: Mateusz Ozga <110818320+mozga-amd@users.noreply.github.com> Date: Thu, 12 Sep 2024 11:47:52 +0200 Subject: [PATCH] Pool2d max/avg kernel in the BWD version (#1494) * Add pool2d instance BWD AVG * Add pool2d instance BWD MAX * Fix: avg review * Fix review: part2 * Fix - enable test when type is compiled * Fix review part3 [ROCm/composable_kernel commit: 448c0f56d81df9664143013005249cb47b391f7b] --- .../impl/device_avgpool2d_bwd_nhwc_nhwc.hpp | 523 ++++++++++++++++++ .../element/unary_element_wise_operation.hpp | 29 +- include/ck/utility/reduction_operator.hpp | 175 +++++- .../gpu/avg_pool2d_bwd.hpp | 80 +++ .../gpu/max_pool_bwd.hpp | 13 +- .../gpu/avg_pool2d_bwd/CMakeLists.txt | 8 + ...vice_avg_pool2d_bwd_nhwc_bf16_instance.cpp | 21 + ...evice_avg_pool2d_bwd_nhwc_f16_instance.cpp | 21 + ...evice_avg_pool2d_bwd_nhwc_f32_instance.cpp | 21 + ...device_avg_pool2d_bwd_nhwc_f8_instance.cpp | 20 + ...ce_avg_pool2d_bwd_nhwc_instance_common.hpp | 38 ++ ...vice_avg_pool2d_bwd_nhwc_int8_instance.cpp | 20 + .../gpu/max_pool_bwd/CMakeLists.txt | 3 +- .../device_max_pool_bwd_int8_instance.cpp | 20 + .../max_pool_bwd_instance_common.hpp | 4 +- profiler/include/profiler/data_type_enum.hpp | 3 +- .../profiler/profile_avg_pool2d_bwd_impl.hpp | 255 +++++++++ .../profiler/profile_max_pool2d_bwd_impl.hpp | 295 ++++++++++ profiler/src/CMakeLists.txt | 3 + profiler/src/profile_avg_pool2d_bwd.cpp | 188 +++++++ profiler/src/profile_max_pool2d_bwd.cpp | 178 ++++++ test/pool/CMakeLists.txt | 6 + test/pool/test_avg_pool2d_bwd.cpp | 133 +++++ test/pool/test_max_pool2d_bwd.cpp | 122 ++++ test/pool/test_pool_fwd_common.hpp | 5 +- 25 files changed, 2168 insertions(+), 16 deletions(-) create mode 100644 include/ck/tensor_operation/gpu/device/impl/device_avgpool2d_bwd_nhwc_nhwc.hpp create mode 100644 library/include/ck/library/tensor_operation_instance/gpu/avg_pool2d_bwd.hpp create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/CMakeLists.txt create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_bf16_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f16_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f32_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f8_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_instance_common.hpp create mode 100644 library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_int8_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/max_pool_bwd/device_max_pool_bwd_int8_instance.cpp create mode 100644 profiler/include/profiler/profile_avg_pool2d_bwd_impl.hpp create mode 100644 profiler/include/profiler/profile_max_pool2d_bwd_impl.hpp mode change 100755 => 100644 profiler/src/CMakeLists.txt create mode 100644 profiler/src/profile_avg_pool2d_bwd.cpp create mode 100644 profiler/src/profile_max_pool2d_bwd.cpp create mode 100644 test/pool/test_avg_pool2d_bwd.cpp create mode 100644 test/pool/test_max_pool2d_bwd.cpp diff --git a/include/ck/tensor_operation/gpu/device/impl/device_avgpool2d_bwd_nhwc_nhwc.hpp b/include/ck/tensor_operation/gpu/device/impl/device_avgpool2d_bwd_nhwc_nhwc.hpp new file mode 100644 index 0000000000..7fca3e2988 --- /dev/null +++ b/include/ck/tensor_operation/gpu/device/impl/device_avgpool2d_bwd_nhwc_nhwc.hpp @@ -0,0 +1,523 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include +#include + +#include "ck/tensor_description/tensor_descriptor.hpp" +#include "ck/tensor_description/tensor_descriptor_helper.hpp" +#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" +#include "ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp" +#include "ck/tensor_operation/gpu/device/device_avgpool_bwd.hpp" +#include "ck/tensor_operation/gpu/grid/gridwise_2d_reduction_threadwise.hpp" +#include "ck/host_utility/device_prop.hpp" +#include "ck/host_utility/kernel_launch.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { + +// In and Din = [N, C, Hi, Wi] +// Out and Dout = [N, C, Ho, Wo] +// Out = AvgPool2dFwd(In) +// Din = AvgPool2dBwd(Dout) +// Pooling dimension = H, W +template +struct DeviceAvgPool2dBwd_NHWC_NHWC : public DeviceAvgPoolBwd<2, + DOutDataType, + DInDataType, + tensor_layout::convolution::NHWC, + tensor_layout::convolution::NHWC> +{ + + static constexpr ck::index_t NDimSpatial = 2; + + static constexpr auto I0 = Number<0>{}; + static constexpr auto I1 = Number<1>{}; + + static constexpr ck::index_t M_BlockTileSize = MThreadClusterSize * MThreadSliceSize; + static constexpr ck::index_t K_BlockTileSize = KThreadClusterSize * KThreadSliceSize; + + static auto + Make2DGridDescriptor_Out_M_K_In_M(const std::vector& dout_n_c_wos_lengths, + const std::vector& din_n_c_wos_length, + const std::vector& dout_n_c_wos_strides, + const std::vector& din_n_c_wos_strides, + const std::vector& window_lengths, + const std::vector& window_strides, + const std::vector& window_dilations, + const std::vector& input_left_pads, + const std::vector& input_right_pads, + const std::vector& tildes) + { + index_t i_ytilde = tildes[0]; + index_t i_xtilde = tildes[1]; + + const index_t N = dout_n_c_wos_lengths[0]; + const index_t C = dout_n_c_wos_lengths[1]; + const index_t Ho = dout_n_c_wos_lengths[2]; + const index_t Wo = dout_n_c_wos_lengths[3]; + + const index_t Hi = din_n_c_wos_length[2]; + const index_t Wi = din_n_c_wos_length[3]; + + const index_t Y = window_lengths[0]; + const index_t X = window_lengths[1]; + + const index_t InLeftPadH = input_left_pads[0]; + const index_t InLeftPadW = input_left_pads[1]; + + const index_t InRightPadH = input_right_pads[0]; + const index_t InRightPadW = input_right_pads[1]; + + const index_t ConvStrideH = window_strides[0]; + const index_t ConvStrideW = window_strides[1]; + + const index_t ConvDilationH = window_dilations[0]; + const index_t ConvDilationW = window_dilations[1]; + + const index_t Ni_stride = dout_n_c_wos_strides[0]; + const index_t Ci_stride = dout_n_c_wos_strides[1]; + const index_t Ho_stride = dout_n_c_wos_strides[2]; + const index_t Wo_stride = dout_n_c_wos_strides[3]; + + const auto GcdStrideDilationH = math::gcd(ConvStrideH, ConvDilationH); + const auto GcdStrideDilationW = math::gcd(ConvStrideW, ConvDilationW); + + const auto YTilde = ConvStrideH / GcdStrideDilationH; + const auto XTilde = ConvStrideW / GcdStrideDilationW; + + const auto YDot = math::integer_divide_ceil(Y, YTilde); + const auto XDot = math::integer_divide_ceil(X, XTilde); + + const auto HTilde = Ho + math::integer_divide_ceil(ConvDilationH * (Y - I1), ConvStrideH); + const auto WTilde = Wo + math::integer_divide_ceil(ConvDilationW * (X - I1), ConvStrideW); + + // only work on Tildes that contribute to non-padding area of input tensor + const auto IHTildeSliceBegin = math::integer_divide_floor( + math::max(I0, InLeftPadH - ConvDilationH * (YTilde - I1)), ConvStrideH); + const auto IWTildeSliceBegin = math::integer_divide_floor( + math::max(I0, InLeftPadW - ConvDilationW * (XTilde - I1)), ConvStrideW); + + const auto IHTildeSliceEnd = + math::min(HTilde, math::integer_divide_ceil(InLeftPadH + Hi - I1, ConvStrideH) + I1); + const auto IWTildeSliceEnd = + math::min(WTilde, math::integer_divide_ceil(InLeftPadW + Wi - I1, ConvStrideW) + I1); + + const auto HTildeSlice = IHTildeSliceEnd - IHTildeSliceBegin; + const auto WTildeSlice = IWTildeSliceEnd - IWTildeSliceBegin; + + // ReduceK is different for each Reduce + const auto YDotSlice = math::integer_divide_ceil(Y - i_ytilde, YTilde); + const auto XDotSlice = math::integer_divide_ceil(X - i_xtilde, XTilde); + + // Problem size of reduction kernel + const index_t MRaw = N * HTildeSlice * WTildeSlice * C; + const index_t MPad = math::integer_least_multiple(MRaw, M_BlockTileSize) - MRaw; + + const index_t KRaw = YDotSlice * XDotSlice; + const index_t KPad = math::integer_least_multiple(KRaw, K_BlockTileSize) - KRaw; + + const auto out_n_ho_wo_c_grid_desc = make_naive_tensor_descriptor( + make_tuple(N, Ho, Wo, C), make_tuple(Ni_stride, Ho_stride, Wo_stride, Ci_stride)); + + // Out[ReduceM, ReduceK] + const auto out_n_hop_wop_c_grid_desc = transform_tensor_descriptor( + out_n_ho_wo_c_grid_desc, + make_tuple(make_pass_through_transform(N), + make_pad_transform(Ho, I0, I0), + make_pad_transform(Wo, I0, I0), + make_pass_through_transform(C)), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{})); + + const auto out_n_ydot_htilde_xdot_wtilde_c_grid_desc = transform_tensor_descriptor( + out_n_hop_wop_c_grid_desc, + make_tuple(make_pass_through_transform(N), + make_embed_transform(make_tuple(YDot, HTilde), + make_tuple(-ConvDilationH / GcdStrideDilationH, I1)), + make_embed_transform(make_tuple(XDot, WTilde), + make_tuple(-ConvDilationW / GcdStrideDilationW, I1)), + make_pass_through_transform(C)), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), + make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{})); + + const auto out_n_ydotslice_htildeslice_xdotslice_wtildeslice_c_grid_desc = + transform_tensor_descriptor( + out_n_ydot_htilde_xdot_wtilde_c_grid_desc, + make_tuple(make_pass_through_transform(N), + make_slice_transform(YDot, I0, YDotSlice), + make_slice_transform(HTilde, IHTildeSliceBegin, HTildeSlice), + make_slice_transform(XDot, I0, XDotSlice), + make_slice_transform(WTilde, IWTildeSliceBegin, WTildeSlice), + make_pass_through_transform(C)), + make_tuple(Sequence<0>{}, + Sequence<1>{}, + Sequence<2>{}, + Sequence<3>{}, + Sequence<4>{}, + Sequence<5>{}), + make_tuple(Sequence<0>{}, + Sequence<1>{}, + Sequence<2>{}, + Sequence<3>{}, + Sequence<4>{}, + Sequence<5>{})); + + const auto out_grid_desc_reducemraw_reducekraw = transform_tensor_descriptor( + out_n_ydotslice_htildeslice_xdotslice_wtildeslice_c_grid_desc, + make_tuple(make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice, C)), + make_merge_transform(make_tuple(YDotSlice, XDotSlice))), + make_tuple(Sequence<0, 2, 4, 5>{}, Sequence<1, 3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + const auto out_grid_desc_reducem_reducek = transform_tensor_descriptor( + out_grid_desc_reducemraw_reducekraw, + make_tuple(make_right_pad_transform(MRaw, MPad), make_right_pad_transform(KRaw, KPad)), + make_tuple(Sequence<0>{}, Sequence<1>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + // In[ReduceM] + const auto in_n_hi_wi_c_grid_desc = + make_naive_tensor_descriptor(make_tuple(N, Hi, Wi, C), + make_tuple(din_n_c_wos_strides[0], + din_n_c_wos_strides[2], + din_n_c_wos_strides[3], + din_n_c_wos_strides[1])); + + const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor( + in_n_hi_wi_c_grid_desc, + make_tuple(make_pass_through_transform(N), + make_pad_transform(Hi, InLeftPadH, InRightPadH), + make_pad_transform(Wi, InLeftPadW, InRightPadW), + make_pass_through_transform(C)), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{})); + + const auto in_n_ytilde_htilde_xtilde_wtilde_c_grid_desc = transform_tensor_descriptor( + in_n_hip_wip_c_grid_desc, + make_tuple(make_pass_through_transform(N), + make_embed_transform(make_tuple(YTilde, HTilde), + make_tuple(ConvDilationH, ConvStrideH)), + make_embed_transform(make_tuple(XTilde, WTilde), + make_tuple(ConvDilationW, ConvStrideW)), + make_pass_through_transform(C)), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), + make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{})); + + const auto in_n_htildeslice_wtildeslice_c_grid_desc = transform_tensor_descriptor( + in_n_ytilde_htilde_xtilde_wtilde_c_grid_desc, + make_tuple(make_pass_through_transform(N), + make_freeze_transform(i_ytilde), + make_slice_transform(HTilde, IHTildeSliceBegin, HTildeSlice), + make_freeze_transform(i_xtilde), + make_slice_transform(WTilde, IWTildeSliceBegin, WTildeSlice), + make_pass_through_transform(C)), + make_tuple(Sequence<0>{}, + Sequence<1>{}, + Sequence<2>{}, + Sequence<3>{}, + Sequence<4>{}, + Sequence<5>{}), + make_tuple(Sequence<0>{}, + Sequence<>{}, + Sequence<1>{}, + Sequence<>{}, + Sequence<2>{}, + Sequence<3>{})); + + const auto in_grid_desc_reducemraw = transform_tensor_descriptor( + in_n_htildeslice_wtildeslice_c_grid_desc, + make_tuple(make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice, C))), + make_tuple(Sequence<0, 1, 2, 3>{}), + make_tuple(Sequence<0>{})); + + const auto in_grid_desc_reducem = + transform_tensor_descriptor(in_grid_desc_reducemraw, + make_tuple(make_right_pad_transform(MRaw, MPad)), + make_tuple(Sequence<0>{}), + make_tuple(Sequence<0>{})); + + return make_tuple(out_grid_desc_reducem_reducek, in_grid_desc_reducem); + } + + using DoutDinGridDesc = decltype(Make2DGridDescriptor_Out_M_K_In_M({0, 0, 0, 0}, + {0, 0, 0, 0}, + {0, 0, 0, 0}, + {0, 0, 0, 0}, + {0, 0}, + {0, 0}, + {0, 0}, + {0, 0}, + {0, 0}, + {0, 0})); + + using DoutGridDesc_M_K = remove_cvref_t>; + using DinGridDesc_M = remove_cvref_t>; + + // FIXME + // for NHWC, the dim C is the fastest dimension, and is not reduced. + // Hence, it is in M dimension for reduction kernel. + static constexpr index_t OutSrcInDstVectorDim = 0; // 0: M, 1: K + + using PassThrough = tensor_operation::element_wise::PassThrough; + using Div = tensor_operation::element_wise::UnaryDivide; + + using gridwise_reduce = GridwiseReduction_mk_to_m_threadwise; + + struct Argument : public BaseArgument + { + Argument(const DOutDataType* p_dout, + DInDataType* p_din, + std::vector dout_n_c_wos_lengths, + std::vector din_n_c_wos_length, + std::vector dout_n_c_wos_strides, + std::vector din_n_c_wos_strides, + std::vector window_lengths, + std::vector window_strides, + std::vector window_dilations, + std::vector input_left_pads, + std::vector input_right_pads) + : p_dout_grid_{p_dout}, + p_din_grid_{p_din}, + dout_n_c_wos_lengths_{dout_n_c_wos_lengths}, + din_n_c_wos_length_{din_n_c_wos_length}, + dout_n_c_wos_strides_{dout_n_c_wos_strides}, + din_n_c_wos_strides_{din_n_c_wos_strides}, + num_reduce_{1}, + div_element_op_{window_lengths[0] * window_lengths[1]} + { + std::vector Tildes(NDimSpatial); + for(int i = 0; i < NDimSpatial; ++i) + { + int GcdStrideDilation = math::gcd(window_strides[i], window_dilations[i]); + Tildes[i] = window_strides[i] / GcdStrideDilation; + num_reduce_ *= Tildes[i]; + } + + for(index_t i_ytilde = 0; i_ytilde < Tildes[0]; ++i_ytilde) + { + for(index_t i_xtilde = 0; i_xtilde < Tildes[1]; ++i_xtilde) + { + const auto YDotSlice = + math::integer_divide_ceil(window_lengths[0] - i_ytilde, Tildes[0]); + const auto XDotSlice = + math::integer_divide_ceil(window_lengths[1] - i_xtilde, Tildes[1]); + + if(YDotSlice * XDotSlice <= 0) + { + continue; + } + + const auto dout_din_grid_desc = + Make2DGridDescriptor_Out_M_K_In_M(dout_n_c_wos_lengths, + din_n_c_wos_length, + dout_n_c_wos_strides, + din_n_c_wos_strides, + window_lengths, + window_strides, + window_dilations, + input_left_pads, + input_right_pads, + {i_ytilde, i_xtilde}); + + dout_grid_desc_m_k_container_.push_back(dout_din_grid_desc[I0]); + din_grid_desc_m_container_.push_back(dout_din_grid_desc[I1]); + } + } + } + + const DOutDataType* p_dout_grid_; + DInDataType* p_din_grid_; + std::vector dout_n_c_wos_lengths_; + std::vector din_n_c_wos_length_; + std::vector dout_n_c_wos_strides_; + std::vector din_n_c_wos_strides_; + + int num_reduce_; + std::vector dout_grid_desc_m_k_container_; + std::vector din_grid_desc_m_container_; + + Div div_element_op_; + }; + + struct Invoker : public BaseInvoker + { + float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{}) + { + float ave_time = 0; + + for(index_t i = 0; i < arg.num_reduce_; i++) + { + const auto kernel = kernel_reduce_threadwise; + + ck::index_t M = arg.dout_grid_desc_m_k_container_[i].GetLength(I0); + const index_t grid_size = (M / M_BlockTileSize); + + ave_time += launch_and_time_kernel(stream_config, + kernel, + dim3(grid_size), + dim3(BlockSize), + 0, + arg.dout_grid_desc_m_k_container_[i], + arg.din_grid_desc_m_container_[i], + PassThrough{}, + arg.div_element_op_, + float(1), + arg.p_dout_grid_, + nullptr, + float(0), + arg.p_din_grid_, + nullptr); + } + + return ave_time; + } + + float Run(const BaseArgument* p_arg, + const StreamConfig& stream_config = StreamConfig{}) override + { + return Run(*dynamic_cast(p_arg), stream_config); + } + }; + + static bool IsSupportedArgument(const Argument& arg) + { + constexpr index_t Rank = NDimSpatial + 2; + int doutFastestDim = -1; + int dinFastestDim = -1; + + for(int i = 0; i < Rank; ++i) + { + if(arg.dout_n_c_wos_strides_[i] == 1) + doutFastestDim = i; + if(arg.din_n_c_wos_strides_[i] == 1) + dinFastestDim = i; + } + if(InSrcOutDstVectorSize != 1 && (dinFastestDim != 1 || doutFastestDim != 1)) + { + return false; + } + if(doutFastestDim == -1 || dinFastestDim == -1) + { + if constexpr(InSrcOutDstVectorSize != 1) + return false; + } + else + { + if(arg.dout_n_c_wos_lengths_[doutFastestDim] % InSrcOutDstVectorSize != 0) + return false; + if(arg.din_n_c_wos_length_[dinFastestDim] % InSrcOutDstVectorSize != 0) + return false; + } + return true; + } + + bool IsSupportedArgument(const BaseArgument* p_arg) override + { + return IsSupportedArgument(*dynamic_cast(p_arg)); + } + + std::unique_ptr + MakeArgumentPointer(const void* p_dout, + void* p_din, + std::vector dout_n_c_wos_lengths, + std::vector din_n_c_wos_length, + std::vector dout_n_c_wos_strides, + std::vector din_n_c_wos_strides, + std::vector window_lengths, + std::vector window_strides, + std::vector window_dilations, + std::vector input_left_pads, + std::vector input_right_pads) override + { + constexpr index_t Rank = NDimSpatial + 2; + + if(dout_n_c_wos_strides.size() != Rank || din_n_c_wos_strides.size() != Rank || + dout_n_c_wos_lengths.size() != Rank || din_n_c_wos_length.size() != Rank) + { + throw std::runtime_error("dimension of [dout|din]_n_c_wos_strides or " + "[dout|din]_n_c_wos_lengths is not equal to Rank"); + } + + if(window_lengths.size() != NDimSpatial || window_strides.size() != NDimSpatial || + window_dilations.size() != NDimSpatial || input_left_pads.size() != NDimSpatial || + input_right_pads.size() != NDimSpatial) + { + throw std::runtime_error( + "dimension of [window_lengths, window_strides, window_dilations, input_left_pads, " + "input_right_pads] is not equal to Rank"); + } + return std::make_unique(static_cast(p_dout), + static_cast(p_din), + dout_n_c_wos_lengths, + din_n_c_wos_length, + dout_n_c_wos_strides, + din_n_c_wos_strides, + window_lengths, + window_strides, + window_dilations, + input_left_pads, + input_right_pads); + } + + std::unique_ptr MakeInvokerPointer() override + { + return std::make_unique(Invoker{}); + } + + std::string GetTypeString() const override + { + auto str = std::stringstream(); + + // clang-format off + str << "DeviceAvgPool2dBwd<" << BlockSize << ","; + str << "M_C" << MThreadClusterSize << "_S" << MThreadSliceSize << ","; + str << "K_C" << KThreadClusterSize << "_S" << KThreadSliceSize << ","; + str <<"InSrcOutDstVectorSize_" << InSrcOutDstVectorSize << ">"; + // clang-format on + + return str.str(); + } +}; + +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp index bf4a1c800f..8079b04b84 100644 --- a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp +++ b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp @@ -355,12 +355,39 @@ struct UnaryDivide __host__ __device__ void operator()(T& y, const T& x) const { static_assert(is_same::value || is_same::value || - is_same::value, + is_same::value || is_same::value, "Data type is not supported by this operation!"); y = x / type_convert(divider_); }; + template <> + __host__ __device__ void operator()(half_t& y, const half_t& x) const + { + float x_ = type_convert(x); + float divider_f_ = type_convert(divider_); + + y = type_convert(x_ / divider_f_); + }; + + template <> + __host__ __device__ void operator()(bhalf_t& y, const bhalf_t& x) const + { + float x_ = type_convert(x); + float divider_f_ = type_convert(divider_); + + y = type_convert(x_ / divider_f_); + }; + + template <> + __host__ __device__ void operator()(f8_t& y, const f8_t& x) const + { + float x_ = type_convert(x); + float divider_f_ = type_convert(divider_); + + y = type_convert(x_ / divider_f_); + }; + int32_t divider_ = 1; }; diff --git a/include/ck/utility/reduction_operator.hpp b/include/ck/utility/reduction_operator.hpp index fffd0ac49e..9f0a6b6854 100644 --- a/include/ck/utility/reduction_operator.hpp +++ b/include/ck/utility/reduction_operator.hpp @@ -52,12 +52,28 @@ struct Add __host__ __device__ inline constexpr void operator()(T& a, T b) const { static_assert(is_same::value || is_same::value || - is_same::value || is_same::value, + is_same::value || is_same::value, "The data type is not supported by the Add accumulator!"); a = a + b; } + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + a = type_convert(a_ + b_); + } + + __host__ __device__ inline constexpr void operator()(half_t& a, half_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + a = type_convert(a_ + b_); + } + __host__ __device__ inline constexpr void operator()(bhalf_t& a, bhalf_t b) const { float a_ = type_convert(a); @@ -112,12 +128,28 @@ struct Mul __host__ __device__ inline constexpr void operator()(T& a, T b) const { static_assert(is_same::value || is_same::value || - is_same::value || is_same::value, + is_same::value || is_same::value, "The data type is not supported by the Mul accumulator!"); a = a * b; } + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + a = type_convert(a_ * b_); + } + + __host__ __device__ inline constexpr void operator()(half_t& a, half_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + a = type_convert(a_ * b_); + } + __host__ __device__ inline constexpr void operator()(bhalf_t& a, bhalf_t b) const { float a_ = type_convert(a); @@ -137,6 +169,16 @@ struct Max float val = NumericLimits::Lowest(); return type_convert(val); } + if constexpr(is_same_v) + { + float val = NumericLimits::Lowest(); + return type_convert(val); + } + if constexpr(is_same_v) + { + float val = NumericLimits::Lowest(); + return type_convert(val); + } else { return NumericLimits::Lowest(); @@ -154,8 +196,7 @@ struct Max __host__ __device__ inline constexpr void operator()(T& a, T b) const { static_assert(is_same::value || is_same::value || - is_same::value || is_same::value || - is_same::value, + is_same::value || is_same::value, "The data type is not supported by the Max accumulator!"); if(a < b) @@ -171,12 +212,29 @@ struct Max a = b; } + __host__ __device__ inline constexpr void operator()(half_t& a, half_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ < b_) + a = b; + } + + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ < b_) + a = b; + } + template __host__ __device__ inline constexpr void operator()(T& a, T b, bool& changed) const { static_assert(is_same::value || is_same::value || - is_same::value || is_same::value || - is_same::value, + is_same::value || is_same::value, "The data type is not supported by the Max accumulator!"); if(a < b) @@ -197,6 +255,30 @@ struct Max changed = true; } } + + __host__ __device__ inline constexpr void operator()(half_t& a, half_t b, bool& changed) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ < b_) + { + a = b; + changed = true; + } + } + + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b, bool& changed) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ < b_) + { + a = b; + changed = true; + } + } }; struct Min @@ -209,6 +291,16 @@ struct Min float val = NumericLimits::Max(); return type_convert(val); } + else if constexpr(is_same_v) + { + float val = NumericLimits::Max(); + return type_convert(val); + } + else if constexpr(is_same_v) + { + float val = NumericLimits::Max(); + return type_convert(val); + } else { return NumericLimits::Max(); @@ -227,8 +319,7 @@ struct Min __host__ __device__ inline constexpr void operator()(T& a, T b) const { static_assert(is_same::value || is_same::value || - is_same::value || is_same::value || - is_same::value, + is_same::value || is_same::value, "The data type is not supported by the Min accumulator!"); if(a > b) @@ -244,6 +335,24 @@ struct Min a = b; } + __host__ __device__ inline constexpr void operator()(half_t& a, half_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ > b_) + a = b; + } + + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ > b_) + a = b; + } + template __host__ __device__ inline constexpr void operator()(T& a, T b, bool& changed) const { @@ -270,6 +379,30 @@ struct Min changed = true; } } + + __host__ __device__ inline constexpr void operator()(half_t& a, half_t b, bool& changed) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ > b_) + { + a = b; + changed = true; + } + } + + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b, bool& changed) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ > b_) + { + a = b; + changed = true; + } + } }; struct AMax @@ -299,6 +432,15 @@ struct AMax a = b; } + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ < b_) + a = b; + } + template __host__ __device__ inline constexpr void operator()(T& a, T b, bool& changed) const { @@ -313,6 +455,18 @@ struct AMax changed = true; } } + + __host__ __device__ inline constexpr void operator()(f8_t& a, f8_t b, bool& changed) const + { + float a_ = type_convert(a); + float b_ = type_convert(b); + + if(a_ < b_) + { + a = b; + changed = true; + } + } }; template @@ -352,7 +506,8 @@ struct InMemoryDataOperationSupportedOnDataType::value || is_same::value || is_same::value || is_same::value || - is_same::value || is_same::value; + is_same::value || is_same::value || + is_same::value; }; template @@ -361,7 +516,7 @@ struct InMemoryDataOperationSupportedOnDataType::value || is_same::value || is_same::value || is_same::value || - is_same::value; + is_same::value || is_same::value; }; } // namespace reduce diff --git a/library/include/ck/library/tensor_operation_instance/gpu/avg_pool2d_bwd.hpp b/library/include/ck/library/tensor_operation_instance/gpu/avg_pool2d_bwd.hpp new file mode 100644 index 0000000000..b2cb946ffe --- /dev/null +++ b/library/include/ck/library/tensor_operation_instance/gpu/avg_pool2d_bwd.hpp @@ -0,0 +1,80 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include "ck/tensor_operation/gpu/device/device_avgpool_bwd.hpp" +#include "ck/library/tensor_operation_instance/device_operation_instance_factory.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +#ifdef CK_ENABLE_BF16 +void add_device_avgpool_2D_bwd_nhwc_bf16_instances( + std::vector>>&); +#endif + +#ifdef CK_ENABLE_FP16 +void add_device_avgpool_2D_bwd_nhwc_f16_instances( + std::vector>>&); +#endif + +#ifdef CK_ENABLE_FP8 +void add_device_avgpool_2D_bwd_nhwc_f8_instances( + std::vector>>&); +#endif + +#ifdef CK_ENABLE_FP32 +void add_device_avgpool_2D_bwd_nhwc_f32_instances( + std::vector>>&); +#endif + +#ifdef CK_ENABLE_INT8 +void add_device_avgpool_2D_bwd_nhwc_int8_instances( + std::vector>>&); +#endif + +template +struct DeviceOperationInstanceFactory< + ck::tensor_operation::device:: + DeviceAvgPoolBwd<2, DOutDataType, DInDataType, InLayout, OutLayout>> +{ + using DeviceOp = DeviceAvgPoolBwd<2, DOutDataType, DInDataType, InLayout, OutLayout>; + + static auto GetInstances() + { + std::vector> op_ptrs; + if constexpr(is_same_v && is_same_v) + { +#ifdef CK_ENABLE_FP16 + if constexpr(is_same_v && is_same_v) + add_device_avgpool_2D_bwd_nhwc_f16_instances(op_ptrs); +#endif +#ifdef CK_ENABLE_BF16 + else if constexpr(is_same_v && is_same_v) + add_device_avgpool_2D_bwd_nhwc_bf16_instances(op_ptrs); +#endif +#ifdef CK_ENABLE_FP32 + else if constexpr(is_same_v && is_same_v) + add_device_avgpool_2D_bwd_nhwc_f32_instances(op_ptrs); +#endif +#ifdef CK_ENABLE_FP8 + else if constexpr(is_same_v && is_same_v) + add_device_avgpool_2D_bwd_nhwc_f8_instances(op_ptrs); +#endif +#ifdef CK_ENABLE_INT8 + else if constexpr(is_same_v && is_same_v) + add_device_avgpool_2D_bwd_nhwc_int8_instances(op_ptrs); +#endif + } + + return op_ptrs; + } +}; + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/include/ck/library/tensor_operation_instance/gpu/max_pool_bwd.hpp b/library/include/ck/library/tensor_operation_instance/gpu/max_pool_bwd.hpp index 63ea4f2891..1dc3544ecb 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/max_pool_bwd.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/max_pool_bwd.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -23,6 +23,11 @@ void add_device_maxpool_bwd_bf16_instances( void add_device_maxpool_bwd_f32_instances( std::vector>>&); #endif +#ifdef CK_ENABLE_INT8 +void add_device_maxpool_bwd_int8_instances( + std::vector>>&); +#endif + template struct DeviceOperationInstanceFactory< ck::tensor_operation::device::DeviceMaxPoolBwd> @@ -32,6 +37,7 @@ struct DeviceOperationInstanceFactory< static auto GetInstances() { std::vector> op_ptrs; + #ifdef CK_ENABLE_FP16 if constexpr(is_same_v && is_same_v && is_same_v) @@ -47,6 +53,11 @@ struct DeviceOperationInstanceFactory< is_same_v) add_device_maxpool_bwd_f32_instances(op_ptrs); #endif +#ifdef CK_ENABLE_INT8 + else if constexpr(is_same_v && is_same_v && + is_same_v) + add_device_maxpool_bwd_int8_instances(op_ptrs); +#endif return op_ptrs; } diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/CMakeLists.txt new file mode 100644 index 0000000000..ad69023465 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/CMakeLists.txt @@ -0,0 +1,8 @@ +set(DEVICE_AVGPOOL_2D_BWD_INSTANCES) +list(APPEND DEVICE_AVGPOOL_2D_BWD_INSTANCES device_avg_pool2d_bwd_nhwc_bf16_instance.cpp + device_avg_pool2d_bwd_nhwc_f16_instance.cpp + device_avg_pool2d_bwd_nhwc_f32_instance.cpp + device_avg_pool2d_bwd_nhwc_f8_instance.cpp + device_avg_pool2d_bwd_nhwc_int8_instance.cpp) +add_instance_library(device_avg_pool2d_bwd_instance ${DEVICE_AVGPOOL_2D_BWD_INSTANCES}) + diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_bf16_instance.cpp b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_bf16_instance.cpp new file mode 100644 index 0000000000..05e7c3ede3 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_bf16_instance.cpp @@ -0,0 +1,21 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_avg_pool2d_bwd_nhwc_instance_common.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_avgpool_2D_bwd_nhwc_bf16_instances( + std::vector>>& instances) +{ + add_device_operation_instances(instances, + device_avgpool_2D_bwd_nhwc_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f16_instance.cpp b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f16_instance.cpp new file mode 100644 index 0000000000..ce865423d5 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f16_instance.cpp @@ -0,0 +1,21 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_avg_pool2d_bwd_nhwc_instance_common.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_avgpool_2D_bwd_nhwc_f16_instances( + std::vector>>& instances) +{ + add_device_operation_instances(instances, + device_avgpool_2D_bwd_nhwc_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f32_instance.cpp b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f32_instance.cpp new file mode 100644 index 0000000000..7c348cea98 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f32_instance.cpp @@ -0,0 +1,21 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_avg_pool2d_bwd_nhwc_instance_common.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_avgpool_2D_bwd_nhwc_f32_instances( + std::vector>>& instances) +{ + add_device_operation_instances(instances, + device_avgpool_2D_bwd_nhwc_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f8_instance.cpp b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f8_instance.cpp new file mode 100644 index 0000000000..34b72471ab --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_f8_instance.cpp @@ -0,0 +1,20 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_avg_pool2d_bwd_nhwc_instance_common.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_avgpool_2D_bwd_nhwc_f8_instances( + std::vector>>& instances) +{ + add_device_operation_instances(instances, device_avgpool_2D_bwd_nhwc_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_instance_common.hpp b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_instance_common.hpp new file mode 100644 index 0000000000..aca6b84653 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_instance_common.hpp @@ -0,0 +1,38 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include "ck/ck.hpp" +#include "ck/tensor_operation/gpu/device/impl/device_avgpool2d_bwd_nhwc_nhwc.hpp" +#include "ck/utility/data_type.hpp" + +#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +using F16 = ck::half_t; +using BF16 = ck::bhalf_t; +using F8 = ck::f8_t; +using I8 = int8_t; +using I32 = int32_t; +using F32 = float; +using NHWC = ck::tensor_layout::convolution::NHWC; + +template +using device_avgpool_2D_bwd_nhwc_instances = std::tuple< + // clang-format off + DeviceAvgPool2dBwd_NHWC_NHWC, + DeviceAvgPool2dBwd_NHWC_NHWC, + DeviceAvgPool2dBwd_NHWC_NHWC, + DeviceAvgPool2dBwd_NHWC_NHWC, + DeviceAvgPool2dBwd_NHWC_NHWC + // clang-format on + >; +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_int8_instance.cpp b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_int8_instance.cpp new file mode 100644 index 0000000000..a55fe7f5f2 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/avg_pool2d_bwd/device_avg_pool2d_bwd_nhwc_int8_instance.cpp @@ -0,0 +1,20 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_avg_pool2d_bwd_nhwc_instance_common.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_avgpool_2D_bwd_nhwc_int8_instances( + std::vector>>& instances) +{ + add_device_operation_instances(instances, device_avgpool_2D_bwd_nhwc_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/max_pool_bwd/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/max_pool_bwd/CMakeLists.txt index d12864398e..a2315175d8 100644 --- a/library/src/tensor_operation_instance/gpu/max_pool_bwd/CMakeLists.txt +++ b/library/src/tensor_operation_instance/gpu/max_pool_bwd/CMakeLists.txt @@ -1,5 +1,6 @@ set(DEVICE_MAXPOOL_BWD_INSTANCES) list(APPEND DEVICE_MAXPOOL_BWD_INSTANCES device_max_pool_bwd_f16_instance.cpp device_max_pool_bwd_bf16_instance.cpp - device_max_pool_bwd_f32_instance.cpp) + device_max_pool_bwd_f32_instance.cpp + device_max_pool_bwd_int8_instance.cpp) add_instance_library(device_max_pool_bwd_instance ${DEVICE_MAXPOOL_BWD_INSTANCES}) diff --git a/library/src/tensor_operation_instance/gpu/max_pool_bwd/device_max_pool_bwd_int8_instance.cpp b/library/src/tensor_operation_instance/gpu/max_pool_bwd/device_max_pool_bwd_int8_instance.cpp new file mode 100644 index 0000000000..5971a9de06 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/max_pool_bwd/device_max_pool_bwd_int8_instance.cpp @@ -0,0 +1,20 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "max_pool_bwd_instance_common.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_maxpool_bwd_int8_instances( + std::vector>>& instances) +{ + add_device_operation_instances(instances, device_maxpool_bwd_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/max_pool_bwd/max_pool_bwd_instance_common.hpp b/library/src/tensor_operation_instance/gpu/max_pool_bwd/max_pool_bwd_instance_common.hpp index 0bba106ee2..f19b69a4ae 100644 --- a/library/src/tensor_operation_instance/gpu/max_pool_bwd/max_pool_bwd_instance_common.hpp +++ b/library/src/tensor_operation_instance/gpu/max_pool_bwd/max_pool_bwd_instance_common.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -17,6 +17,8 @@ namespace instance { using I32 = int32_t; using F16 = ck::half_t; using BF16 = ck::bhalf_t; +using I8 = int8_t; +using F8 = ck::f8_t; using F32 = float; template diff --git a/profiler/include/profiler/data_type_enum.hpp b/profiler/include/profiler/data_type_enum.hpp index c046c7fabb..bbfb04df7d 100644 --- a/profiler/include/profiler/data_type_enum.hpp +++ b/profiler/include/profiler/data_type_enum.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -14,6 +14,7 @@ enum struct DataTypeEnum Int8x4 = 4, BFloat16 = 5, Double = 6, + Float8 = 7, Unknown = 100, }; diff --git a/profiler/include/profiler/profile_avg_pool2d_bwd_impl.hpp b/profiler/include/profiler/profile_avg_pool2d_bwd_impl.hpp new file mode 100644 index 0000000000..caf24f016a --- /dev/null +++ b/profiler/include/profiler/profile_avg_pool2d_bwd_impl.hpp @@ -0,0 +1,255 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include + +#include "ck/ck.hpp" +#include "ck/library/tensor_operation_instance/gpu/avg_pool2d_bwd.hpp" +#include "ck/library/utility/check_err.hpp" +#include "ck/library/utility/device_memory.hpp" +#include "ck/library/utility/host_tensor.hpp" +#include "ck/library/utility/host_tensor_generator.hpp" +#include "ck/library/utility/literals.hpp" +#include "ck/library/reference_tensor_operation/cpu/reference_avgpool_bwd.hpp" + +namespace ck { +namespace profiler { + +template +std::vector f_tensor_strides_nchw( + ck::index_t N, ck::index_t C, ck::index_t H, ck::index_t W, TensorLayout layout) +{ + using namespace ck::literals; + (void)N; + if constexpr(ck::is_same::value) + return {C * H * W, 1_uz, W * C, C}; + else + throw std::runtime_error("not supported yet"); +}; + +template +bool profile_avg_pool2d_bwd_impl(int do_verification, + int init_method, + bool do_log, + bool time_kernel, + std::vector in_length, + std::vector window_spatial_lengths, + std::vector window_strides, + std::vector window_dilations, + std::vector input_left_pads, + std::vector input_right_pads) +{ + constexpr index_t InOutRank = 4; + constexpr index_t WindowRank = 2; + + if(in_length.size() != InOutRank || window_spatial_lengths.size() != WindowRank || + window_strides.size() != WindowRank || window_dilations.size() != WindowRank || + input_left_pads.size() != WindowRank || input_right_pads.size() != WindowRank) + { + std::cout << "Parameter is incorrect" << std::endl; + return false; + } + + std::vector out_length(InOutRank); + + const int N = in_length[0]; + const int C = in_length[1]; + + out_length[0] = N; + out_length[1] = C; + + // Calculate Ho, Wo + for(unsigned i = 2; i < InOutRank; ++i) + { + const int idx = i - 2; + auto pad1 = input_left_pads[idx]; + auto pad2 = input_right_pads[idx]; + auto windows_size = window_spatial_lengths[idx]; + auto windows_stride = window_strides[idx]; + auto windows_dilation = window_dilations[idx]; + auto eff = (windows_size - 1) * windows_dilation + 1; + out_length[i] = (in_length[i] + pad1 + pad2 - eff) / windows_stride + 1; + } + + const int Hi = in_length[2]; + const int Wi = in_length[3]; + const int Ho = out_length[2]; + const int Wo = out_length[3]; + + auto f_host_tensor_descriptor = + [](std::size_t N_, std::size_t C_, std::size_t H, std::size_t W) { + using namespace ck::literals; + + return HostTensorDescriptor({N_, C_, H, W}, {C_ * H * W, 1_uz, W * C_, C_}); + }; + + Tensor out_n_c_ho_wo_host(f_host_tensor_descriptor(N, C, Ho, Wo)); + Tensor in_n_c_hi_wi_device(f_host_tensor_descriptor(N, C, Hi, Wi)); + Tensor in_n_c_hi_wi_host(f_host_tensor_descriptor(N, C, Hi, Wi)); + + switch(init_method) + { + case 0: { + out_n_c_ho_wo_host.GenerateTensorValue(GeneratorTensor_1{}); + break; + } + case 1: { + out_n_c_ho_wo_host.GenerateTensorValue(GeneratorTensor_2{-5, 5}); + break; + } + default: { + out_n_c_ho_wo_host.GenerateTensorValue(GeneratorTensor_3{-0.5, 0.5}); + } + } + + DeviceMem dout_device_buf(sizeof(DOutDataType) * + out_n_c_ho_wo_host.mDesc.GetElementSpaceSize()); + DeviceMem din_device_buf(sizeof(DInDataType) * in_n_c_hi_wi_device.mDesc.GetElementSpaceSize()); + + dout_device_buf.ToDevice(out_n_c_ho_wo_host.mData.data()); + + using DeviceOp = ck::tensor_operation::device:: + DeviceAvgPoolBwd<2, DOutDataType, DInDataType, DOutLayout, DInLayout>; + + // get device op instances + const auto instance_ptrs = + ck::tensor_operation::device::instance::DeviceOperationInstanceFactory< + DeviceOp>::GetInstances(); + + std::cout << "found " << instance_ptrs.size() << " instances" << std::endl; + + std::string best_instance_name; + float best_avg_time = std::numeric_limits::max(); + float best_gb_per_sec = 0; + + if(do_verification) + { + using ReferencePoolingBwdInstance = + ck::tensor_operation::host::ReferenceAvgPoolBwd<2, DInDataType, DOutDataType>; + + ReferencePoolingBwdInstance ref_pooling_bwd; + auto ref_pooling_bwd_argument = ref_pooling_bwd.MakeArgument(in_n_c_hi_wi_host, + out_n_c_ho_wo_host, + window_spatial_lengths, + window_strides, + window_dilations, + input_left_pads, + input_right_pads); + + auto ref_invoker = ref_pooling_bwd.MakeInvoker(); + ref_invoker.Run(ref_pooling_bwd_argument); + } + + int num_kernel = 0; + bool pass = true; + bool instance_found = false; + for(auto& inst_ptr : instance_ptrs) + { + auto argument_ptr = inst_ptr->MakeArgumentPointer( + static_cast(dout_device_buf.GetDeviceBuffer()), + static_cast(din_device_buf.GetDeviceBuffer()), + {N, C, Ho, Wo}, + {N, C, Hi, Wi}, + f_tensor_strides_nchw(N, C, Ho, Wo, DOutLayout{}), + f_tensor_strides_nchw(N, C, Hi, Wi, DInLayout{}), + window_spatial_lengths, + window_strides, + window_dilations, + input_left_pads, + input_right_pads); + + if(inst_ptr->IsSupportedArgument(argument_ptr.get())) + { + ++num_kernel; + instance_found = true; + } + else + { + if(time_kernel) + { + std::cout << inst_ptr->GetTypeString() << " skipped due to unsupported argument: "; + LogRange(std::cout << "doutput lengths = ", out_length, ", ") << std::endl; + } + + continue; + } + + din_device_buf.SetZero(); + + auto invoker_ptr = inst_ptr->MakeInvokerPointer(); + float avg_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, time_kernel}); + + std::size_t num_bytes = out_n_c_ho_wo_host.mDesc.GetElementSize() * sizeof(DOutDataType) + + in_n_c_hi_wi_device.mDesc.GetElementSize() * sizeof(DInDataType); + + float gb_per_sec = num_bytes / 1.E6 / avg_time; + + if(time_kernel) + { + std::cout << "Perf: " << std::setw(10) << avg_time << " ms, " << gb_per_sec << " GB/s, " + << inst_ptr->GetTypeString() << std::endl; + } + + if(avg_time < best_avg_time) + { + best_instance_name = inst_ptr->GetTypeString(); + best_avg_time = avg_time; + best_gb_per_sec = gb_per_sec; + } + + if(do_verification) + { + din_device_buf.FromDevice(in_n_c_hi_wi_device.mData.data()); + bool local_pass = ck::utils::check_err(in_n_c_hi_wi_device.mData, + in_n_c_hi_wi_host.mData, + "Error: Incorrect results", + 1e-3, + 1e-3); + + if(do_log) + { + LogRangeAsType( + std::cout << "in_n_c_hi_wi_device: ", in_n_c_hi_wi_device.mData, ",") + << std::endl; + + LogRangeAsType( + std::cout << "in_n_c_hi_wi_host: ", in_n_c_hi_wi_host.mData, ",") + << std::endl; + } + + if(!local_pass) + { + std::cout << inst_ptr->GetTypeString() << " failed verification: "; + LogRange(std::cout << "doutput lengths = [", out_length, ", ") << "]." << std::endl; + pass &= local_pass; + } + else + { + if(time_kernel) + { + std::cout << "pass" << std::endl; + } + } + } + } + + if(time_kernel) + { + LogRange(std::cout << "length = ", out_length, ",") << std::endl; + std::cout << "best perf = " << best_avg_time << " ms, " << best_gb_per_sec << " GB/s, " + << best_instance_name << std::endl; + } + + if(num_kernel == 0) + { + std::cout << "Error: No kernel is applicable" << std::endl; + return false; + } + + return pass && instance_found; +} + +} // namespace profiler +} // namespace ck diff --git a/profiler/include/profiler/profile_max_pool2d_bwd_impl.hpp b/profiler/include/profiler/profile_max_pool2d_bwd_impl.hpp new file mode 100644 index 0000000000..7a712f21f2 --- /dev/null +++ b/profiler/include/profiler/profile_max_pool2d_bwd_impl.hpp @@ -0,0 +1,295 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include + +#include "ck/ck.hpp" +#include "ck/library/tensor_operation_instance/gpu/pool3d_fwd.hpp" +#include "ck/library/tensor_operation_instance/gpu/max_pool_bwd.hpp" +#include "ck/library/utility/check_err.hpp" +#include "ck/library/utility/device_memory.hpp" +#include "ck/library/utility/host_tensor.hpp" +#include "ck/library/utility/host_tensor_generator.hpp" +#include "ck/library/utility/literals.hpp" +#include "ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp" +#include "ck/library/reference_tensor_operation/cpu/reference_maxpool_bwd.hpp" + +namespace ck { +namespace profiler { + +template +bool profile_max_pool2d_bwd_impl(int do_verification, + int init_method, + bool do_log, + bool time_kernel, + std::vector in_length, + std::vector window_spatial_lengths, + std::vector window_strides, + std::vector window_dilations, + std::vector input_left_pads, + std::vector input_right_pads) +{ + // AtomicAdd only support f32 for now. ComputeDataType must be float32 + using ComputeDataType = float; + + constexpr index_t InOutRank = 4; + constexpr index_t WindowRank = 2; + + using PassThrough = ck::tensor_operation::element_wise::PassThrough; + + if(in_length.size() != InOutRank || window_spatial_lengths.size() != WindowRank || + window_strides.size() != WindowRank || window_dilations.size() != WindowRank || + input_left_pads.size() != WindowRank || input_right_pads.size() != WindowRank) + { + std::cout << "Parameter is incorrect" << std::endl; + return false; + } + + std::vector out_length(InOutRank); + + int N = in_length[0]; + int C = in_length[1]; + + out_length[0] = N; + out_length[1] = C; + + // Calculate Ho, Wo + for(unsigned i = 2; i < InOutRank; ++i) + { + const int idx = i - 2; + auto pad1 = input_left_pads[idx]; + auto pad2 = input_right_pads[idx]; + auto windows_size = window_spatial_lengths[idx]; + auto windows_stride = window_strides[idx]; + auto windows_dilation = window_dilations[idx]; + auto eff = (windows_size - 1) * windows_dilation + 1; + out_length[i] = (in_length[i] + pad1 + pad2 - eff) / windows_stride + 1; + } + + int Hi = in_length[2]; + int Wi = in_length[3]; + int Ho = out_length[2]; + int Wo = out_length[3]; + + auto f_host_tensor_descriptor = + [](std::size_t N_, std::size_t C_, std::size_t H, std::size_t W) { + using namespace ck::literals; + + return HostTensorDescriptor({N_, C_, H, W}, {C_ * H * W, 1_uz, W * C_, C_}); + }; + + Tensor in_n_c_hi_wi(f_host_tensor_descriptor(N, C, Hi, Wi)); + Tensor out_n_c_ho_wo(f_host_tensor_descriptor(N, C, Ho, Wo)); + Tensor out_indices_n_c_ho_wo(f_host_tensor_descriptor(N, C, Ho, Wo)); + Tensor dout_n_c_ho_wo(f_host_tensor_descriptor(N, C, Ho, Wo)); + Tensor din_n_c_hi_wi_host(f_host_tensor_descriptor(N, C, Hi, Wi)); + + Tensor din_n_c_hi_wi_device(f_host_tensor_descriptor(N, C, Hi, Wi)); + + switch(init_method) + { + case 0: { + in_n_c_hi_wi.GenerateTensorValue(GeneratorTensor_1{}); + dout_n_c_ho_wo.GenerateTensorValue(GeneratorTensor_1{}); + break; + } + case 1: { + in_n_c_hi_wi.GenerateTensorValue(GeneratorTensor_2{-5, 5}); + dout_n_c_ho_wo.GenerateTensorValue(GeneratorTensor_2{-5, 5}); + break; + } + default: { + in_n_c_hi_wi.GenerateTensorValue(GeneratorTensor_3{-0.5, 0.5}); + dout_n_c_ho_wo.GenerateTensorValue(GeneratorTensor_3{-0.5, 0.5}); + } + } + + DeviceMem indices_device_buf(sizeof(IndexDataType) * + out_indices_n_c_ho_wo.mDesc.GetElementSpaceSize()); + DeviceMem dout_device_buf(sizeof(DOutDataType) * dout_n_c_ho_wo.mDesc.GetElementSpaceSize()); + DeviceMem din_device_buf(sizeof(DInDataType) * + din_n_c_hi_wi_device.mDesc.GetElementSpaceSize()); + + // Generate index data from forwarding + { + using ReferencePoolingFwdInstance = + ck::tensor_operation::host::ReferencePoolingFwd; + + ReferencePoolingFwdInstance ref_pooling_fwd; + auto ref_pooling_fwd_argument = ref_pooling_fwd.MakeArgument(in_n_c_hi_wi, + out_n_c_ho_wo, + out_indices_n_c_ho_wo, + window_spatial_lengths, + window_strides, + window_dilations, + input_left_pads, + input_right_pads); + auto ref_pooling_fwd_invoker = ref_pooling_fwd.MakeInvoker(); + ref_pooling_fwd_invoker.Run(ref_pooling_fwd_argument); + } + + indices_device_buf.ToDevice(out_indices_n_c_ho_wo.mData.data()); + dout_device_buf.ToDevice(dout_n_c_ho_wo.mData.data()); + + using DeviceOp = + ck::tensor_operation::device::DeviceMaxPoolBwd; + + // get device op instances + const auto instance_ptrs = + ck::tensor_operation::device::instance::DeviceOperationInstanceFactory< + DeviceOp>::GetInstances(); + + std::cout << "found " << instance_ptrs.size() << " instances" << std::endl; + + std::string best_instance_name; + float best_avg_time = std::numeric_limits::max(); + float best_gb_per_sec = 0; + + if(do_verification) + { + using ReferencePoolingBwdInstance = + ck::tensor_operation::host::ReferenceMaxPoolBwd; + + ReferencePoolingBwdInstance ref_pooling_bwd; + auto ref_pooling_bwd_argument = ref_pooling_bwd.MakeArgument( + dout_n_c_ho_wo, out_indices_n_c_ho_wo, din_n_c_hi_wi_host, PassThrough{}); + auto ref_invoker = ref_pooling_bwd.MakeInvoker(); + ref_invoker.Run(ref_pooling_bwd_argument); + } + + int num_kernel = 0; + + bool pass = true; + bool instance_found = false; + + for(auto& inst_ptr : instance_ptrs) + { + auto argument_ptr = inst_ptr->MakeArgumentPointer( + static_cast(dout_device_buf.GetDeviceBuffer()), + static_cast(indices_device_buf.GetDeviceBuffer()), + static_cast(din_device_buf.GetDeviceBuffer()), + dout_n_c_ho_wo.mDesc.GetElementSpaceSize(), + din_n_c_hi_wi_device.mDesc.GetElementSpaceSize(), + window_spatial_lengths, + window_strides, + window_dilations); + + if(inst_ptr->IsSupportedArgument(argument_ptr.get())) + { + ++num_kernel; + instance_found = true; + } + else + { + if(time_kernel) + { + std::cout << inst_ptr->GetTypeString() << " skipped due to unsupported argument: "; + LogRange(std::cout << "doutput lengths = ", out_length, ", ") << std::endl; + } + + continue; + } + + size_t workspace_sz = inst_ptr->GetWorkSpaceSize(argument_ptr.get()); + DeviceMem workspace_device_buf(workspace_sz); + inst_ptr->SetWorkSpacePointer(argument_ptr.get(), workspace_device_buf.GetDeviceBuffer()); + + auto invoker_ptr = inst_ptr->MakeInvokerPointer(); + float avg_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, time_kernel}); + + std::size_t num_bytes = + dout_n_c_ho_wo.mDesc.GetElementSize() * sizeof(DOutDataType) + + out_indices_n_c_ho_wo.mDesc.GetElementSize() * sizeof(IndexDataType) + + din_n_c_hi_wi_device.mDesc.GetElementSize() * sizeof(DInDataType); + + float gb_per_sec = num_bytes / 1.E6 / avg_time; + + if(time_kernel) + std::cout << "Perf: " << std::setw(10) << avg_time << " ms, " << gb_per_sec << " GB/s, " + << inst_ptr->GetTypeString() << std::endl; + + if(avg_time < best_avg_time) + { + best_instance_name = inst_ptr->GetTypeString(); + best_avg_time = avg_time; + best_gb_per_sec = gb_per_sec; + } + + if(do_verification) + { + din_device_buf.FromDevice(din_n_c_hi_wi_device.mData.data()); + + bool local_pass = ck::utils::check_err(din_n_c_hi_wi_device.mData, + din_n_c_hi_wi_host.mData, + "Error: Incorrect results", + 1e-3, + 1e-3); + + if(do_log) + { + LogRangeAsType( + std::cout << "out_indices_n_c_ho_wo: ", out_indices_n_c_ho_wo.mData, ",") + << std::endl; + + LogRangeAsType( + std::cout << "din_n_c_hi_wi_device: ", din_n_c_hi_wi_device.mData, ",") + << std::endl; + + LogRangeAsType( + std::cout << "din_n_c_hi_wi_host: ", din_n_c_hi_wi_host.mData, ",") + << std::endl; + } + + if(!local_pass) + { + std::cout << inst_ptr->GetTypeString() << " failed verification: "; + LogRange(std::cout << "doutput lengths = [", out_length, ", ") << "]." << std::endl; + pass &= local_pass; + } + else + { + if(time_kernel) + { + std::cout << "pass" << std::endl; + } + } + } + } + + if(time_kernel) + { + LogRange(std::cout << "length = ", out_length, ",") << std::endl; + std::cout << "best perf = " << best_avg_time << " ms, " << best_gb_per_sec << " GB/s, " + << best_instance_name << std::endl; + } + + if(num_kernel == 0) + { + std::cout << "Error: No kernel is applicable" << std::endl; + return false; + } + + return pass && instance_found; +} + +} // namespace profiler +} // namespace ck diff --git a/profiler/src/CMakeLists.txt b/profiler/src/CMakeLists.txt old mode 100755 new mode 100644 index 554808cac5..8d91c2b5e0 --- a/profiler/src/CMakeLists.txt +++ b/profiler/src/CMakeLists.txt @@ -13,6 +13,8 @@ set(PROFILER_SOURCES profile_max_pool3d_fwd.cpp profile_avg_pool3d_bwd.cpp profile_max_pool3d_bwd.cpp + profile_avg_pool2d_bwd.cpp + profile_max_pool2d_bwd.cpp profile_softmax.cpp profile_batchnorm_fwd.cpp profile_batchnorm_bwd.cpp @@ -101,6 +103,7 @@ target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_reduce_instance) target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batchnorm_instance) target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_pool2d_fwd_instance) target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_pool3d_fwd_instance) +target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_avg_pool2d_bwd_instance) target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_avg_pool3d_bwd_instance) target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_max_pool_bwd_instance) target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_image_to_column_instance) diff --git a/profiler/src/profile_avg_pool2d_bwd.cpp b/profiler/src/profile_avg_pool2d_bwd.cpp new file mode 100644 index 0000000000..61d0413d43 --- /dev/null +++ b/profiler/src/profile_avg_pool2d_bwd.cpp @@ -0,0 +1,188 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include +#include +#include + +#include "profiler/data_type_enum.hpp" +#include "profiler/profile_avg_pool2d_bwd_impl.hpp" +#include "profiler_operation_registry.hpp" + +using ck::index_t; + +struct maxPoolbwdArgParser +{ + std::unordered_map> long_opts = {{"length", {}}, + {"wsize", {}}, + {"wstride", {}}, + {"wdilation", {}}, + {"pad1", {}}, + {"pad2", {}}}; + + bool parse_opt(int argc, char* argv[], const std::string& key, int i) + { + if(std::string("--") + key == argv[i]) + { + int pos = i; + while(++i < argc && argv[i][0] != '-') {} + int end = i; + for(int j = pos + 1; j < end; j++) + { + long_opts[key].push_back(std::stoi(argv[j])); + } + return true; + } + return false; + } + + void operator()(int argc, char* argv[]) + { + for(auto& kv : long_opts) + { + for(int i = 1; i < argc; i++) + { + if(parse_opt(argc, argv, kv.first, i)) + break; + } + } + } +}; + +void print_help_avg_pool2d_bwd() +{ + std::cout << "arg1: data type (0: fp16; 1: fp32; 3: int8; 5: bf16, 7: Float8)\n" + << "arg2: verification (0: no; 1: yes)\n" + << "arg3: initialization (0: no init; 1: integer value; 2: decimal value)\n" + << "arg4: print tensor value (0: no; 1: yes)\n" + << "arg5: time kernel (0=no, 1=yes)\n" + << "--length: input tensor length for NCHW(e.g, --length 2 32 30 30) \n" + << "--wsize: window size for YX (e.g, --wsize 2 2) \n" + << "--wstride: window stride for HW (e.g, --wstride 2 2) \n" + << "--wdilation: window dilation for HW (e.g, --wdilation 1 1) \n" + << "--pad1: left side of padding in HW (e.g, --pad1 1 1) \n" + << "--pad2: right side of padding in HW (e.g, --pad2 1 1) \n" + << "eg: ckProfiler avg_pool2d_bwd 0 1 2 0 --length 2 32 30 30 --wsize 2 2 " + "--wstride 2 2 --wdilation 1 1 --pad1 1 1 --pad2 1 1" + << std::endl; +} + +int profile_avg_pool2d_bwd(int argc, char* argv[]) +{ + ck::DataTypeEnum data_type = ck::DataTypeEnum::Float8; + bool do_verification = true; + int init_method = 2; + bool do_log = false; + bool time_kernel = true; + + std::vector in_length = {2, 32, 30, 30}; + std::vector wsize = {2, 2}; + std::vector wstride = {2, 2}; + std::vector wdilation = {1, 1}; + std::vector pad1 = {1, 1}; + std::vector pad2 = {1, 1}; + + if(argc != 2 && argc != 33) + { + print_help_avg_pool2d_bwd(); + return 0; + } + else if(argc == 33) + { + data_type = static_cast(std::stoi(argv[2])); + do_verification = std::stoi(argv[3]); + init_method = std::stoi(argv[4]); + do_log = std::stoi(argv[5]); + time_kernel = std::stoi(argv[6]); + + maxPoolbwdArgParser arg_parser; + arg_parser(argc, argv); + in_length = arg_parser.long_opts["length"]; + wsize = arg_parser.long_opts["wsize"]; + wstride = arg_parser.long_opts["wstride"]; + wdilation = arg_parser.long_opts["wdilation"]; + pad1 = arg_parser.long_opts["pad1"]; + pad2 = arg_parser.long_opts["pad2"]; + } + + using F16 = ck::half_t; + using BF16 = ck::bhalf_t; + using F8 = ck::f8_t; + using F32 = float; + using I8 = int8_t; + using NHWC = ck::tensor_layout::convolution::NHWC; + + if(data_type == ck::DataTypeEnum::Half) + { + ck::profiler::profile_avg_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::BFloat16) + { + ck::profiler::profile_avg_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::Float) + { + ck::profiler::profile_avg_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::Float8) + { + ck::profiler::profile_avg_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::Int8) + { + ck::profiler::profile_avg_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else + { + throw std::runtime_error("not implemented yet"); + } + + return 0; +} + +REGISTER_PROFILER_OPERATION("avg_pool2d_bwd", "avg_pool2d bwd", profile_avg_pool2d_bwd); diff --git a/profiler/src/profile_max_pool2d_bwd.cpp b/profiler/src/profile_max_pool2d_bwd.cpp new file mode 100644 index 0000000000..26e84c880e --- /dev/null +++ b/profiler/src/profile_max_pool2d_bwd.cpp @@ -0,0 +1,178 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include +#include +#include + +#include "profiler/data_type_enum.hpp" +#include "profiler/profile_max_pool2d_bwd_impl.hpp" +#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" +#include "profiler_operation_registry.hpp" + +using ck::index_t; + +struct maxPoolbwdArgParser +{ + std::unordered_map> long_opts = {{"length", {}}, + {"wsize", {}}, + {"wstride", {}}, + {"wdilation", {}}, + {"pad1", {}}, + {"pad2", {}}}; + + bool parse_opt(int argc, char* argv[], const std::string& key, int i) + { + if(std::string("--") + key == argv[i]) + { + int pos = i; + while(++i < argc && argv[i][0] != '-') {} + int end = i; + for(int j = pos + 1; j < end; j++) + { + long_opts[key].push_back(std::stoi(argv[j])); + } + return true; + } + return false; + } + + void operator()(int argc, char* argv[]) + { + for(auto& kv : long_opts) + { + for(int i = 1; i < argc; i++) + { + if(parse_opt(argc, argv, kv.first, i)) + break; + } + } + } +}; + +void print_help_max_pool2d_bwd() +{ + std::cout << "arg1: data type (0: fp16; 1: fp32; 3: int8; 5: bf16)\n" + << "arg2: verification (0: no; 1: yes)\n" + << "arg3: initialization (0: no init; 1: integer value; 2: decimal value)\n" + << "arg4: print tensor value (0: no; 1: yes)\n" + << "arg5: time kernel (0=no, 1=yes)\n" + << "--length: input tensor length for NCHW(e.g, --length 2 32 30 30) \n" + << "--wsize: window size for YX (e.g, --wsize 2 2) \n" + << "--wstride: window stride for HW (e.g, --wstride 2 2) \n" + << "--wdilation: window dilation for HW (e.g, --wdilation 1 1) \n" + << "--pad1: left side of padding in HW (e.g, --pad1 1 1) \n" + << "--pad2: right side of padding in HW (e.g, --pad2 1 1) \n" + << "eg: ckProfiler max_pool2d_bwd 0 1 2 0 --length 2 32 30 30 --wsize 2 2 " + "--wstride 2 2 --wdilation 1 1 --pad1 1 1 --pad2 1 1" + << std::endl; +} + +int profile_max_pool2d_bwd(int argc, char* argv[]) +{ + ck::DataTypeEnum data_type = ck::DataTypeEnum::Half; + bool do_verification = true; + int init_method = 2; + bool do_log = false; + bool time_kernel = true; + + std::vector in_length = {2, 32, 30, 30}; + std::vector wsize = {2, 2}; + std::vector wstride = {2, 2}; + std::vector wdilation = {1, 1}; + std::vector pad1 = {1, 1}; + std::vector pad2 = {1, 1}; + + if(argc != 2 && argc != 33) + { + print_help_max_pool2d_bwd(); + return 0; + } + else if(argc == 33) + { + data_type = static_cast(std::stoi(argv[2])); + do_verification = std::stoi(argv[3]); + init_method = std::stoi(argv[4]); + do_log = std::stoi(argv[5]); + time_kernel = std::stoi(argv[6]); + + // parse the long options + maxPoolbwdArgParser arg_parser; + arg_parser(argc, argv); + in_length = arg_parser.long_opts["length"]; + wsize = arg_parser.long_opts["wsize"]; + wstride = arg_parser.long_opts["wstride"]; + wdilation = arg_parser.long_opts["wdilation"]; + pad1 = arg_parser.long_opts["pad1"]; + pad2 = arg_parser.long_opts["pad2"]; + } + + using F16 = ck::half_t; + using BF16 = ck::bhalf_t; + using F32 = float; + using I8 = int8_t; + using I32 = int32_t; + + if(data_type == ck::DataTypeEnum::Half) + { + ck::profiler::profile_max_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::BFloat16) + { + ck::profiler::profile_max_pool2d_bwd_impl( + do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::Float) + { + ck::profiler::profile_max_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else if(data_type == ck::DataTypeEnum::Int8) + { + ck::profiler::profile_max_pool2d_bwd_impl(do_verification, + init_method, + do_log, + time_kernel, + in_length, + wsize, + wstride, + wdilation, + pad1, + pad2); + } + else + { + + throw std::runtime_error("not implemented yet"); + } + + return 0; +} + +REGISTER_PROFILER_OPERATION("max_pool2d_bwd", "max_pool2d bwd", profile_max_pool2d_bwd); diff --git a/test/pool/CMakeLists.txt b/test/pool/CMakeLists.txt index 0118a7591b..06eb8b85ed 100644 --- a/test/pool/CMakeLists.txt +++ b/test/pool/CMakeLists.txt @@ -4,10 +4,14 @@ add_gtest_executable(test_avg_pool3d_bwd test_avg_pool3d_bwd.cpp) add_gtest_executable(test_max_pool3d_bwd test_max_pool3d_bwd.cpp) add_gtest_executable(test_avg_pool3d_fwd test_avg_pool3d_fwd.cpp) add_gtest_executable(test_max_pool3d_fwd test_max_pool3d_fwd.cpp) +add_gtest_executable(test_avg_pool2d_bwd test_avg_pool2d_bwd.cpp) +add_gtest_executable(test_max_pool2d_bwd test_max_pool2d_bwd.cpp) add_gtest_executable(test_avg_pool2d_fwd test_avg_pool2d_fwd.cpp) add_gtest_executable(test_max_pool2d_fwd test_max_pool2d_fwd.cpp) target_link_libraries(test_avg_pool3d_bwd PRIVATE utility device_avg_pool3d_bwd_instance) +target_link_libraries(test_avg_pool2d_bwd PRIVATE utility device_avg_pool2d_bwd_instance) +target_link_libraries(test_max_pool2d_bwd PRIVATE utility device_max_pool_bwd_instance) target_link_libraries(test_max_pool3d_bwd PRIVATE utility device_max_pool_bwd_instance) target_link_libraries(test_avg_pool3d_fwd PRIVATE utility device_pool3d_fwd_instance) target_link_libraries(test_max_pool3d_fwd PRIVATE utility device_pool3d_fwd_instance) @@ -18,5 +22,7 @@ add_dependencies(test_pool test_avg_pool3d_bwd) add_dependencies(test_pool test_max_pool3d_bwd) add_dependencies(test_pool test_avg_pool3d_fwd) add_dependencies(test_pool test_max_pool3d_fwd) +add_dependencies(test_pool test_avg_pool2d_bwd) +add_dependencies(test_pool test_max_pool2d_bwd) add_dependencies(test_pool test_avg_pool2d_fwd) add_dependencies(test_pool test_max_pool2d_fwd) diff --git a/test/pool/test_avg_pool2d_bwd.cpp b/test/pool/test_avg_pool2d_bwd.cpp new file mode 100644 index 0000000000..0866325fc6 --- /dev/null +++ b/test/pool/test_avg_pool2d_bwd.cpp @@ -0,0 +1,133 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "gtest/gtest.h" +#include "profiler/profile_avg_pool2d_bwd_impl.hpp" +#include "test_pool_fwd_common.hpp" + +template +class AvgPool2dBWDTest : public ::testing::Test +{ + protected: + using InDataType = std::tuple_element_t<0, T>; + using OutDataType = std::tuple_element_t<1, T>; + + static std::vector params; + + void Run() + { + for(auto param : this->params) + { + bool success = + ck::profiler::profile_avg_pool2d_bwd_impl( + true, + 2, + false, + false, + param.length_, + param.window_spatial_lengths_, + param.window_strides_, + param.window_dilations_, + param.input_left_pads_, + param.input_right_pads_); + EXPECT_TRUE(success); + } + } +}; + +template +std::vector AvgPool2dBWDTest::params = { + {{1, 1, 1, 1}, {1, 1}, {1, 1}, {1, 1}, {0, 0}, {0, 0}}, + {{1, 1, 64, 64}, {64, 64}, {1, 1}, {1, 1}, {0, 0}, {0, 0}}, + {{1, 5, 7, 7}, {2, 2}, {2, 2}, {1, 1}, {2, 2}, {0, 0}}, + {{1, 1, 8, 8}, {2, 2}, {2, 2}, {1, 1}, {2, 2}, {0, 0}}, + {{1, 1, 8, 8}, {2, 2}, {1, 1}, {1, 1}, {1, 1}, {0, 0}}, + {{2, 32, 30, 30}, {2, 2}, {2, 2}, {1, 1}, {1, 1}, {1, 1}}, + {{1, 2, 30, 30}, {2, 2}, {2, 2}, {1, 1}, {0, 0}, {0, 0}}}; + +using Avg_Pool_2D_f32_types = ::testing::Types>; +using Avg_Pool_2D_int8_types = ::testing::Types>; +using Avg_Pool_2D_f16_types = ::testing::Types>; +using Avg_Pool_2D_bf16_types = ::testing::Types>; +using Avg_Pool_2D_f8_types = ::testing::Types>; + +template +class AvgPool2D_f32 : public AvgPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_FP32) + { + GTEST_SKIP() << "Skipping AvgPool2D_f32 tests because CK_ENABLE_FP32 is not enabled"; + } + } +}; + +template +class AvgPool2D_int8 : public AvgPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_INT8) + { + GTEST_SKIP() << "Skipping AvgPool2D_int8 tests because CK_ENABLE_INT8 is not enabled"; + } + } +}; + +template +class AvgPool2D_f16 : public AvgPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_FP16) + { + GTEST_SKIP() << "Skipping AvgPool2D_f16 because CK_ENABLE_FP16 is not enabled"; + } + } +}; + +template +class AvgPool2D_bf16 : public AvgPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_BF16) + { + GTEST_SKIP() << "Skipping AvgPool2D_bf16 tests because CK_ENABLE_BF16 is not enabled"; + } + } +}; + +template +class AvgPool2D_f8 : public AvgPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_FP8) + { + GTEST_SKIP() << "Skipping AvgPool2D_f8 tests because CK_ENABLE_FP8 is not enabled"; + } + } +}; + +TYPED_TEST_SUITE(AvgPool2D_f32, Avg_Pool_2D_f32_types); +TYPED_TEST_SUITE(AvgPool2D_int8, Avg_Pool_2D_int8_types); +TYPED_TEST_SUITE(AvgPool2D_f16, Avg_Pool_2D_f16_types); +TYPED_TEST_SUITE(AvgPool2D_bf16, Avg_Pool_2D_bf16_types); +TYPED_TEST_SUITE(AvgPool2D_f8, Avg_Pool_2D_f8_types); + +TYPED_TEST(AvgPool2D_f32, AvgPool2DTest_f32) { this->Run(); } + +TYPED_TEST(AvgPool2D_int8, AvgPool2DTest_int8) { this->Run(); } + +TYPED_TEST(AvgPool2D_f16, AvgPool2DTest_f16) { this->Run(); } + +TYPED_TEST(AvgPool2D_bf16, AvgPool2DTest_bf16) { this->Run(); } + +TYPED_TEST(AvgPool2D_f8, AvgPool2DTest_f8) { this->Run(); } diff --git a/test/pool/test_max_pool2d_bwd.cpp b/test/pool/test_max_pool2d_bwd.cpp new file mode 100644 index 0000000000..eae8f2c4d5 --- /dev/null +++ b/test/pool/test_max_pool2d_bwd.cpp @@ -0,0 +1,122 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "gtest/gtest.h" +#include "profiler/profile_max_pool2d_bwd_impl.hpp" +#include "test_pool_fwd_common.hpp" + +template +class MaxPool2dBWDTest : public ::testing::Test +{ + protected: + using DOutDataType = std::tuple_element_t<0, T>; + using DInDataType = std::tuple_element_t<1, T>; + using IndexDataType = std::tuple_element_t<2, T>; + + using InDataType = DInDataType; + using OutDataType = DOutDataType; + + static std::vector params; + + void Run() + { + for(auto param : this->params) + { + bool success = + ck::profiler::profile_max_pool2d_bwd_impl(true, + 2, + false, + false, + param.length_, + param.window_spatial_lengths_, + param.window_strides_, + param.window_dilations_, + param.input_left_pads_, + param.input_right_pads_); + EXPECT_TRUE(success); + } + } +}; + +template +std::vector MaxPool2dBWDTest::params = { + {{1, 1, 1, 1}, {1, 1}, {1, 1}, {1, 1}, {0, 0}, {0, 0}}, + {{2, 16, 64, 64}, {64, 64}, {1, 1}, {1, 1}, {0, 0}, {0, 0}}, + {{2, 16, 64, 64}, {4, 4}, {4, 4}, {2, 2}, {0, 0}, {0, 0}}, + {{2, 32, 30, 30}, {2, 2}, {2, 2}, {1, 1}, {1, 1}, {1, 1}}, + {{2, 2, 30, 30}, {2, 2}, {2, 2}, {1, 1}, {1, 1}, {1, 1}}}; + +using Max_Pool_2D_f32_types = ::testing::Types>; +using Max_Pool_2D_int8_types = ::testing::Types>; +using Max_Pool_2D_f16_types = ::testing::Types>; +using Max_Pool_2D_bf16_types = ::testing::Types>; + +template +class MaxPool2D_f32 : public MaxPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_FP32) + { + GTEST_SKIP() << "Skipping MaxPool2D_f32 tests because CK_ENABLE_FP32 is not enabled"; + } + } +}; + +template +class MaxPool2D_int8 : public MaxPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_INT8) + { + GTEST_SKIP() << "Skipping MaxPool2D_int8 tests because CK_ENABLE_INT8 is not enabled"; + } + } +}; + +template +class MaxPool2D_f16 : public MaxPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_FP16) + { + GTEST_SKIP() << "Skipping MaxPool2D_f16 because CK_ENABLE_FP16 is not enabled"; + } + } +}; + +template +class MaxPool2D_bf16 : public MaxPool2dBWDTest +{ + protected: + void SetUp() override + { + if(!CK_ENABLE_BF16) + { + GTEST_SKIP() << "Skipping MaxPool2D_bf16 tests because CK_ENABLE_BF16 is not enabled"; + } + } +}; + +TYPED_TEST_SUITE(MaxPool2D_f32, Max_Pool_2D_f32_types); +TYPED_TEST_SUITE(MaxPool2D_int8, Max_Pool_2D_int8_types); +TYPED_TEST_SUITE(MaxPool2D_f16, Max_Pool_2D_f16_types); +TYPED_TEST_SUITE(MaxPool2D_bf16, Max_Pool_2D_bf16_types); + +TYPED_TEST(MaxPool2D_f32, MaxPool2DTest_f32) { this->Run(); } + +TYPED_TEST(MaxPool2D_int8, MaxPool2DTest_int8) { this->Run(); } + +TYPED_TEST(MaxPool2D_f16, MaxPool2DTest_f16) { this->Run(); } + +TYPED_TEST(MaxPool2D_bf16, MaxPool2DTest_bf16) { this->Run(); } diff --git a/test/pool/test_pool_fwd_common.hpp b/test/pool/test_pool_fwd_common.hpp index 5917a27e56..0302c3c337 100644 --- a/test/pool/test_pool_fwd_common.hpp +++ b/test/pool/test_pool_fwd_common.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" @@ -8,8 +8,11 @@ using F16 = ck::half_t; using BF16 = ck::bhalf_t; using F32 = float; using I32 = int32_t; +using I8 = int8_t; +using F8 = ck::f8_t; using ck::index_t; using NDHWC = ck::tensor_layout::convolution::NDHWC; +using NHWC = ck::tensor_layout::convolution::NHWC; struct PoolingParam {