CharlesShang · MatthewHowe · Nov 13, 2020 · Nov 13, 2020 · Nov 13, 2020 · CCZUHUALOOKENGLH
diff --git a/DCN/__init__.py b/DCN/__init__.py
@@ -0,0 +1 @@
+from .dcn_v2 import *
diff --git a/dcn_v2.py → DCN/dcn_v2.py b/dcn_v2.py → DCN/dcn_v2.py
diff --git a/src/cpu/dcn_v2_cpu.cpp → DCN/src/cpu/dcn_v2_cpu.cpp b/src/cpu/dcn_v2_cpu.cpp → DCN/src/cpu/dcn_v2_cpu.cpp
@@ -1,5 +1,6 @@
 #include <vector>
 #include "cpu/dcn_v2_im2col_cpu.h"
+#include <iostream>
 
 #include <ATen/ATen.h>
 //#include <ATen/cuda/CUDAContext.h>
@@ -12,8 +13,12 @@
 
 // author: Charles Shang
 // https://github.com/torch/cunn/blob/master/lib/THCUNN/generic/SpatialConvolutionMM.cu
+
 // modified from the CUDA version for CPU use by Daniel K. Suhendro
 
+// edit by: James Bockman and Matthew Howe
+// modified for torch implementation to remove use of deprecated torch access to Blas
+
 at::Tensor
 dcn_v2_cpu_forward(const at::Tensor &input,
                     const at::Tensor &weight,
@@ -60,9 +65,10 @@ dcn_v2_cpu_forward(const at::Tensor &input,
     const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1;
     const int width_out = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1;
 
-    auto ones = at::ones({height_out, width_out}, input.options());
+    // auto ones = at::ones({height_out, width_out}, input.options());
+    auto ones = at::ones({bias.sizes()[0], height_out, width_out}, input.options());
     auto columns = at::empty({channels * kernel_h * kernel_w, 1 * height_out * width_out}, input.options());
-    auto output = at::empty({batch, channels_out, height_out, width_out}, input.options());
+    auto output = at::zeros({batch, channels_out, height_out, width_out}, input.options());
 
     using scalar_t = float;
     for (int b = 0; b < batch; b++)
@@ -71,37 +77,35 @@ dcn_v2_cpu_forward(const at::Tensor &input,
         auto offset_n = offset.select(0, b);
         auto mask_n = mask.select(0, b);
         auto output_n = output.select(0, b);
+        // std::cout << "output_n: " << output_n << "output.select(0,b): " << output.select(0,b) << "\n"; 
 
         // Do Bias first:
         // M,N,K are dims of matrix A and B
         // (see http://docs.nvidia.com/cuda/cublas/#cublas-lt-t-gt-gemm)
         // (N x 1) (1 x M)
-        long m_ = channels_out;
-        long n_ = height_out * width_out;
-        long k_ = 1;
-        THFloatBlas_gemm('t', 'n', n_, m_, k_, 1.0f,
-                         ones.contiguous().data<scalar_t>(), k_,
-                         bias.contiguous().data<scalar_t>(), k_, 0.0f,
-                         output_n.data<scalar_t>(), n_);
-
-        modulated_deformable_im2col_cpu(input_n.data<scalar_t>(),
-                                         offset_n.data<scalar_t>(),
-                                         mask_n.data<scalar_t>(),
+
+        // torch implementation
+        auto ones_T = at::transpose(ones.contiguous(), 2, 0);
+        ones_T = at::mul(ones_T, bias.contiguous());
+        ones_T = at::transpose(ones_T, 2, 0);
+        output_n = at::add(output_n, ones_T);
+
+        modulated_deformable_im2col_cpu(input_n.data_ptr<scalar_t>(),
+                                         offset_n.data_ptr<scalar_t>(),
+                                         mask_n.data_ptr<scalar_t>(),
                                          1, channels, height, width,
                                          height_out, width_out, kernel_h, kernel_w,
                                          pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w,
                                          deformable_group,
-                                         columns.data<scalar_t>());
+                                         columns.data_ptr<scalar_t>());
 
         //(k * m)  x  (m * n)
         // Y = WC
-        long m = channels_out;
-        long n = height_out * width_out;
-        long k = channels * kernel_h * kernel_w;
-        THFloatBlas_gemm('n', 'n', n, m, k, 1.0f,
-                         columns.data<scalar_t>(), n,
-                         weight.data<scalar_t>(), k, 1.0f,
-                         output_n.data<scalar_t>(), n);
+
+        // torch implementation
+        auto weight_flat = weight.view({channels_out, channels * kernel_h * kernel_w});
+        auto product = at::matmul(weight_flat, columns);
+        output.select(0, b) = at::add(output_n, product.view({channels_out, height_out, width_out}));
     }
     return output;
 }
@@ -148,7 +152,7 @@ std::vector<at::Tensor> dcn_v2_cpu_backward(const at::Tensor &input,
     const int width_out = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1;
 
     auto ones = at::ones({height_out, width_out}, input.options());
-    auto columns = at::empty({channels * kernel_h * kernel_w, 1 * height_out * width_out}, input.options());
+    auto columns = at::zeros({channels * kernel_h * kernel_w, 1 * height_out * width_out}, input.options());
     auto output = at::empty({batch, channels_out, height_out, width_out}, input.options());
 
     auto grad_input = at::zeros_like(input);
@@ -169,65 +173,57 @@ std::vector<at::Tensor> dcn_v2_cpu_backward(const at::Tensor &input,
         auto grad_offset_n = grad_offset.select(0, b);
         auto grad_mask_n = grad_mask.select(0, b);
 
-        long m = channels * kernel_h * kernel_w;
-        long n = height_out * width_out;
-        long k = channels_out;
 
-        THFloatBlas_gemm('n', 't', n, m, k, 1.0f,
-                         grad_output_n.data<scalar_t>(), n,
-                         weight.data<scalar_t>(), m, 0.0f,
-                         columns.data<scalar_t>(), n);
+
+        // Torch implementation
+        auto weight_flat = weight.view({channels_out, channels*kernel_h*kernel_w});
+        weight_flat = at::transpose(weight_flat, 1, 0);
+        auto grad_output_n_flat = grad_output_n.view({channels_out, height_out*width_out});
+        columns = at::matmul(weight_flat, grad_output_n_flat);
 
         // gradient w.r.t. input coordinate data
-        modulated_deformable_col2im_coord_cpu(columns.data<scalar_t>(),
-                                               input_n.data<scalar_t>(),
-                                               offset_n.data<scalar_t>(),
-                                               mask_n.data<scalar_t>(),
+        modulated_deformable_col2im_coord_cpu(columns.data_ptr<scalar_t>(),
+                                               input_n.data_ptr<scalar_t>(),
+                                               offset_n.data_ptr<scalar_t>(),
+                                               mask_n.data_ptr<scalar_t>(),
                                                1, channels, height, width,
                                                height_out, width_out, kernel_h, kernel_w,
                                                pad_h, pad_w, stride_h, stride_w,
                                                dilation_h, dilation_w, deformable_group,
-                                               grad_offset_n.data<scalar_t>(),
-                                               grad_mask_n.data<scalar_t>());
+                                               grad_offset_n.data_ptr<scalar_t>(),
+                                               grad_mask_n.data_ptr<scalar_t>());
         // gradient w.r.t. input data
-        modulated_deformable_col2im_cpu(columns.data<scalar_t>(),
-                                         offset_n.data<scalar_t>(),
-                                         mask_n.data<scalar_t>(),
+        modulated_deformable_col2im_cpu(columns.data_ptr<scalar_t>(),
+                                         offset_n.data_ptr<scalar_t>(),
+                                         mask_n.data_ptr<scalar_t>(),
                                          1, channels, height, width,
                                          height_out, width_out, kernel_h, kernel_w,
                                          pad_h, pad_w, stride_h, stride_w,
                                          dilation_h, dilation_w, deformable_group,
-                                         grad_input_n.data<scalar_t>());
+                                         grad_input_n.data_ptr<scalar_t>());
 
         // gradient w.r.t. weight, dWeight should accumulate across the batch and group
-        modulated_deformable_im2col_cpu(input_n.data<scalar_t>(),
-                                         offset_n.data<scalar_t>(),
-                                         mask_n.data<scalar_t>(),
+        modulated_deformable_im2col_cpu(input_n.data_ptr<scalar_t>(),
+                                         offset_n.data_ptr<scalar_t>(),
+                                         mask_n.data_ptr<scalar_t>(),
                                          1, channels, height, width,
                                          height_out, width_out, kernel_h, kernel_w,
                                          pad_h, pad_w, stride_h, stride_w,
                                          dilation_h, dilation_w, deformable_group,
-                                         columns.data<scalar_t>());
-
-        long m_ = channels_out;
-        long n_ = channels * kernel_h * kernel_w;
-        long k_ = height_out * width_out;
-
-        THFloatBlas_gemm('t', 'n', n_, m_, k_, 1.0f,
-                         columns.data<scalar_t>(), k_,
-                         grad_output_n.data<scalar_t>(), k_, 1.0f,
-                         grad_weight.data<scalar_t>(), n_);
-
-        // gradient w.r.t. bias
-        // long m_ = channels_out;
-        // long k__ = height_out * width_out;
-        THFloatBlas_gemv('t', k_, m_, 1.0f,
-                         grad_output_n.data<scalar_t>(), k_,
-                         ones.data<scalar_t>(), 1, 1.0f,
-                         grad_bias.data<scalar_t>(), 1);
+                                         columns.data_ptr<scalar_t>());
+
+        // Torch implementation
+        auto product = at::matmul(grad_output_n_flat, at::transpose(columns, 1, 0));
+        grad_weight = at::add(grad_weight, product.view({channels_out, channels, kernel_h, kernel_w}));
+
+
+        // Torch implementation
+        auto ones_flat = ones.view({height_out*width_out});
+        product = at::matmul(grad_output_n_flat, ones_flat);
+        grad_bias = at::add(grad_bias, product);
     }
 
     return {
         grad_input, grad_offset, grad_mask, grad_weight, grad_bias
     };
-}
+}
diff --git a/src/cpu/dcn_v2_im2col_cpu.cpp → DCN/src/cpu/dcn_v2_im2col_cpu.cpp b/src/cpu/dcn_v2_im2col_cpu.cpp → DCN/src/cpu/dcn_v2_im2col_cpu.cpp
diff --git a/src/cpu/dcn_v2_im2col_cpu.h → DCN/src/cpu/dcn_v2_im2col_cpu.h b/src/cpu/dcn_v2_im2col_cpu.h → DCN/src/cpu/dcn_v2_im2col_cpu.h
diff --git a/src/cpu/dcn_v2_psroi_pooling_cpu.cpp → DCN/src/cpu/dcn_v2_psroi_pooling_cpu.cpp b/src/cpu/dcn_v2_psroi_pooling_cpu.cpp → DCN/src/cpu/dcn_v2_psroi_pooling_cpu.cpp
diff --git a/src/cpu/vision.h → DCN/src/cpu/vision.h b/src/cpu/vision.h → DCN/src/cpu/vision.h
diff --git a/src/cuda/dcn_v2_cuda.cu → DCN/src/cuda/dcn_v2_cuda.cu b/src/cuda/dcn_v2_cuda.cu → DCN/src/cuda/dcn_v2_cuda.cu
@@ -104,7 +104,7 @@ dcn_v2_cuda_forward(const at::Tensor &input,
     const int block = 128;
     const int grid = (batch + block - 1) / block;
 
-    createBatchGemmBuffer<<<grid, block, 0, THCState_getCurrentStream(state)>>>(
+    createBatchGemmBuffer<<<grid, block, 0, c10::cuda::getCurrentCUDAStream()>>>(
         input_b, output_b,
         columns_b, ones_b,
         weight_b, bias_b,
@@ -136,7 +136,7 @@ dcn_v2_cuda_forward(const at::Tensor &input,
                             output_b, n_,
                             batch);
 
-    modulated_deformable_im2col_cuda(THCState_getCurrentStream(state),
+    modulated_deformable_im2col_cuda(c10::cuda::getCurrentCUDAStream(),
                                      input.data<scalar_t>(),
                                      offset.data<scalar_t>(),
                                      mask.data<scalar_t>(),
@@ -276,7 +276,7 @@ std::vector<at::Tensor> dcn_v2_cuda_backward(const at::Tensor &input,
                          columns.data<scalar_t>(), n);
 
         // gradient w.r.t. input coordinate data
-        modulated_deformable_col2im_coord_cuda(THCState_getCurrentStream(state),
+        modulated_deformable_col2im_coord_cuda(c10::cuda::getCurrentCUDAStream(),
                                                columns.data<scalar_t>(),
                                                input_n.data<scalar_t>(),
                                                offset_n.data<scalar_t>(),
@@ -288,7 +288,7 @@ std::vector<at::Tensor> dcn_v2_cuda_backward(const at::Tensor &input,
                                                grad_offset_n.data<scalar_t>(),
                                                grad_mask_n.data<scalar_t>());
         // gradient w.r.t. input data
-        modulated_deformable_col2im_cuda(THCState_getCurrentStream(state),
+        modulated_deformable_col2im_cuda(c10::cuda::getCurrentCUDAStream(),
                                          columns.data<scalar_t>(),
                                          offset_n.data<scalar_t>(),
                                          mask_n.data<scalar_t>(),
@@ -299,7 +299,7 @@ std::vector<at::Tensor> dcn_v2_cuda_backward(const at::Tensor &input,
                                          grad_input_n.data<scalar_t>());
 
         // gradient w.r.t. weight, dWeight should accumulate across the batch and group
-        modulated_deformable_im2col_cuda(THCState_getCurrentStream(state),
+        modulated_deformable_im2col_cuda(c10::cuda::getCurrentCUDAStream(),
                                          input_n.data<scalar_t>(),
                                          offset_n.data<scalar_t>(),
                                          mask_n.data<scalar_t>(),
@@ -321,15 +321,21 @@ std::vector<at::Tensor> dcn_v2_cuda_backward(const at::Tensor &input,
         // gradient w.r.t. bias
         // long m_ = channels_out;
         // long k__ = height_out * width_out;
-        THCudaBlas_Sgemv(state,
-                         't',
-                         k_, m_, 1.0f,
-                         grad_output_n.data<scalar_t>(), k_,
-                         ones.data<scalar_t>(), 1, 1.0f,
-                         grad_bias.data<scalar_t>(), 1);
+        // THCudaBlas_Sgemm(state,
+        //                  't', 'n',
+        //                  k_, m_, 1, 1.0f,
+        //                  grad_output_n.data<scalar_t>(), k_,
+        //                  ones.data<scalar_t>(), 1, 1.0f,
+        //                  grad_bias.data<scalar_t>(), 1);
+        THCudaBlas_Sgemm(state,
+            'N', 'N', 1, m_, k_, 1.0f,
+            ones.data<scalar_t>(), 1,
+            grad_output_n.data<scalar_t>(), k_,
+            1.0f,
+            grad_bias.data<scalar_t>(), 1);
     }
 
     return {
         grad_input, grad_offset, grad_mask, grad_weight, grad_bias
     };
-}
+}
diff --git a/src/cuda/dcn_v2_im2col_cuda.cu → DCN/src/cuda/dcn_v2_im2col_cuda.cu b/src/cuda/dcn_v2_im2col_cuda.cu → DCN/src/cuda/dcn_v2_im2col_cuda.cu
diff --git a/src/cuda/dcn_v2_im2col_cuda.h → DCN/src/cuda/dcn_v2_im2col_cuda.h b/src/cuda/dcn_v2_im2col_cuda.h → DCN/src/cuda/dcn_v2_im2col_cuda.h
diff --git a/src/cuda/dcn_v2_psroi_pooling_cuda.cu → DCN/src/cuda/dcn_v2_psroi_pooling_cuda.cu b/src/cuda/dcn_v2_psroi_pooling_cuda.cu → DCN/src/cuda/dcn_v2_psroi_pooling_cuda.cu
diff --git a/src/cuda/vision.h → DCN/src/cuda/vision.h b/src/cuda/vision.h → DCN/src/cuda/vision.h
diff --git a/src/dcn_v2.h → DCN/src/dcn_v2.h b/src/dcn_v2.h → DCN/src/dcn_v2.h
diff --git a/src/vision.cpp → DCN/src/vision.cpp b/src/vision.cpp → DCN/src/vision.cpp
diff --git a/testcpu.py → DCN/testcpu.py b/testcpu.py → DCN/testcpu.py
diff --git a/testcuda.py → DCN/testcuda.py b/testcuda.py → DCN/testcuda.py
@@ -132,7 +132,7 @@ def check_pooling_zero_offset():
 
 
 def check_gradient_dpooling():
-    input = torch.randn(2, 3, 5, 5).cuda() * 0.01
+    input = torch.randn(2, 3, 5, 5).cuda().float() * 0.01
     N = 4
     batch_inds = torch.randint(2, (N, 1)).cuda().float()
     x = torch.rand((N, 1)).cuda().float() * 15

diff --git a/__init__.py b/__init__.py
diff --git a/setup.py b/setup.py
@@ -17,7 +17,7 @@
 
 def get_extensions():
     this_dir = os.path.dirname(os.path.abspath(__file__))
-    extensions_dir = os.path.join(this_dir, "src")
+    extensions_dir = os.path.join(this_dir, "DCN", "src")
 
     main_file = glob.glob(os.path.join(extensions_dir, "*.cpp"))
     source_cpu = glob.glob(os.path.join(extensions_dir, "cpu", "*.cpp"))
@@ -68,4 +68,4 @@ def get_extensions():
     # install_requires=requirements,
     ext_modules=get_extensions(),
     cmdclass={"build_ext": torch.utils.cpp_extension.BuildExtension},
-)
+)