From 7307e40badbd0f63e45eb66b76cea6123dd989cc Mon Sep 17 00:00:00 2001
From: Cao Yuhang <yhcao6@gmail.com>
Date: Wed, 3 Apr 2019 16:59:02 +0800
Subject: [PATCH] solve deformable_col2im: invalid argument by increasing
 channel

---
 mmdet/ops/dcn/src/deform_conv_cuda_kernel.cu | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/mmdet/ops/dcn/src/deform_conv_cuda_kernel.cu b/mmdet/ops/dcn/src/deform_conv_cuda_kernel.cu
index 80b73bc..fd56016 100644
--- a/mmdet/ops/dcn/src/deform_conv_cuda_kernel.cu
+++ b/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>
-- 
GitLab