diff --git a/onnxruntime/core/providers/cpu/tensor/upsample.cc b/onnxruntime/core/providers/cpu/tensor/upsample.cc index d11eef247f..9c0e66a61f 100644 --- a/onnxruntime/core/providers/cpu/tensor/upsample.cc +++ b/onnxruntime/core/providers/cpu/tensor/upsample.cc @@ -29,6 +29,7 @@ ONNX_CPU_OPERATOR_VERSIONED_TYPED_KERNEL( KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), Upsample); + template void UpsampleNearest2x(int64_t batch_size, int64_t num_channels, @@ -594,7 +595,7 @@ Status Upsample::BaseCompute(OpKernelContext* context, if (dims.size() != scales.size()) return Status(ONNXRUNTIME, INVALID_ARGUMENT, - is_resize ? "Resize: input tensor's dimension does not match the scales." + is_resize_ ? "Resize: input tensor's dimension does not match the scales." : "Upsample: input tensor's dimension does not match the scales."); if (roi.size() != 2 * X->Shape().GetDims().size()) @@ -614,7 +615,7 @@ Status Upsample::BaseCompute(OpKernelContext* context, switch (mode_) { case UpsampleMode::NN: return UpsampleNearest(X->template Data(), Y->template MutableData(), X->Shape(), Y->Shape(), scales, roi, - is_resize, use_extrapolation_, extrapolation_value_, use_nearest2x_optimization, + is_resize_, use_extrapolation_, extrapolation_value_, use_nearest2x_optimization_, get_original_coordinate_, get_nearest_pixel_); case UpsampleMode::LINEAR: { //The correct behavior of 'linear' mode for an N-D input is not clear right now, @@ -623,7 +624,7 @@ Status Upsample::BaseCompute(OpKernelContext* context, std::ostringstream oss; oss << "'Linear' mode only support 2-D inputs ('Bilinear') or 4-D inputs " "with the corresponding outermost 2 scale values being 1 in the "; - oss << (is_resize ? "Resize operator" : "Upsample operator"); + oss << (is_resize_ ? "Resize operator" : "Upsample operator"); return Status(ONNXRUNTIME, FAIL, oss.str()); } @@ -659,7 +660,7 @@ Status Upsample::BaseCompute(OpKernelContext* context, return Status::OK(); } default: - return Status(ONNXRUNTIME, FAIL, is_resize ? "Resize: unexpected mode" : "Upsample: unexpected mode"); + return Status(ONNXRUNTIME, FAIL, is_resize_ ? "Resize: unexpected mode" : "Upsample: unexpected mode"); } } diff --git a/onnxruntime/core/providers/cpu/tensor/upsample.h b/onnxruntime/core/providers/cpu/tensor/upsample.h index 4beef0ffd6..4617e30809 100644 --- a/onnxruntime/core/providers/cpu/tensor/upsample.h +++ b/onnxruntime/core/providers/cpu/tensor/upsample.h @@ -25,13 +25,32 @@ enum UpsampleMode { CUBIC = 2, // cubic interpolation }; +enum ResizeCoordinateTransformationMode { + HALF_PIXEL = 0, + ASYMMETRIC = 1, + PYTORCH_HALF_PIXEL = 2, + TF_HALF_PIXEL_FOR_NN = 3, + ALIGN_CORNERS = 4, + TF_CROP_AND_RESIZE = 5, + CoordinateTransformationModeCount = 6, +}; + +enum ResizeNearestMode { + SIMPLE = 0, // For resize op 10 + ROUND_PREFER_FLOOR = 1, + ROUND_PREFER_CEIL = 2, + FLOOR = 3, + CEIL = 4, + NearestModeCount = 5, +}; + class UpsampleBase { protected: UpsampleBase(OpKernelInfo info) : scales_cached_(false), roi_cached_(false), use_extrapolation_(false) { int start; int end; info.GetKernelDef().SinceVersion(&start, &end); - is_resize = (start >= 10); + is_resize_ = (start >= 10); std::string mode; ORT_ENFORCE(info.GetAttr("mode", &mode).IsOK()); @@ -47,14 +66,18 @@ class UpsampleBase { extrapolation_value_ = info.GetAttrOrDefault("extrapolation_value", 0.0f); // Coordinate transformation mode attr was introduced in version 11, before that asymmetric mode was the only available transformation mode - std::string coordinate_transform_mode = start > 10 - ? info.GetAttrOrDefault("coordinate_transformation_mode", "half_pixel") - : "asymmetric"; - get_original_coordinate_ = GetOriginalCoordinateFromResizedCoordinate(coordinate_transform_mode); - use_extrapolation_ = need_roi_input_ = coordinate_transform_mode == "tf_crop_and_resize" ? true : false; + std::string coordinate_transform_mode_name = start > 10 + ? info.GetAttrOrDefault("coordinate_transformation_mode", "half_pixel") + : "asymmetric"; + coordinate_transform_mode_ = StringToCoordinateTransformationMode(coordinate_transform_mode_name); + get_original_coordinate_ = GetOriginalCoordinateFromResizedCoordinate(coordinate_transform_mode_); + use_extrapolation_ = need_roi_input_ = (coordinate_transform_mode_ == TF_CROP_AND_RESIZE); - std::string nearest_mode = info.GetAttrOrDefault("nearest_mode", "round_prefer_floor"); - get_nearest_pixel_ = GetNearestPixelFromOriginal(nearest_mode, start); + std::string nearest_mode_name = (mode_ == NN && start >= 11) + ? info.GetAttrOrDefault("nearest_mode", "round_prefer_floor") + : ""; + nearest_mode_ = StringToNearstMode(nearest_mode_name); + get_nearest_pixel_ = GetNearestPixelFromOriginal(nearest_mode_); cubic_coeff_a_ = info.GetAttrOrDefault("cubic_coeff_a", -0.75f); exclude_outside_ = info.GetAttrOrDefault("exclude_outside", 0) == 0 ? false : true; @@ -65,7 +88,7 @@ class UpsampleBase { // after version 11 update, this optimization is no longer applicable for all the available modes... // TODO : needs more testing to enable this for version 11 - use_nearest2x_optimization = start > 10 ? false : true; + use_nearest2x_optimization_ = start > 10 ? false : true; if (start > 10) { roi_input_idx_ = 1; @@ -98,13 +121,15 @@ class UpsampleBase { } } + UpsampleMode mode_; + ResizeCoordinateTransformationMode coordinate_transform_mode_; GetOriginalCoordinateFunc get_original_coordinate_; + ResizeNearestMode nearest_mode_; GetNearestPixelFunc get_nearest_pixel_; float cubic_coeff_a_; bool exclude_outside_; float extrapolation_value_; - UpsampleMode mode_; - bool use_nearest2x_optimization = false; + bool use_nearest2x_optimization_ = false; std::vector scales_; std::vector roi_; @@ -112,7 +137,7 @@ class UpsampleBase { bool roi_cached_; bool need_roi_input_; bool use_extrapolation_; - bool is_resize = false; + bool is_resize_ = false; int roi_input_idx_ = -1; int scales_input_idx_ = -1; @@ -125,7 +150,6 @@ class UpsampleBase { if (mode == UpsampleModeLinear) { return UpsampleMode::LINEAR; } - if (mode == UpsampleModeCubic) { return UpsampleMode::CUBIC; } @@ -133,78 +157,114 @@ class UpsampleBase { UpsampleModeNN + "(default) or " + UpsampleModeLinear + " or " + UpsampleModeCubic + "."); } + ResizeCoordinateTransformationMode StringToCoordinateTransformationMode( + const std::string& coordinate_transform_mode_name) { + if (coordinate_transform_mode_name == "asymmetric") { + return ASYMMETRIC; + } + if (coordinate_transform_mode_name == "pytorch_half_pixel") { + return PYTORCH_HALF_PIXEL; + } + if (coordinate_transform_mode_name == "tf_half_pixel_for_nn") { + return TF_HALF_PIXEL_FOR_NN; + } + if (coordinate_transform_mode_name == "align_corners") { + return ALIGN_CORNERS; + } + if (coordinate_transform_mode_name == "tf_crop_and_resize") { + return TF_CROP_AND_RESIZE; + } + if (coordinate_transform_mode_name == "half_pixel") { + return HALF_PIXEL; + } + ORT_THROW("coordinate_transform_mode:[" + coordinate_transform_mode_name + "] is not supportted!"); + } + GetOriginalCoordinateFunc GetOriginalCoordinateFromResizedCoordinate( - const std::string& coordinate_transform_mode) { - if (coordinate_transform_mode == "asymmetric") { - return [](float x_resized, float x_scale, float, float, float, float) { - return x_resized / x_scale; - }; - } else if (coordinate_transform_mode == "pytorch_half_pixel") { - return [](float x_resized, float x_scale, float length_resized, float, float, float) { - return length_resized > 1 ? (x_resized + 0.5f) / x_scale - 0.5f : 0.0f; - }; - } else if (coordinate_transform_mode == "tf_half_pixel_for_nn") { - return [](float x_resized, float x_scale, float, float, float, float) { - return (x_resized + 0.5f) / x_scale; - }; - } else if (coordinate_transform_mode == "align_corners") { - return [](float x_resized, float, float length_resized, float length_original, float, float) { - return length_resized == 1 ? 0 : x_resized * (length_original - 1) / (length_resized - 1); - }; - } else if (coordinate_transform_mode == "tf_crop_and_resize") { - return [](float x_resized, float, float length_resized, float length_original, float roi_start, float roi_end) { - auto orig = length_resized > 1 - ? roi_start * (length_original - 1) + (x_resized * (roi_end - roi_start) * (length_original - 1)) / (length_resized - 1) - : 0.5 * (roi_start + roi_end) * (length_original - 1); - return static_cast(orig); - }; - } else { // "half_pixel" - return [](float x_resized, float x_scale, float, float, float, float) { - return ((x_resized + 0.5f) / x_scale) - 0.5f; - }; + ResizeCoordinateTransformationMode coordinate_transform_mode) { + switch (coordinate_transform_mode) { + case ASYMMETRIC: + return [](float x_resized, float x_scale, float, float, float, float) { + return x_resized / x_scale; + }; + case PYTORCH_HALF_PIXEL: + return [](float x_resized, float x_scale, float length_resized, float, float, float) { + return length_resized > 1 ? (x_resized + 0.5f) / x_scale - 0.5f : 0.0f; + }; + case TF_HALF_PIXEL_FOR_NN: + return [](float x_resized, float x_scale, float, float, float, float) { + return (x_resized + 0.5f) / x_scale; + }; + case ALIGN_CORNERS: + return [](float x_resized, float, float length_resized, float length_original, float, float) { + return length_resized == 1 ? 0 : x_resized * (length_original - 1) / (length_resized - 1); + }; + case TF_CROP_AND_RESIZE: + return [](float x_resized, float, float length_resized, float length_original, float roi_start, float roi_end) { + auto orig = length_resized > 1 + ? roi_start * (length_original - 1) + (x_resized * (roi_end - roi_start) * (length_original - 1)) / (length_resized - 1) + : 0.5 * (roi_start + roi_end) * (length_original - 1); + return static_cast(orig); + }; + default: // "half_pixel" + return [](float x_resized, float x_scale, float, float, float, float) { + return ((x_resized + 0.5f) / x_scale) - 0.5f; + }; } } - GetNearestPixelFunc GetNearestPixelFromOriginal( - const std::string& nearest_mode, int opset_version) { - // versions older than 11 did not have nearest_mode attr. Use the original logic in this case - // to maintain backward compatibility - if (opset_version < 11) { - return [](float x_original, bool isDownSample) { - if (isDownSample) { - return static_cast(std::ceil(x_original)); - } else { - return static_cast(x_original); - } - }; + ResizeNearestMode StringToNearstMode(const std::string& nearst_mode_name) { + if (nearst_mode_name == "round_prefer_floor") { + return ROUND_PREFER_FLOOR; + } else if (nearst_mode_name == "round_prefer_ceil") { + return ROUND_PREFER_CEIL; + } else if (nearst_mode_name == "floor") { + return FLOOR; + } else if (nearst_mode_name == "ceil") { + return CEIL; + } else if (nearst_mode_name == "") { + return SIMPLE; } + ORT_THROW("nearst_mode:[" + nearst_mode_name + "] is not supportted!"); + } - // if opset version >=11 choose the rounding mode based on nearest_mode attr - if (nearest_mode == "round_prefer_ceil") { - return [](float x_original, bool) { - return static_cast(std::round(x_original)); - }; - } else if (nearest_mode == "floor") { - return [](float x_original, bool) { - return static_cast(std::floor(x_original)); - }; - } else if (nearest_mode == "ceil") { - return [](float x_original, bool) { - return static_cast(std::ceil(x_original)); - }; - } else { // default is round_prefer_floor - return [](float x_original, bool) { - // for half way cases prefer floor - if (x_original == static_cast(x_original) + 0.5f) { + GetNearestPixelFunc GetNearestPixelFromOriginal(ResizeNearestMode nearest_mode) { + switch (nearest_mode) { + case SIMPLE: + // versions older than 11 did not have nearest_mode attr. Use the original logic in this case + // to maintain backward compatibility + return [](float x_original, bool isDownSample) { + if (isDownSample) { + return static_cast(std::ceil(x_original)); + } else { + return static_cast(x_original); + } + }; + case ROUND_PREFER_CEIL: + return [](float x_original, bool) { + return static_cast(std::round(x_original)); + }; + case FLOOR: + return [](float x_original, bool) { return static_cast(std::floor(x_original)); - } - return static_cast(std::round(x_original)); - }; + }; + case CEIL: + return [](float x_original, bool) { + return static_cast(std::ceil(x_original)); + }; + default: // default is round_prefer_floor + return [](float x_original, bool) { + // for half way cases prefer floor + if (x_original == static_cast(x_original) + 0.5f) { + return static_cast(std::floor(x_original)); + } + return static_cast(std::round(x_original)); + }; } } - void ScalesValidation(const std::vector& scales, const UpsampleMode mode) const { - if (!is_resize) { + void ScalesValidation(const std::vector& scales, const UpsampleMode mode) const { + if (!is_resize_) { for (auto& scale : scales) { ORT_ENFORCE(scale >= 1, "Scale value should be greater than or equal to 1."); } @@ -218,7 +278,7 @@ class UpsampleBase { ORT_ENFORCE(scales.size() == 2 || (scales.size() == 4 && scales[0] == 1 && scales[1] == 1), "'Linear' mode and 'Cubic' mode only support 2-D inputs ('Bilinear', 'Bicubic') or 4-D inputs " "with the corresponding outermost 2 scale values being 1 in the ", - is_resize ? "Resize operator" : "Upsample operator"); + is_resize_ ? "Resize operator" : "Upsample operator"); } } @@ -257,7 +317,7 @@ class UpsampleBase { output_dims[i] = static_cast(scales[i] * input_dims[i]); } } -}; +}; // UpsampleBase template class Upsample : public UpsampleBase, public OpKernel { diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index b33be3c7e0..c0604f036c 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -557,11 +557,11 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kO class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, double, MaxPool); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, MaxPool); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, NonMaxSuppression); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, Resize); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, Resize); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, Resize); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int32_t, Resize); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, Resize); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, float, Resize); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, double, Resize); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, Resize); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, int32_t, Resize); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, uint8_t, Resize); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, ReverseSequence); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, RoiAlign); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, RoiAlign); @@ -652,6 +652,11 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, MaxPool); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, MaxPool); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, MaxPool); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Resize); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Resize); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, bool, Equal); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Equal); @@ -1007,11 +1012,11 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1103,6 +1108,11 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cuda/tensor/resize.cc b/onnxruntime/core/providers/cuda/tensor/resize.cc index 96f05017c3..4940c872eb 100644 --- a/onnxruntime/core/providers/cuda/tensor/resize.cc +++ b/onnxruntime/core/providers/cuda/tensor/resize.cc @@ -6,15 +6,27 @@ namespace onnxruntime { namespace cuda { #define REGISTER_KERNEL_TYPED(T) \ - ONNX_OPERATOR_TYPED_KERNEL_EX( \ + ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \ Resize, \ kOnnxDomain, \ - 10, \ + 10, 10, \ T, \ kCudaExecutionProvider, \ KernelDefBuilder() \ .InputMemoryType(1) \ .TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + Resize); \ + ONNX_OPERATOR_TYPED_KERNEL_EX( \ + Resize, \ + kOnnxDomain, \ + 11, \ + T, \ + kCudaExecutionProvider, \ + KernelDefBuilder() \ + .InputMemoryType(1) \ + .InputMemoryType(2) \ + .InputMemoryType(3) \ + .TypeConstraint("T", DataTypeImpl::GetTensorType()), \ Resize); REGISTER_KERNEL_TYPED(float) diff --git a/onnxruntime/core/providers/cuda/tensor/resize_impl.cu b/onnxruntime/core/providers/cuda/tensor/resize_impl.cu index 55d7fcaf01..1442363110 100644 --- a/onnxruntime/core/providers/cuda/tensor/resize_impl.cu +++ b/onnxruntime/core/providers/cuda/tensor/resize_impl.cu @@ -3,147 +3,245 @@ namespace onnxruntime { namespace cuda { + +using onnxruntime::ResizeCoordinateTransformationMode; +using onnxruntime::ResizeNearestMode; +using onnxruntime::UpsampleMode; + +__device__ int NearestPixel_SIMPLE(float x_original, bool is_down_sampling) { + if (is_down_sampling) { + return static_cast(ceil(x_original)); + } else { + return static_cast(x_original); + } +} + +__device__ int NearestPixel_ROUND_PREFER_FLOOR(float x_original, bool) { + if (x_original == static_cast(x_original) + 0.5f) { + return static_cast(floor(x_original)); + } + return static_cast(round(x_original)); +} + +__device__ int NearestPixel_ROUND_PREFER_CEIL(float x_original, bool) { + return static_cast(round(x_original)); +} + +__device__ int NearestPixel_FLOOR(float x_original, bool) { + return static_cast(floor(x_original)); +} + +__device__ int NearestPixel_CEIL(float x_original, bool) { + return static_cast(ceil(x_original)); +} + +using CudaFunctionNearestPixel = int (*)(float, bool); +__device__ CudaFunctionNearestPixel func_NearestPixel_SIMPLE = NearestPixel_SIMPLE; +__device__ CudaFunctionNearestPixel func_NearestPixel_ROUND_PREFER_FLOOR = NearestPixel_ROUND_PREFER_FLOOR; +__device__ CudaFunctionNearestPixel func_NearestPixel_ROUND_PREFER_CEIL = NearestPixel_ROUND_PREFER_CEIL; +__device__ CudaFunctionNearestPixel func_NearestPixel_FLOOR = NearestPixel_FLOOR; +__device__ CudaFunctionNearestPixel func_NearestPixel_CEIL = NearestPixel_CEIL; + +CudaFunctionNearestPixel GetDeviceNearstPixelFunction(ResizeNearestMode nearest_mode) { + static bool already_copied = false; + static std::mutex s_mutext; + static CudaFunctionNearestPixel s_nearest_pixel[ResizeNearestMode::NearestModeCount]; + if (!already_copied) { + std::lock_guard lock(s_mutext); + if (!already_copied) { + CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::SIMPLE], + func_NearestPixel_SIMPLE, sizeof(CudaFunctionNearestPixel))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::ROUND_PREFER_FLOOR], + func_NearestPixel_ROUND_PREFER_FLOOR, sizeof(CudaFunctionNearestPixel))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::ROUND_PREFER_CEIL], + func_NearestPixel_ROUND_PREFER_CEIL, sizeof(CudaFunctionNearestPixel))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::FLOOR], + func_NearestPixel_FLOOR, sizeof(CudaFunctionNearestPixel))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::CEIL], + func_NearestPixel_CEIL, sizeof(CudaFunctionNearestPixel))); + already_copied = true; + } + } + return s_nearest_pixel[nearest_mode]; +} + +__device__ float TransformCoordinate_ASYMMETRIC(float x_resized, float x_scale, float, float, float, float) { + return x_resized / x_scale; +} + +__device__ float TransformCoordinate_HALF_PIXEL(float x_resized, float x_scale, float, float, float, float) { + return ((x_resized + 0.5f) / x_scale) - 0.5f; +} + +__device__ float TransformCoordinate_PYTORCH_HALF_PIXEL( + float x_resized, float x_scale, float length_resized, float, float, float) { + return length_resized > 1 ? (x_resized + 0.5f) / x_scale - 0.5f : 0.0f; +} + +__device__ float TransformCoordinate_TF_HALF_PIXEL_FOR_NN( + float x_resized, float x_scale, float, float, float, float) { + return (x_resized + 0.5f) / x_scale; +} + +__device__ float TransformCoordinate_ALIGN_CORNERS( + float x_resized, float, float length_resized, float length_original, float, float) { + return length_resized == 1 ? 0 : x_resized * (length_original - 1) / (length_resized - 1); +} + +__device__ float TransformCoordinate_TF_CROP_AND_RESIZE( + float x_resized, float, float length_resized, float length_original, float roi_start, float roi_end) { + auto orig = length_resized > 1 + ? roi_start * (length_original - 1) + (x_resized * (roi_end - roi_start) * (length_original - 1)) / (length_resized - 1) + : 0.5 * (roi_start + roi_end) * (length_original - 1); + return static_cast(orig); +} + +using CudaFunctionOriginalCoordinate = float (*)(float, float, float, float, float, float); + +__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_ASYMMETRIC = TransformCoordinate_ASYMMETRIC; +__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_HALF_PIXEL = TransformCoordinate_HALF_PIXEL; +__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_PYTORCH_HALF_PIXEL = TransformCoordinate_PYTORCH_HALF_PIXEL; +__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_ALIGN_CORNERS = TransformCoordinate_ALIGN_CORNERS; +__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_TF_HALF_PIXEL_FOR_NN = TransformCoordinate_TF_HALF_PIXEL_FOR_NN; +__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_TF_CROP_AND_RESIZE = TransformCoordinate_TF_CROP_AND_RESIZE; + +CudaFunctionOriginalCoordinate GetDeviceOriginalCoordinateFunc(ResizeCoordinateTransformationMode coordinate_transform_mode) { + static bool already_copied = false; + static std::mutex s_mutext; + static CudaFunctionOriginalCoordinate s_coordinate_tranforms[ResizeCoordinateTransformationMode::CoordinateTransformationModeCount]; + if (!already_copied) { + std::lock_guard lock(s_mutext); + if (!already_copied) { + CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::HALF_PIXEL], + func_TransformCoordinate_HALF_PIXEL, sizeof(CudaFunctionOriginalCoordinate))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::ASYMMETRIC], + func_TransformCoordinate_ASYMMETRIC, sizeof(CudaFunctionOriginalCoordinate))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::PYTORCH_HALF_PIXEL], + func_TransformCoordinate_PYTORCH_HALF_PIXEL, sizeof(CudaFunctionOriginalCoordinate))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::ALIGN_CORNERS], + func_TransformCoordinate_ALIGN_CORNERS, sizeof(CudaFunctionOriginalCoordinate))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::TF_HALF_PIXEL_FOR_NN], + func_TransformCoordinate_TF_HALF_PIXEL_FOR_NN, sizeof(CudaFunctionOriginalCoordinate))); + CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::TF_CROP_AND_RESIZE], + func_TransformCoordinate_TF_CROP_AND_RESIZE, sizeof(CudaFunctionOriginalCoordinate))); + already_copied = true; + } + } + return s_coordinate_tranforms[coordinate_transform_mode]; +} + template -__global__ void _ResizeNearestKernel(const size_t rank, - const int64_t* input_pitches, - const fast_divmod* output_div_pitches, - const float* scales, - const T* input_data, - T* output_data, - const size_t N) { +__global__ void _ResizeNearestKernel( + const size_t rank, + const int64_t* input_shape, + const int64_t* output_shape, + const int64_t* input_pitches, + const fast_divmod* output_div_pitches, + const float* scales, + const float* roi, + const T* input_data, + T* output_data, + const size_t N, + bool extrapolation_enabled, + float extrapolation_value, + CudaFunctionOriginalCoordinate transform_coordinate, + CudaFunctionNearestPixel calc_nearest_pixel) { CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N); CUDA_LONG input_index = 0; CUDA_LONG output_index = id; int div, mod; + bool extrapolation_occured = false; for (int dim = 0; dim < rank; ++dim) { output_div_pitches[dim].divmod(output_index, div, mod); output_index = mod; - if (scales[dim] <= 1) { //downsample - div = std::ceil(div / scales[dim]); - } else { //upsample - div = div / scales[dim]; + float orig_coord = transform_coordinate(static_cast(div), scales[dim], static_cast(output_shape[dim]), + static_cast(input_shape[dim]), roi[dim], roi[dim + rank]); + if (extrapolation_enabled && !extrapolation_occured) { + extrapolation_occured = (orig_coord < 0.f || orig_coord > static_cast(input_shape[dim] - 1)); } + div = calc_nearest_pixel(orig_coord, scales[dim] < 1); + if (div >= input_shape[dim]) div = input_shape[dim] - 1; + if (div < 0) div = 0; input_index += input_pitches[dim] * div; } - output_data[id] = input_data[input_index]; + output_data[id] = extrapolation_occured ? static_cast(extrapolation_value) : input_data[input_index]; } -// The following method supports a 4-D input in 'Linear mode' -// that amounts to 'Bilinear' Upsampling/Resizing in the sense that it assumes -// the scale values for the outermost 2 dimensions are 1. -// This is the common use-case where the 4-D input (batched multi-channel images) -// is usually of shape [N, C, H, W] and the scales are [1.0, 1.0, height_scale, width_scale] +struct BilinearMappingInfo { + int origin_; + float weight_; + int extrapolate_; +}; + template -__global__ void _ResizeBilinear4DInputKernel(const int64_t input_dim2, - const int64_t* input_pitches, - const fast_divmod* output_div_pitches, - const float* scales, - const T* input_data, - T* output_data, - const size_t N) { - CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N); - CUDA_LONG input_index = 0; - - // For bilinear mode, scales[0]=scales[1]=1 - int mod; - int index_of_dim0, index_of_dim1, index_of_dim2, index_of_dim3; - output_div_pitches[0].divmod(id, index_of_dim0, mod); - output_div_pitches[1].divmod(mod, index_of_dim1, mod); - output_div_pitches[2].divmod(mod, index_of_dim2, mod); - index_of_dim3 = mod; - int index_of_input_dim2, index_of_input_dim3; - float x_offset_0, y_offset_0, x_offset_1, y_offset_1; - index_of_input_dim2 = static_cast(index_of_dim2 / scales[2]); - index_of_input_dim3 = static_cast(index_of_dim3 / scales[3]); - input_index = index_of_dim0 * input_pitches[0] + - index_of_dim1 * input_pitches[1] + - index_of_input_dim2 * input_pitches[2] + - index_of_input_dim3; - - T x00 = input_data[input_index]; - T x10, x01, x11; - - bool end_of_dim2 = false, end_of_dim3 = false; - if (index_of_input_dim2 == (input_dim2 - 1)) { - // It's the end in dimension 2 - x01 = x00; - end_of_dim2 = true; - } else { - x01 = input_data[input_index + input_pitches[2]]; +__global__ void _ResizeBilinearCoordinateMapping( + int64_t input_height, int64_t input_width, + int64_t output_height, int64_t output_width, + float scale_height, float scale_width, + float roi_height_start, float roi_height_end, + float roi_width_start, float roi_width_end, + const size_t SumHW, bool extrapolation_enabled, + CudaFunctionOriginalCoordinate transform_coordinate, + BilinearMappingInfo* dims_mapping) { + CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, SumHW); + if (id < output_height) { // y = id + float input_y = transform_coordinate(static_cast(id), scale_height, + static_cast(output_height), static_cast(input_height), + roi_height_start, roi_height_end); + dims_mapping[id].extrapolate_ = (int)(extrapolation_enabled && (input_y < 0 || input_y > static_cast(input_height - 1))); + input_y = max(0.0f, min(input_y, static_cast(input_height - 1))); + int y_int = static_cast(input_y); + dims_mapping[id].origin_ = y_int; + dims_mapping[id].weight_ = (y_int >= input_height - 1) ? 0.5f : input_y - y_int; + } else { //x = id - output_height + float input_x = transform_coordinate(static_cast(id - output_height), scale_width, + static_cast(output_width), static_cast(input_width), + roi_width_start, roi_width_end); + dims_mapping[id].extrapolate_ = (int)(extrapolation_enabled && (input_x < 0 || input_x > static_cast(input_width - 1))); + input_x = max(0.0f, min(input_x, static_cast(input_width - 1))); + int x_int = static_cast(input_x); + dims_mapping[id].origin_ = x_int; + dims_mapping[id].weight_ = (x_int >= input_width - 1) ? 0.5f : input_x - x_int; } - - if (index_of_input_dim3 == (input_pitches[2] - 1)) { - // It's the end in dimension 3 - x10 = x00; - x11 = x01; - end_of_dim3 = true; - } else { - x10 = input_data[input_index + 1]; - x11 = end_of_dim2 ? x10 : input_data[input_index + input_pitches[2] + 1]; - } - - y_offset_0 = end_of_dim2 ? 0.5f : index_of_dim2 / scales[2] - index_of_input_dim2; - y_offset_1 = 1.0f - y_offset_0; - x_offset_0 = end_of_dim3 ? 0.5f : index_of_dim3 / scales[3] - index_of_input_dim3; - x_offset_1 = 1.0f - x_offset_0; - - output_data[id] = - x00 * static_cast(y_offset_1 * x_offset_1) + - x01 * static_cast(y_offset_0 * x_offset_1) + - x10 * static_cast(y_offset_1 * x_offset_0) + - x11 * static_cast(y_offset_0 * x_offset_0); } -// The following method supports a 2-D input in 'Linear mode' +// The following method supports a N-D input in 'Linear mode'. Last two dimension is [H, W]. +// the scale values for the outer dimensions except last two are 1. template -__global__ void _ResizeBilinear2DInputKernel(const int64_t input_dim0, - const int64_t* input_pitches, - const fast_divmod* output_div_pitches, - const float* scales, - const T* input_data, - T* output_data, - const size_t N) { +__global__ void _ResizeBilinearKernel( + int64_t input_height, int64_t input_width, + int64_t output_height, int64_t output_width, + fast_divmod div_output_width, fast_divmod div_output_image, + const T* input_data, T* output_data, const size_t N, + float extrapolation_value, + BilinearMappingInfo* dims_mapping) { CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N); - CUDA_LONG input_index = 0; + int bxc, output_image_index; + div_output_image.divmod(id, bxc, output_image_index); + CUDA_LONG input_index = bxc * input_height * input_width; + int output_y, output_x; + div_output_width.divmod(output_image_index, output_y, output_x); - int mod; - int index_of_dim0, index_of_dim1; - output_div_pitches[0].divmod(id, index_of_dim0, mod); - index_of_dim1 = mod; - int index_of_input_dim0, index_of_input_dim1; - float x_offset_0, y_offset_0, x_offset_1, y_offset_1; - index_of_input_dim0 = static_cast(index_of_dim0 / scales[0]); - index_of_input_dim1 = static_cast(index_of_dim1 / scales[1]); - input_index = index_of_input_dim0 * input_pitches[0] + index_of_input_dim1; + if (dims_mapping[output_y].extrapolate_ || dims_mapping[output_x + output_height].extrapolate_) { + output_data[id] = extrapolation_value; + return; + } + float y_offset_0 = dims_mapping[output_y].weight_; + int y_int = dims_mapping[output_y].origin_; + float x_offset_0 = dims_mapping[output_x + output_height].weight_; + int x_int = dims_mapping[output_x + output_height].origin_; + input_index += y_int * input_width + x_int; T x00 = input_data[input_index]; - T x10, x01, x11; - - bool end_of_dim0 = false, end_of_dim1 = false; - if (index_of_input_dim0 == (input_dim0 - 1)) { - // It's the end in dimension 0 - x01 = x00; - end_of_dim0 = true; - } else { - x01 = input_data[input_index + input_pitches[0]]; - } - - if (index_of_input_dim1 == (input_pitches[0] - 1)) { - // It's the end in dimension 1 - x10 = x00; - x11 = x01; - end_of_dim1 = true; - } else { - x10 = input_data[input_index + 1]; - x11 = end_of_dim0 ? x10 : input_data[input_index + input_pitches[0] + 1]; - } - - y_offset_0 = end_of_dim0 ? 0.5f : index_of_dim0 / scales[0] - index_of_input_dim0; - y_offset_1 = 1.0f - y_offset_0; - x_offset_0 = end_of_dim1 ? 0.5f : index_of_dim1 / scales[1] - index_of_input_dim1; - x_offset_1 = 1.0f - x_offset_0; + bool end_of_h = (y_int >= input_height - 1); + bool end_of_w = (x_int >= input_width - 1); + T x10 = end_of_w ? x00 : input_data[input_index + 1]; + T x01 = end_of_h ? x00 : input_data[input_index + input_width]; + T x11 = end_of_w ? x01 : (end_of_h ? x10 : input_data[input_index + input_width + 1]); + float y_offset_1 = 1.0f - y_offset_0; + float x_offset_1 = 1.0f - x_offset_0; output_data[id] = x00 * static_cast(y_offset_1 * x_offset_1) + x01 * static_cast(y_offset_0 * x_offset_1) + @@ -152,41 +250,216 @@ __global__ void _ResizeBilinear2DInputKernel(const int64_t input_dim0, } template -void ResizeImpl(const onnxruntime::UpsampleMode upsample_mode, - const size_t rank, - const int64_t input_dim2, - const int64_t* input_pitches, - const fast_divmod* output_div_pitches, - const float* scales_vals, - const T* input_data, - T* output_data, - const size_t N) { +__device__ __forceinline__ float CubicInterpolationRowwise( + const T* image, int x, int y, int input_height, int input_width, + float coeff0, float coeff1, float coeff2, float coeff3) { + int row_index = max(0, min(y, input_height - 1)) * input_width; + return coeff0 * static_cast(image[row_index + max(0, min(x - 1, input_width - 1))]) + + coeff1 * static_cast(image[row_index + max(0, min(x, input_width - 1))]) + + coeff2 * static_cast(image[row_index + max(0, min(x + 1, input_width - 1))]) + + coeff3 * static_cast(image[row_index + max(0, min(x + 2, input_width - 1))]); +} + +struct CubicMappingInfo { + int origin_; + int extrapolate_; + float coeff0_; + float coeff1_; + float coeff2_; + float coeff3_; +}; + +template +__global__ void _ResizeCubicCoordinateMapping( + int64_t input_height, int64_t input_width, + int64_t output_height, int64_t output_width, + float scale_height, float scale_width, + float roi_height_start, float roi_height_end, + float roi_width_start, float roi_width_end, + const size_t SumHW, bool extrapolation_enabled, + float cubic_coeff_a, bool exclude_outside, + CudaFunctionOriginalCoordinate transform_coordinate, + CubicMappingInfo* dims_mapping) { + CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, SumHW); + auto& dm = dims_mapping[id]; + bool is_y_axis = (id < output_height); + int max_input_coord = static_cast(is_y_axis ? input_height : input_width); + + float input_coordinat = transform_coordinate( + static_cast(is_y_axis ? id : id - output_height), + (is_y_axis ? scale_height : scale_width), + static_cast(is_y_axis ? output_height : output_width), + static_cast(max_input_coord), + (is_y_axis ? roi_height_start : roi_width_start), + (is_y_axis ? roi_height_end : roi_width_end)); + int coord_int = static_cast(floor(input_coordinat)); + float s_coord = abs(input_coordinat - coord_int); + float coeff_sum = 1.0f; + float coeff_0 = static_cast(((cubic_coeff_a * (s_coord + 1) - 5 * cubic_coeff_a) * (s_coord + 1) + 8 * cubic_coeff_a) * (s_coord + 1) - 4 * cubic_coeff_a); + float coeff_1 = static_cast(((cubic_coeff_a + 2) * s_coord - (cubic_coeff_a + 3)) * s_coord * s_coord + 1); + float coeff_2 = static_cast(((cubic_coeff_a + 2) * (1 - s_coord) - (cubic_coeff_a + 3)) * (1 - s_coord) * (1 - s_coord) + 1); + float coeff_3 = static_cast(((cubic_coeff_a * (2 - s_coord) - 5 * cubic_coeff_a) * (2 - s_coord) + 8 * cubic_coeff_a) * (2 - s_coord) - 4 * cubic_coeff_a); + if (exclude_outside) { + coeff_0 = (coord_int - 1 < 0 || coord_int - 1 >= max_input_coord) ? 0.0 : coeff_0; + coeff_1 = (coord_int + 0 < 0 || coord_int + 0 >= max_input_coord) ? 0.0 : coeff_1; + coeff_2 = (coord_int + 1 < 0 || coord_int + 1 >= max_input_coord) ? 0.0 : coeff_2; + coeff_3 = (coord_int + 2 < 0 || coord_int + 2 >= max_input_coord) ? 0.0 : coeff_3; + coeff_sum = coeff_0 + coeff_1 + coeff_2 + coeff_3; + } + dm.origin_ = coord_int; + dm.coeff0_ = coeff_0 / coeff_sum; + dm.coeff1_ = coeff_1 / coeff_sum; + dm.coeff2_ = coeff_2 / coeff_sum; + dm.coeff3_ = coeff_3 / coeff_sum; + dm.extrapolate_ = (int)(extrapolation_enabled && (input_coordinat < 0 || input_coordinat > static_cast(max_input_coord - 1))); +} + +template +__global__ void _ResizeBiCubicKernel( + int64_t input_height, int64_t input_width, int64_t output_height, int64_t output_width, + fast_divmod div_output_width, fast_divmod div_output_image, + const T* input_data, T* output_data, const size_t N, float extrapolation_value, + CubicMappingInfo* dims_mapping) { + CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N); + int bxc, output_image_index, output_x, output_y; + div_output_image.divmod(id, bxc, output_image_index); + CUDA_LONG input_index = bxc * input_height * input_width; + div_output_width.divmod(output_image_index, output_y, output_x); + + CubicMappingInfo& y_info = dims_mapping[output_y]; + CubicMappingInfo& x_info = dims_mapping[output_x + output_height]; + if (y_info.extrapolate_ || x_info.extrapolate_) { + output_data[id] = extrapolation_value; + return; + } + + float w0 = x_info.coeff0_; + float w1 = x_info.coeff1_; + float w2 = x_info.coeff2_; + float w3 = x_info.coeff3_; + int x_int = x_info.origin_; + int y_int = y_info.origin_; + const T* image = input_data + input_index; + output_data[id] = y_info.coeff0_ * CubicInterpolationRowwise(image, x_int, y_int - 1, input_height, input_width, w0, w1, w2, w3) + + y_info.coeff1_ * CubicInterpolationRowwise(image, x_int, y_int, input_height, input_width, w0, w1, w2, w3) + + y_info.coeff2_ * CubicInterpolationRowwise(image, x_int, y_int + 1, input_height, input_width, w0, w1, w2, w3) + + y_info.coeff3_ * CubicInterpolationRowwise(image, x_int, y_int + 2, input_height, input_width, w0, w1, w2, w3); +} + +size_t CalcResizeBufferSize(const onnxruntime::UpsampleMode upsample_mode, + const std::vector& output_dims) { + switch (upsample_mode) { + case UpsampleMode::NN: + return 0; + case UpsampleMode::LINEAR: + return sizeof(BilinearMappingInfo) * std::accumulate(output_dims.rbegin(), output_dims.rbegin() + 2, 0); + case UpsampleMode::CUBIC: + return sizeof(CubicMappingInfo) * std::accumulate(output_dims.rbegin(), output_dims.rbegin() + 2, 0); + } + return 0; +} + +template +void ResizeImpl( + const UpsampleMode upsample_mode, + const int rank, + CudaKernel::CudaAsyncBuffer& input_shape, + CudaKernel::CudaAsyncBuffer& output_shape, + CudaKernel::CudaAsyncBuffer& input_strides, + CudaKernel::CudaAsyncBuffer& output_div_pitches, + CudaKernel::CudaAsyncBuffer& scales_vals, + CudaKernel::CudaAsyncBuffer& roi_vals, + const T* input_data, + T* output_data, + const size_t N, + bool extrapolation_enabled, + float extrapolation_value, + float cubic_coeff_a, + bool exclude_outside, + ResizeCoordinateTransformationMode coordinate_transform_mode, + ResizeNearestMode nearest_mode, + void* dims_mapping) { int blocksPerGrid = (int)(ceil(static_cast(N) / GridDim::maxThreadsPerBlock)); - if (onnxruntime::UpsampleMode::NN == upsample_mode) { - _ResizeNearestKernel<<>>( - rank, input_pitches, output_div_pitches, scales_vals, - input_data, output_data, N); - } else if (onnxruntime::UpsampleMode::LINEAR == upsample_mode && rank == 4) { - _ResizeBilinear4DInputKernel<<>>( - input_dim2, input_pitches, output_div_pitches, scales_vals, - input_data, output_data, N); - } else if (onnxruntime::UpsampleMode::LINEAR == upsample_mode && rank == 2) { - _ResizeBilinear2DInputKernel<<>>( - input_dim2, input_pitches, output_div_pitches, scales_vals, - input_data, output_data, N); + CudaFunctionOriginalCoordinate transform_coordinate = GetDeviceOriginalCoordinateFunc(coordinate_transform_mode); + CudaFunctionNearestPixel calc_nearest_pixel = GetDeviceNearstPixelFunction(nearest_mode); + fast_divmod div_output_image = (rank > 2) ? output_div_pitches.CpuPtr()[rank - 3] : fast_divmod(gsl::narrow_cast(N)); + int64_t output_height = output_shape.CpuPtr()[rank - 2]; + int64_t output_width = output_shape.CpuPtr()[rank - 1]; + int blocksPerDimsMappingGrid = (int)(ceil(static_cast(output_height + output_width) / 32)); + + switch (upsample_mode) { + case UpsampleMode::NN: + input_shape.CopyToGpu(); + output_shape.CopyToGpu(); + roi_vals.CopyToGpu(); + scales_vals.CopyToGpu(); + input_strides.CopyToGpu(); + output_div_pitches.CopyToGpu(); + _ResizeNearestKernel<<>>( + rank, input_shape.GpuPtr(), output_shape.GpuPtr(), + input_strides.GpuPtr(), output_div_pitches.GpuPtr(), + scales_vals.GpuPtr(), roi_vals.GpuPtr(), + input_data, output_data, N, + extrapolation_enabled, extrapolation_value, + transform_coordinate, calc_nearest_pixel); + return; + case UpsampleMode::LINEAR: + _ResizeBilinearCoordinateMapping<<>>( + input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1], + output_height, output_width, + scales_vals.CpuPtr()[rank - 2], scales_vals.CpuPtr()[rank - 1], + roi_vals.CpuPtr()[rank - 2], roi_vals.CpuPtr()[rank - 2 + rank], + roi_vals.CpuPtr()[rank - 1], roi_vals.CpuPtr()[rank - 1 + rank], + output_height + output_width, extrapolation_enabled, transform_coordinate, + reinterpret_cast(dims_mapping)); + _ResizeBilinearKernel<<>>( + input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1], + output_height, output_width, + output_div_pitches.CpuPtr()[rank - 2], div_output_image, + input_data, output_data, N, extrapolation_value, + reinterpret_cast(dims_mapping)); + return; + case UpsampleMode::CUBIC: + _ResizeCubicCoordinateMapping<<>>( + input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1], + output_height, output_width, + scales_vals.CpuPtr()[rank - 2], scales_vals.CpuPtr()[rank - 1], + roi_vals.CpuPtr()[rank - 2], roi_vals.CpuPtr()[rank - 2 + rank], + roi_vals.CpuPtr()[rank - 1], roi_vals.CpuPtr()[rank - 1 + rank], + output_height + output_width, extrapolation_enabled, + cubic_coeff_a, exclude_outside, transform_coordinate, + reinterpret_cast(dims_mapping)); + _ResizeBiCubicKernel<<>>( + input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1], + output_height, output_width, + output_div_pitches.CpuPtr()[rank - 2], div_output_image, + input_data, output_data, N, extrapolation_value, + reinterpret_cast(dims_mapping)); + // CUDA_CALL(cudaGetLastError()); + return; } } -#define SPECIALIZED_IMPL(T) \ - template void ResizeImpl(const onnxruntime::UpsampleMode upsample_mode, \ - const size_t rank, \ - const int64_t input_dim2, \ - const int64_t* input_pitches, \ - const fast_divmod* output_div_pitches, \ - const float* scales_vals, \ - const T* input_data, \ - T* output_data, \ - const size_t N); +#define SPECIALIZED_IMPL(T) \ + template void ResizeImpl( \ + const UpsampleMode upsample_mode, \ + const int rank, \ + CudaKernel::CudaAsyncBuffer& input_shape, \ + CudaKernel::CudaAsyncBuffer& output_shape, \ + CudaKernel::CudaAsyncBuffer& input_strides, \ + CudaKernel::CudaAsyncBuffer& output_div_pitches, \ + CudaKernel::CudaAsyncBuffer& scales_vals, \ + CudaKernel::CudaAsyncBuffer& roi_vals, \ + const T* input_data, \ + T* output_data, \ + const size_t N, \ + bool extrapolation_enabled, \ + float extrapolation_value, \ + float cubic_coeff_a, \ + bool exclude_outside, \ + ResizeCoordinateTransformationMode coordinate_transform_mode, \ + ResizeNearestMode nearest_mode, \ + void* dims_mapping); SPECIALIZED_IMPL(float) SPECIALIZED_IMPL(double) diff --git a/onnxruntime/core/providers/cuda/tensor/resize_impl.h b/onnxruntime/core/providers/cuda/tensor/resize_impl.h index 7b17707d5b..9248713d4d 100644 --- a/onnxruntime/core/providers/cuda/tensor/resize_impl.h +++ b/onnxruntime/core/providers/cuda/tensor/resize_impl.h @@ -6,20 +6,34 @@ #include "core/providers/cuda/shared_inc/cuda_utils.h" #include "core/common/common.h" #include "core/providers/cpu/tensor/resize.h" +#include "core/providers/cuda/cuda_common.h" namespace onnxruntime { namespace cuda { +size_t CalcResizeBufferSize(const onnxruntime::UpsampleMode upsample_mode, + const std::vector& output_dims); + template -void ResizeImpl(const onnxruntime::UpsampleMode upsample_mode, - const size_t rank, - const int64_t input_dim2, - const int64_t* input_pitches, - const fast_divmod* output_div_pitches, - const float* scales_vals, - const T* input_data, - T* output_data, - const size_t N); +void ResizeImpl( + const onnxruntime::UpsampleMode upsample_mode, + const int rank, + CudaKernel::CudaAsyncBuffer& input_shape, + CudaKernel::CudaAsyncBuffer& output_shape, + CudaKernel::CudaAsyncBuffer& input_strides, + CudaKernel::CudaAsyncBuffer& output_div_pitches, + CudaKernel::CudaAsyncBuffer& scales_vals, + CudaKernel::CudaAsyncBuffer& roi, + const T* input_data, + T* output_data, + const size_t N, + bool extrapolation_enabled, + float extrapolation_value, + float cubic_coeff_a, + bool exclude_outside, + onnxruntime::ResizeCoordinateTransformationMode coordinate_transform_mode, + onnxruntime::ResizeNearestMode nearest_mode, + void* dims_mapping); } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/tensor/upsample.cc b/onnxruntime/core/providers/cuda/tensor/upsample.cc index 86529df81b..68ab2cebde 100644 --- a/onnxruntime/core/providers/cuda/tensor/upsample.cc +++ b/onnxruntime/core/providers/cuda/tensor/upsample.cc @@ -31,34 +31,26 @@ REGISTER_KERNEL_TYPED(int32_t) REGISTER_KERNEL_TYPED(uint8_t) template -Status Upsample::BaseCompute(OpKernelContext* context, const std::vector& scales) const { +Status Upsample::BaseCompute(OpKernelContext* context, + const std::vector& roi, + const std::vector& scales, + const std::vector& output_dims) const { const Tensor* X = context->Input(0); - - ORT_ENFORCE(nullptr != X); const std::vector& X_dims = X->Shape().GetDims(); auto rank = X_dims.size(); + + ORT_ENFORCE(output_dims.size() == rank, "Rank of input and output tensor should be same."); if (rank == 0) - return Status(ONNXRUNTIME, INVALID_ARGUMENT, - is_resize ? "Resize: input tensor cannot be scalar." : "Upsample: input tensor cannot be scalar."); - + return Status(ONNXRUNTIME, INVALID_ARGUMENT, + is_resize_ ? "Resize: input tensor cannot be scalar." : "Upsample: input tensor cannot be scalar."); if (rank != scales.size()) - return Status(ONNXRUNTIME, INVALID_ARGUMENT, - is_resize ? "Resize: input tensor's dimension does not match the scales." : - "Upsample: input tensor's dimension does not match the scales."); + return Status(ONNXRUNTIME, INVALID_ARGUMENT, + is_resize_ ? "Resize: input tensor's dimension does not match the scales." : "Upsample: input tensor's dimension does not match the scales."); + if (roi.size() != 2 * X->Shape().GetDims().size()) + return Status(ONNXRUNTIME, INVALID_ARGUMENT, + "Resize: size of roi array should be 2 * N where N is the rank of input tensor X."); - if (UpsampleMode::LINEAR == mode_ && rank != 4 && rank != 2) { - std::ostringstream oss; - oss << "'Linear' mode only support 2-D inputs ('Bilinear') or 4-D inputs " - "with the corresponding outermost 2 scale values being 1 in the "; - oss << (is_resize ? "Resize operator" : "Upsample operator"); - return Status(ONNXRUNTIME, FAIL, oss.str()); - } - - std::vector Y_dims; - for (std::size_t i = 0; i < rank; i++) { - Y_dims.push_back(static_cast(scales[i] * X_dims[i])); - } - Tensor* Y = context->Output(0, Y_dims); + Tensor* Y = context->Output(0, output_dims); typedef typename ToCudaType::MappedType CudaT; // kernel @@ -66,7 +58,7 @@ Status Upsample::BaseCompute(OpKernelContext* context, const std::vector input_strides(this, rank); gsl::span input_stride_span = input_strides.CpuSpan(); - TensorPitches output_pitches(Y_dims); + TensorPitches output_pitches(output_dims); CudaAsyncBuffer output_div_pitches(this, rank); gsl::span div_strides_span = output_div_pitches.CpuSpan(); @@ -74,24 +66,29 @@ Status Upsample::BaseCompute(OpKernelContext* context, const std::vector(output_pitches[i])); } - input_strides.CopyToGpu(); - output_div_pitches.CopyToGpu(); - size_t output_count = Y->Shape().Size(); - if (is_resize) { + if (is_resize_) { + CudaAsyncBuffer input_shape(this, X_dims); + CudaAsyncBuffer output_shape(this, output_dims); + CudaAsyncBuffer roi_vals(this, roi); CudaAsyncBuffer scales_vals(this, scales); - scales_vals.CopyToGpu(); - ResizeImpl(mode_, - rank, - (UpsampleMode::LINEAR == mode_) ? (rank == 2 ? X_dims[0] : X_dims[2]) : 0, - input_strides.GpuPtr(), - output_div_pitches.GpuPtr(), - scales_vals.GpuPtr(), + + size_t temp_buffer_size = CalcResizeBufferSize(mode_, output_dims); + auto dims_mapping_buffer = GetScratchBuffer(temp_buffer_size); + void* dims_mapping = reinterpret_cast(dims_mapping_buffer.get()); + ResizeImpl(mode_, (int)rank, input_shape, output_shape, + input_strides, output_div_pitches, scales_vals, roi_vals, reinterpret_cast(X->template Data()), reinterpret_cast(Y->template MutableData()), - output_count); + output_count, use_extrapolation_, extrapolation_value_, + cubic_coeff_a_, exclude_outside_, + coordinate_transform_mode_, nearest_mode_, + dims_mapping); } else { + input_strides.CopyToGpu(); + output_div_pitches.CopyToGpu(); + CudaAsyncBuffer scales_div(this, rank); gsl::span scales_div_span = scales_div.CpuSpan(); @@ -102,7 +99,7 @@ Status Upsample::BaseCompute(OpKernelContext* context, const std::vector::BaseCompute(OpKernelContext* context, const std::vector Status Upsample::ComputeInternal(OpKernelContext* context) const { - // Opset 7 - if (OpKernel::Node().InputDefs().size() == 1 || scales_cached_) { - return BaseCompute(context, scales_); + const Tensor* X = context->Input(0); + ORT_ENFORCE(X != nullptr); + + std::vector output_dims(X->Shape().GetDims().size()); + std::vector roi_array(X->Shape().GetDims().size() * 2, 0.0f); + if (!roi_cached_) { + if (need_roi_input_) { + ORT_ENFORCE(roi_input_idx_ > 0, "Invalid roi input index."); + ParseRoiData(context->Input(roi_input_idx_), roi_array); + } + } + const std::vector& roi = roi_cached_ ? roi_ : roi_array; + + if (OpKernel::Node().InputDefs().size() == 1) { + // Compute output shape from scales and input dims + ComputeOutputShape(scales_, X->Shape().GetDims(), output_dims); + return BaseCompute(context, roi, scales_, output_dims); } - // Opset 9 - const Tensor* scales = context->Input(1); + const auto* scales = context->Input(scales_input_idx_); + const auto* sizes = context->Input(sizes_input_idx_); ORT_ENFORCE(scales != nullptr); - int64_t scales_size = scales->Shape().Size(); - std::vector scales_arrary(scales_size); - ParseScalesData(scales, scales_arrary); - return BaseCompute(context, scales_arrary); + + if (scales_cached_) { + ORT_ENFORCE(sizes == nullptr, "Only one of scales or sizes must be provided as input."); + ComputeOutputShape(scales_, X->Shape().GetDims(), output_dims); + return BaseCompute(context, roi, scales_, output_dims); + } + + std::vector scales_array(X->Shape().GetDims().size()); + if (scales != nullptr && scales->Shape().Size() != 0) { + // use scales input data + ORT_ENFORCE(sizes == nullptr, "Only one of scales or sizes must be provided as input."); + ParseScalesData(scales, scales_array); + ComputeOutputShape(scales_array, X->Shape().GetDims(), output_dims); + } else { + // When sizes input is available directly populate it into the output_dims array. + ORT_ENFORCE(sizes != nullptr && sizes->Shape().Size() != 0, + "Either scales or sizes MUST be provided as input."); + ORT_ENFORCE(sizes->Shape().Size() == output_dims.size(), + "Resize: input tensor's rank does not match the output tensor's rank."); + memcpy(output_dims.data(), sizes->template Data(), sizes->Shape().Size() * sizeof(int64_t)); + ParseScalesDataFromOutputSize(output_dims, X->Shape().GetDims(), scales_array); + } + + return BaseCompute(context, roi, scales_array, output_dims); } } // namespace cuda diff --git a/onnxruntime/core/providers/cuda/tensor/upsample.h b/onnxruntime/core/providers/cuda/tensor/upsample.h index 1c098bbaee..3b6818625d 100644 --- a/onnxruntime/core/providers/cuda/tensor/upsample.h +++ b/onnxruntime/core/providers/cuda/tensor/upsample.h @@ -17,7 +17,8 @@ class Upsample : public UpsampleBase, public CudaKernel { } Status ComputeInternal(OpKernelContext* context) const override; - Status BaseCompute(OpKernelContext* context, const std::vector& scales) const; + Status BaseCompute(OpKernelContext* context, const std::vector& roi, const std::vector& scales, + const std::vector& output_dims) const; }; } // namespace cuda diff --git a/onnxruntime/core/providers/cuda/tensor/upsample_impl.h b/onnxruntime/core/providers/cuda/tensor/upsample_impl.h index 151cd9cf18..0ff2000908 100644 --- a/onnxruntime/core/providers/cuda/tensor/upsample_impl.h +++ b/onnxruntime/core/providers/cuda/tensor/upsample_impl.h @@ -21,17 +21,5 @@ void UpampleImpl(const onnxruntime::UpsampleMode upsample_mode, T* output_data, const size_t N); -template -void ResizeImpl( - const onnxruntime::UpsampleMode upsample_mode, - int64_t batch_size, - int64_t num_channels, - int64_t input_height, - int64_t input_width, - float height_scale, - float width_scale, - const T* Xdata, - T* Ydata, - const size_t N); } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/test/providers/cpu/tensor/resize_op_test.cc b/onnxruntime/test/providers/cpu/tensor/resize_op_test.cc index 5076734106..0ae4b93bb1 100644 --- a/onnxruntime/test/providers/cpu/tensor/resize_op_test.cc +++ b/onnxruntime/test/providers/cpu/tensor/resize_op_test.cc @@ -441,7 +441,7 @@ TEST(ResizeOpTest, ResizeOpCubicDownSampleTest) { test.Run(); } -TEST(ResizeOpTest, ResizeOpLineartDownSampleTest_exclude_outside) { +TEST(ResizeOpTest, ResizeOpCubicDownSampleTest_exclude_outside) { OpTester test("Resize", 11); std::vector roi{}; std::vector scales{0.8f, 0.8f};