From 102d01b206247ae53634d1a6868e1749d4397898 Mon Sep 17 00:00:00 2001
From: George Wu <jywu@microsoft.com>
Date: Wed, 29 Jun 2022 17:32:59 -0700
Subject: [PATCH] update roialign cuda impl to onnx opset16 (#12036)

* roialign opset16

* fix

* fix
---
 .../core/providers/cuda/object_detection/roialign.cc     | 1 +
 .../providers/cuda/object_detection/roialign_impl.cu     | 9 ++++++---
 .../core/providers/cuda/object_detection/roialign_impl.h | 1 +
 3 files changed, 8 insertions(+), 3 deletions(-)

diff --git a/onnxruntime/core/providers/cuda/object_detection/roialign.cc b/onnxruntime/core/providers/cuda/object_detection/roialign.cc
index 9c513d39ff..fd9daffd7c 100644
--- a/onnxruntime/core/providers/cuda/object_detection/roialign.cc
+++ b/onnxruntime/core/providers/cuda/object_detection/roialign.cc
@@ -59,6 +59,7 @@ Status RoiAlign<T>::ComputeInternal(OpKernelContext* context) const {
         num_roi_cols,
         reinterpret_cast<typename ToCudaType<T>::MappedType*>(Y.template MutableData<T>()),
         this->mode_ == RoiAlignMode::avg,
+        this->half_pixel_,
         batch_indices_ptr->template Data<int64_t>());
   }
 
diff --git a/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu b/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu
index b4b5171f80..537ad0a8b9 100644
--- a/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu
+++ b/onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu
@@ -94,6 +94,7 @@ __global__ void RoIAlignForward(
     int64_t roi_cols,
     T* top_data,
     const bool is_mode_avg,
+    const bool half_pixel,
     const int64_t* batch_indices_ptr) {
   for (size_t index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; index += blockDim.x * gridDim.x) {
     // (n, c, ph, pw) is an element in the pooled output
@@ -106,9 +107,8 @@ __global__ void RoIAlignForward(
     const T* offset_bottom_rois = bottom_rois + n * roi_cols;
     const auto roi_batch_ind = batch_indices_ptr[n];
 
-    bool continuous_coordinate = false;
     // Do not using rounding; this implementation detail is critical
-    T roi_offset = continuous_coordinate ? T(0.5) : T(0);
+    T roi_offset = half_pixel ? T(0.5) : T(0);
     T roi_start_w = offset_bottom_rois[0] * spatial_scale - roi_offset;
     T roi_start_h = offset_bottom_rois[1] * spatial_scale - roi_offset;
     T roi_end_w = offset_bottom_rois[2] * spatial_scale - roi_offset;
@@ -116,7 +116,7 @@ __global__ void RoIAlignForward(
 
     T roi_width = roi_end_w - roi_start_w;
     T roi_height = roi_end_h - roi_start_h;
-    if (!continuous_coordinate) { // backward compatiblity
+    if (!half_pixel) { // backward compatiblity
       // Force malformed ROIs to be 1x1
       roi_width = max(roi_width, (T)1.);
       roi_height = max(roi_height, (T)1.);
@@ -188,6 +188,7 @@ void RoiAlignImpl(
   int64_t roi_cols,
   T* top_data,
   const bool is_mode_avg,
+  const bool half_pixel,
   const int64_t* batch_indices_ptr) {
     int blocksPerGrid = (int)(ceil(static_cast<float>(nthreads) / GridDim::maxThreadsPerBlock)); 
     RoIAlignForward<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
@@ -204,6 +205,7 @@ void RoiAlignImpl(
       roi_cols,
       top_data,
       is_mode_avg,
+      half_pixel,
       batch_indices_ptr);    
 }
 
@@ -223,6 +225,7 @@ void RoiAlignImpl(
         int64_t roi_cols,                       \
         T* top_data,                            \
         const bool is_mode_avg,                 \
+        const bool half_pixel,                  \
         const int64_t* batch_indices_ptr);
 
 SPECIALIZED_IMPL(float)
diff --git a/onnxruntime/core/providers/cuda/object_detection/roialign_impl.h b/onnxruntime/core/providers/cuda/object_detection/roialign_impl.h
index 4ff9d4cb1d..3fd2f18043 100644
--- a/onnxruntime/core/providers/cuda/object_detection/roialign_impl.h
+++ b/onnxruntime/core/providers/cuda/object_detection/roialign_impl.h
@@ -24,6 +24,7 @@ void RoiAlignImpl(
     int64_t roi_cols,
     T* top_data,
     const bool is_mode_avg,
+    const bool half_pixel,
     const int64_t* batch_indices_ptr);
 
 }  // namespace cuda