From eed9dea3ec85379cee02c970b33d64df19b4eabd Mon Sep 17 00:00:00 2001 From: Changming Sun Date: Mon, 22 Mar 2021 14:39:45 -0700 Subject: [PATCH] Make our CUDA code be compatible with the latest VS2019 update (#7062) --- .../cuda/object_detection/roialign_impl.cu | 4 ++-- .../core/providers/cuda/tensor/resize_impl.cu | 14 +++++++------- 2 files changed, 9 insertions(+), 9 deletions(-) diff --git a/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu b/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu index 937007f57b..b4b5171f80 100644 --- a/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu +++ b/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu @@ -130,9 +130,9 @@ __global__ void RoIAlignForward( // We use roi_bin_grid to sample the grid and mimic integral int roi_bin_grid_h = (sampling_ratio > 0) ? sampling_ratio - : ceil(roi_height / pooled_height); // e.g., = 2 + : _Ceil(roi_height / pooled_height); // e.g., = 2 int roi_bin_grid_w = - (sampling_ratio > 0) ? sampling_ratio : ceil(roi_width / pooled_width); + (sampling_ratio > 0) ? sampling_ratio : _Ceil(roi_width / pooled_width); // We do average (integral) pooling inside a bin const T count = roi_bin_grid_h * roi_bin_grid_w; // e.g. = 4 diff --git a/onnxruntime/core/providers/cuda/tensor/resize_impl.cu b/onnxruntime/core/providers/cuda/tensor/resize_impl.cu index 36d138f107..a73a3e7598 100644 --- a/onnxruntime/core/providers/cuda/tensor/resize_impl.cu +++ b/onnxruntime/core/providers/cuda/tensor/resize_impl.cu @@ -10,7 +10,7 @@ using onnxruntime::UpsampleMode; __device__ int NearestPixel_SIMPLE(float x_original, bool is_down_sampling) { if (is_down_sampling) { - return static_cast(ceil(x_original)); + return static_cast(_Ceil(x_original)); } else { return static_cast(x_original); } @@ -18,21 +18,21 @@ __device__ int NearestPixel_SIMPLE(float x_original, bool is_down_sampling) { __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(_Floor(x_original)); } - return static_cast(round(x_original)); + return static_cast(roundf(x_original)); } __device__ int NearestPixel_ROUND_PREFER_CEIL(float x_original, bool) { - return static_cast(round(x_original)); + return static_cast(roundf(x_original)); } __device__ int NearestPixel_FLOOR(float x_original, bool) { - return static_cast(floor(x_original)); + return static_cast(_Floor(x_original)); } __device__ int NearestPixel_CEIL(float x_original, bool) { - return static_cast(ceil(x_original)); + return static_cast(_Ceil(x_original)); } using CudaFunctionNearestPixel = int (*)(float, bool); @@ -535,7 +535,7 @@ __global__ void _ResizeCubicCoordinateMapping( 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)); + 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);