Merge pull request #464 from yhcao6/dcn_cpp_extension

solve deformable_col2im: invalid argument by increasing channel

mmdet/ops/dcn/src/deform_conv_cuda_kernel.cu

@@ -73,10 +73,11 @@ using namespace at;
        i += blockDim.x * gridDim.x)
 
 const int CUDA_NUM_THREADS = 1024;
+const int kMaxGridNum = 65535;
 
 inline int GET_BLOCKS(const int N)
 {
-  return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS;
+  return std::min(kMaxGridNum, (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS);
 }
 
 template <typename scalar_t>