Encapsulate and standardize roi_align

test/tracing/frcnn/test_frcnn_tracing.cpp

@@ -1,7 +1,7 @@
 #include <ATen/ATen.h>
 #include <torch/script.h>
 #include <torch/torch.h>
-#include <torchvision/ROIAlign.h>
+#include <torchvision/roi_align.h>
 #include <torchvision/cpu/vision_cpu.h>
 #include <torchvision/nms.h>
 

torchvision/csrc/cpu/ROIAlign_cpu.cpp → torchvision/csrc/cpu/roi_align_kernel.cpp

@@ -1,5 +1,6 @@
-#include <ATen/TensorUtils.h>
-#include "vision_cpu.h"
+#include "roi_align_kernel.h"
+
+namespace {
 
 // implementation taken from Caffe2
 template <typename T>
@@ -111,7 +112,7 @@ void pre_calc_for_bilinear_interpolate(
 }
 
 template <typename T>
-void ROIAlignForward(
+void roi_align_forward_kernel_impl(
     int nthreads,
     const T* input,
     const T& spatial_scale,
@@ -277,7 +278,7 @@ inline void add(T* address, const T& val) {
 }
 
 template <typename T>
-void ROIAlignBackward(
+void roi_align_backward_kernel_impl(
     int nthreads,
     const T* grad_output,
     const T& spatial_scale,
@@ -382,9 +383,11 @@ void ROIAlignBackward(
       } // ix
     } // iy
   } // for
-} // ROIAlignBackward
+}
+
+} // namespace
 
-at::Tensor ROIAlign_forward_cpu(
+at::Tensor roi_align_forward_cpu(
     const at::Tensor& input,
     const at::Tensor& rois,
     double spatial_scale,
@@ -398,7 +401,7 @@ at::Tensor ROIAlign_forward_cpu(
 
   at::TensorArg input_t{input, "input", 1}, rois_t{rois, "rois", 2};
 
-  at::CheckedFrom c = "ROIAlign_forward_cpu";
+  at::CheckedFrom c = "roi_align_forward_cpu";
   at::checkAllSameType(c, {input_t, rois_t});
 
   auto num_rois = rois.size(0);
@@ -416,8 +419,8 @@ at::Tensor ROIAlign_forward_cpu(
 
   auto input_ = input.contiguous(), rois_ = rois.contiguous();
   AT_DISPATCH_FLOATING_TYPES_AND_HALF(
-      input.scalar_type(), "ROIAlign_forward", [&] {
-        ROIAlignForward<scalar_t>(
+      input.scalar_type(), "roi_align_forward", [&] {
+        roi_align_forward_kernel_impl<scalar_t>(
             output_size,
             input_.data_ptr<scalar_t>(),
             spatial_scale,
@@ -434,7 +437,7 @@ at::Tensor ROIAlign_forward_cpu(
   return output;
 }
 
-at::Tensor ROIAlign_backward_cpu(
+at::Tensor roi_align_backward_cpu(
     const at::Tensor& grad,
     const at::Tensor& rois,
     double spatial_scale,
@@ -451,7 +454,7 @@ at::Tensor ROIAlign_backward_cpu(
 
   at::TensorArg grad_t{grad, "grad", 1}, rois_t{rois, "rois", 2};
 
-  at::CheckedFrom c = "ROIAlign_backward_cpu";
+  at::CheckedFrom c = "roi_align_backward_cpu";
   at::checkAllSameType(c, {grad_t, rois_t});
 
   at::Tensor grad_input =
@@ -470,8 +473,8 @@ at::Tensor ROIAlign_backward_cpu(
 
   auto rois_ = rois.contiguous();
   AT_DISPATCH_FLOATING_TYPES_AND_HALF(
-      grad.scalar_type(), "ROIAlign_forward", [&] {
-        ROIAlignBackward<scalar_t>(
+      grad.scalar_type(), "roi_align_forward", [&] {
+        roi_align_backward_kernel_impl<scalar_t>(
             grad.numel(),
             grad.data_ptr<scalar_t>(),
             spatial_scale,

torchvision/csrc/cpu/roi_align_kernel.h

@@ -0,0 +1,26 @@
+#pragma once
+
+#include <ATen/ATen.h>
+#include "../macros.h"
+
+VISION_API at::Tensor roi_align_forward_cpu(
+    const at::Tensor& input,
+    const at::Tensor& rois,
+    double spatial_scale,
+    int64_t pooled_height,
+    int64_t pooled_width,
+    int64_t sampling_ratio,
+    bool aligned);
+
+VISION_API at::Tensor roi_align_backward_cpu(
+    const at::Tensor& grad,
+    const at::Tensor& rois,
+    double spatial_scale,
+    int64_t pooled_height,
+    int64_t pooled_width,
+    int64_t batch_size,
+    int64_t channels,
+    int64_t height,
+    int64_t width,
+    int64_t sampling_ratio,
+    bool aligned);

torchvision/csrc/cpu/vision_cpu.h

@@ -4,28 +4,6 @@
 
 // TODO: Delete this file once all the methods are gone
 
-VISION_API at::Tensor ROIAlign_forward_cpu(
-    const at::Tensor& input,
-    const at::Tensor& rois,
-    double spatial_scale,
-    int64_t pooled_height,
-    int64_t pooled_width,
-    int64_t sampling_ratio,
-    bool aligned);
-
-VISION_API at::Tensor ROIAlign_backward_cpu(
-    const at::Tensor& grad,
-    const at::Tensor& rois,
-    double spatial_scale,
-    int64_t pooled_height,
-    int64_t pooled_width,
-    int64_t batch_size,
-    int64_t channels,
-    int64_t height,
-    int64_t width,
-    int64_t sampling_ratio,
-    bool aligned);
-
 VISION_API std::tuple<at::Tensor, at::Tensor> ROIPool_forward_cpu(
     const at::Tensor& input,
     const at::Tensor& rois,

torchvision/csrc/cuda/ROIAlign_cuda.cu → torchvision/csrc/cuda/roi_align_kernel.cu

@@ -1,10 +1,11 @@
-#include <ATen/ATen.h>
-#include <ATen/TensorUtils.h>
 #include <ATen/cuda/CUDAContext.h>
 #include <c10/cuda/CUDAGuard.h>
 #include <THC/THCAtomics.cuh>
 
 #include "cuda_helpers.h"
+#include "roi_align_kernel.h"
+
+namespace {
 
 template <typename T>
 __device__ T bilinear_interpolate(
@@ -61,7 +62,7 @@ __device__ T bilinear_interpolate(
 }
 
 template <typename T>
-__global__ void RoIAlignForward(
+__global__ void roi_align_forward_kernel_impl(
     int nthreads,
     const T* input,
     const T spatial_scale,
@@ -197,7 +198,7 @@ __device__ void bilinear_interpolate_gradient(
 }
 
 template <typename T>
-__global__ void RoIAlignBackward(
+__global__ void roi_align_backward_kernel_impl(
     int nthreads,
     const T* grad_output,
     const T spatial_scale,
@@ -308,9 +309,11 @@ __global__ void RoIAlignBackward(
       } // ix
     } // iy
   } // CUDA_1D_KERNEL_LOOP
-} // RoIAlignBackward
+}
+
+} // namespace
 
-at::Tensor ROIAlign_forward_cuda(
+at::Tensor roi_align_forward_cuda(
     const at::Tensor& input,
     const at::Tensor& rois,
     double spatial_scale,
@@ -325,7 +328,7 @@ at::Tensor ROIAlign_forward_cuda(
 
   at::TensorArg input_t{input, "input", 1}, rois_t{rois, "rois", 2};
 
-  at::CheckedFrom c = "ROIAlign_forward_cuda";
+  at::CheckedFrom c = "roi_align_forward_cuda";
   at::checkAllSameGPU(c, {input_t, rois_t});
   at::checkAllSameType(c, {input_t, rois_t});
 
@@ -354,8 +357,8 @@ at::Tensor ROIAlign_forward_cuda(
 
   auto input_ = input.contiguous(),
        rois_ = rois.contiguous();
-  AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.scalar_type(), "ROIAlign_forward", [&] {
-    RoIAlignForward<scalar_t><<<grid, block, 0, stream>>>(
+  AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.scalar_type(), "roi_align_forward", [&] {
+    roi_align_forward_kernel_impl<scalar_t><<<grid, block, 0, stream>>>(
         output_size,
         input_.data_ptr<scalar_t>(),
         spatial_scale,
@@ -373,7 +376,7 @@ at::Tensor ROIAlign_forward_cuda(
   return output;
 }
 
-at::Tensor ROIAlign_backward_cuda(
+at::Tensor roi_align_backward_cuda(
     const at::Tensor& grad,
     const at::Tensor& rois,
     double spatial_scale,
@@ -390,7 +393,7 @@ at::Tensor ROIAlign_backward_cuda(
 
   at::TensorArg grad_t{grad, "grad", 1}, rois_t{rois, "rois", 2};
 
-  at::CheckedFrom c = "ROIAlign_backward_cuda";
+  at::CheckedFrom c = "roi_align_backward_cuda";
   at::checkAllSameGPU(c, {grad_t, rois_t});
   at::checkAllSameType(c, {grad_t, rois_t});
 
@@ -418,8 +421,8 @@ at::Tensor ROIAlign_backward_cuda(
   int w_stride = grad.stride(3);
 
   auto rois_ = rois.contiguous();
-  AT_DISPATCH_FLOATING_TYPES_AND_HALF(grad.scalar_type(), "ROIAlign_backward", [&] {
-    RoIAlignBackward<scalar_t><<<grid, block, 0, stream>>>(
+  AT_DISPATCH_FLOATING_TYPES_AND_HALF(grad.scalar_type(), "roi_align_backward", [&] {
+    roi_align_backward_kernel_impl<scalar_t><<<grid, block, 0, stream>>>(
         grad.numel(),
         grad.data_ptr<scalar_t>(),
         spatial_scale,

torchvision/csrc/cuda/roi_align_kernel.h

@@ -0,0 +1,26 @@
+#pragma once
+
+#include <ATen/ATen.h>
+#include "../macros.h"
+
+VISION_API at::Tensor roi_align_forward_cuda(
+    const at::Tensor& input,
+    const at::Tensor& rois,
+    double spatial_scale,
+    int64_t pooled_height,
+    int64_t pooled_width,
+    int64_t sampling_ratio,
+    bool aligned);
+
+VISION_API at::Tensor roi_align_backward_cuda(
+    const at::Tensor& grad,
+    const at::Tensor& rois,
+    double spatial_scale,
+    int64_t pooled_height,
+    int64_t pooled_width,
+    int64_t batch_size,
+    int64_t channels,
+    int64_t height,
+    int64_t width,
+    int64_t sampling_ratio,
+    bool aligned);

torchvision/csrc/cuda/vision_cuda.h

@@ -4,28 +4,6 @@
 
 // TODO: Delete this file once all the methods are gone
 
-VISION_API at::Tensor ROIAlign_forward_cuda(
-    const at::Tensor& input,
-    const at::Tensor& rois,
-    double spatial_scale,
-    int64_t pooled_height,
-    int64_t pooled_width,
-    int64_t sampling_ratio,
-    bool aligned);
-
-VISION_API at::Tensor ROIAlign_backward_cuda(
-    const at::Tensor& grad,
-    const at::Tensor& rois,
-    double spatial_scale,
-    int64_t pooled_height,
-    int64_t pooled_width,
-    int64_t batch_size,
-    int64_t channels,
-    int64_t height,
-    int64_t width,
-    int64_t sampling_ratio,
-    bool aligned);
-
 VISION_API std::tuple<at::Tensor, at::Tensor> ROIPool_forward_cuda(
     const at::Tensor& input,
     const at::Tensor& rois,

torchvision/csrc/ROIAlign.h → torchvision/csrc/roi_align.cpp

@@ -1,19 +1,10 @@
-#pragma once
+#include "roi_align.h"
+#include <torch/extension.h>
 
-#include "cpu/vision_cpu.h"
-
-#ifdef WITH_CUDA
-#include "autocast.h"
-#include "cuda/vision_cuda.h"
-#endif
-#ifdef WITH_HIP
-#include "autocast.h"
-#include "hip/vision_cuda.h"
+#if defined(WITH_CUDA) || defined(WITH_HIP)
+#include <ATen/autocast_mode.h>
 #endif
 
-// TODO: put this stuff in torchvision namespace
-
-// roi_align dispatch nexus
 at::Tensor roi_align(
     const at::Tensor& input, // Input feature map.
     const at::Tensor& rois, // List of ROIs to pool over.
@@ -39,7 +30,7 @@ at::Tensor roi_align(
 }
 
 #if defined(WITH_CUDA) || defined(WITH_HIP)
-at::Tensor ROIAlign_autocast(
+at::Tensor roi_align_autocast(
     const at::Tensor& input,
     const at::Tensor& rois,
     double spatial_scale,
@@ -90,6 +81,8 @@ at::Tensor _roi_align_backward(
       aligned);
 }
 
+namespace {
+
 class ROIAlignFunction : public torch::autograd::Function<ROIAlignFunction> {
  public:
   static torch::autograd::variable_list forward(
@@ -189,7 +182,9 @@ class ROIAlignBackwardFunction
   }
 };
 
-at::Tensor ROIAlign_autograd(
+} // namespace
+
+at::Tensor roi_align_autograd(
     const at::Tensor& input,
     const at::Tensor& rois,
     double spatial_scale,
@@ -207,7 +202,7 @@ at::Tensor ROIAlign_autograd(
       aligned)[0];
 }
 
-at::Tensor ROIAlign_backward_autograd(
+at::Tensor roi_align_backward_autograd(
     const at::Tensor& grad,
     const at::Tensor& rois,
     double spatial_scale,