solve deformable_col2im: invalid argument by increasing channel

yhcao6 · yhcao6 · commit 7307e40badbd · 2019-04-03T16:59:02.000+08:00
diff --git 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>

Original file line number	Diff line number	Diff line change
`@@ -73,10 +73,11 @@ using namespace at;`
`73`	`73`	`i += blockDim.x * gridDim.x)`
`74`	`74`
`75`	`75`	`const int CUDA_NUM_THREADS = 1024;`
	`76`	`+const int kMaxGridNum = 65535;`
`76`	`77`
`77`	`78`	`inline int GET_BLOCKS(const int N)`
`78`	`79`	`{`
`79`		`- return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS;`
	`80`	`+ return std::min(kMaxGridNum, (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS);`
`80`	`81`	`}`
`81`	`82`
`82`	`83`	`template <typename scalar_t>`