Repository: pytorch/vision
Pull request: Fixes crash in deformable convolutions (2598)
Fixes #2598

Description:

~~[x] Adjusted `num_kernels` and `batch_size` according to `kMaxGridNum * CUDA_NUM_THREADS`~~
- [x] Redefined `kMaxGridNum` as max grid according to current CUDA device
- [x] Added test to check the code from issue and compared grads CPU/CUDA

~~This PR adjusts `num_kernels` and `batch_size` to avoid the situation such that `num_kernels` > `kMaxGridNum * CUDA_NUM_THREADS`. If it is not possible to adjust `batch_size` (e.g. batch_size=0), an exception is raised.~~

Problem statement:
ops.deform_conv2d causes CUDA illegal memory access
## 🐛 Bug
I try to test the speed of deformable conv2d. But always encountered memory error.

## To Reproduce

```
$ ipython
Python 3.8.5 (default, Jul 27 2020, 08:42:51) 
Type 'copyright', 'credits' or 'license' for more information
IPython 7.17.0 -- An enhanced Interactive Python. Type '?' for help.

In [1]: import torch
   ...: import torchvision as tv
   ...: weight = torch.randn(9,9,3,3).cuda()
   ...: weight.requires_grad = True
   ...: img = torch.randn(8, 9, 1000, 110).cuda()
   ...: def test():
   ...:     offset = torch.randn(8,18,1000,110).cuda()
   ...:     out = tv.ops.deform_conv2d(img, offset, weight, padding=1)
   ...:     out.mean().backward()
   ...: 

In [2]: import os

In [3]: os.environ["CUDA_LAUNCH_BLOCKING"] = "1"

In [4]: timeit test()
---------------------------------------------------------------------------
RuntimeError                              Traceback (most recent call last)
<ipython-input-4-a1086d7a4706> in <module>
----> 1 get_ipython().run_line_magic('timeit', 'test()')

/usr/lib/python3.8/site-packages/IPython/core/interactiveshell.py in run_line_magic(self, magic_name, line, _stack_depth)
   2324                 kwargs['local_ns'] = self.get_local_scope(stack_depth)
   2325             with self.builtin_trap:
-> 2326                 result = fn(*args, **kwargs)
   2327             return result
   2328 

<decorator-gen-60> in timeit(self, line, cell, local_ns)

/usr/lib/python3.8/site-packages/IPython/core/magic.py in <lambda>(f, *a, **k)
    185     # but it's overkill for just that one bit of state.
    186     def magic_deco(arg):
--> 187         call = lambda f, *a, **k: f(*a, **k)
    188 
    189         if callable(arg):

/usr/lib/python3.8/site-packages/IPython/core/magics/execution.py in timeit(self, line, cell, local_ns)
   1171                     break
   1172 
-> 1173         all_runs = timer.repeat(repeat, number)
   1174         best = min(all_runs) / number
   1175         worst = max(all_runs) / number

/usr/lib/python3.8/timeit.py in repeat(self, repeat, number)
    203         r = []
    204         for i in range(repeat):
--> 205             t = self.timeit(number)
    206             r.append(t)
    207         return r

/usr/lib/python3.8/site-packages/IPython/core/magics/execution.py in timeit(self, number)
    167         gc.disable()
    168         try:
--> 169             timing = self.inner(it, self.timer)
    170         finally:
    171             if gcold:

<magic-timeit> in inner(_it, _timer)

<ipython-input-1-a97200bb984a> in test()
      5 img = torch.randn(8, 9, 1000, 110).cuda()
      6 def test():
----> 7     offset = torch.randn(8,18,1000,110).cuda()
      8     out = tv.ops.deform_conv2d(img, offset, weight, padding=1)
      9     out.mean().backward()

RuntimeError: CUDA error: an illegal memory access was encountered

In [5]: 
```

## Environment
PyTorch version: 1.6.0
Is debug build: False
CUDA used to build PyTorch: 11.0

OS: Arch Linux (x86_64)
GCC version: (GCC) 10.1.0
Clang version: 10.0.1 
CMake version: version 3.18.1

Python version: 3.8 (64-bit runtime)
Is CUDA available: True
CUDA runtime version: 11.0.2
GPU models and configuration: GPU 0: GeForce GTX 1050 Ti
Nvidia driver version: 450.57
cuDNN version: Probably one of the following:
/usr/lib/libcudnn.so.8.0.2
/usr/lib/libcudnn_adv_infer.so.8.0.2
/usr/lib/libcudnn_adv_train.so.8.0.2
/usr/lib/libcudnn_cnn_infer.so.8.0.2
/usr/lib/libcudnn_cnn_train.so.8.0.2
/usr/lib/libcudnn_ops_infer.so.8.0.2
/usr/lib/libcudnn_ops_train.so.8.0.2

Versions of relevant libraries:
[pip3] numpy==1.19.1
[pip3] torch==1.6.0
[pip3] torch-cluster==1.4.5
[pip3] torch-geometric==1.3.2
[pip3] torch-scatter==1.4.0
[pip3] torch-sparse==0.4.3
[pip3] torchvision==0.7.0a0
[conda] Could not collect

Hints from issue discussion:
Hi,

Thanks for the report.

I just tried your snippet on my machine (using PyTorch 1.7.0.dev20200727 and torchvision from source id 0.8.0a0+1aef87d) and it worked fine.

Can you double-check that the version of PyTorch and torchvision that you are using in your Python interpreter are indeed 1.6.0 and 0.7.0 ? The 0.7.0a0 is a bit suspicious, and is maybe an RC candidate for torchvision.
That's weired.
I changed a machine and used a nightly build. But still has the same problem.
```
Collecting environment information...
PyTorch version: 1.7.0.dev20200819
Is debug build: False
CUDA used to build PyTorch: 10.2
OS: Arch Linux (x86_64)
GCC version: (GCC) 10.1.0
Clang version: 10.0.1 
CMake version: version 3.18.1
Python version: 3.8 (64-bit runtime)
Is CUDA available: True
CUDA runtime version: 11.0.194
GPU models and configuration: GPU 0: GeForce GTX 1080 Ti
Nvidia driver version: 450.57
cuDNN version: Probably one of the following:
/usr/lib/libcudnn.so.8.0.2
/usr/lib/libcudnn_adv_infer.so.8.0.2
/usr/lib/libcudnn_adv_train.so.8.0.2
/usr/lib/libcudnn_cnn_infer.so.8.0.2
/usr/lib/libcudnn_cnn_train.so.8.0.2
/usr/lib/libcudnn_ops_infer.so.8.0.2
/usr/lib/libcudnn_ops_train.so.8.0.2
Versions of relevant libraries:
[pip3] numpy==1.19.1
[pip3] torch==1.7.0.dev20200819
[pip3] torch-cluster==1.4.5
[pip3] torch-geometric==1.3.2
[pip3] torch-scatter==1.4.0
[pip3] torch-sparse==0.4.3
[pip3] torchvision==0.8.0.dev20200820
[conda] Could not collect
```
However, if i set the Batch size of test data to 1, everything is fine.
This is what I tried. The only difference was that I set `requires_grad` inside tensor creation (this was an experiment on CPU, I also tried one on CUDA)

```python
In [1]: import torch

In [2]: import torchvision as tv

In [3]: weight = torch.randn(9,9,3,3, requires_grad=True)

In [4]: img = torch.randn(8, 9, 20, 30)

In [5]: def test():
   ...:     offset = torch.randn(8,18,20,30)
   ...:     out = tv.ops.deform_conv2d(img, offset, weight, padding=1)
   ...:     out.mean().backward()
   ...:

In [6]: test()

In [7]: test()

In [8]: test()

In [9]: test()

In [10]: test()

In [11]: test()
```

@vfdev-5 can you try seeing if you can reproduce this issue?
@fmassa @godspeed1989 I could reproduce the exception on CUDA by calling twice `test()`
```python
import torch
import torchvision as tv

print(torch.__version__, tv.__version__, torch.version.cuda)

weight = torch.randn(9,9,3,3).cuda()
weight.requires_grad = True
img = torch.randn(8, 9, 1000, 110).cuda()

def test():
    offset = torch.randn(8,18, 1000, 110).cuda()
    out = tv.ops.deform_conv2d(img, offset, weight, padding=1)
    out.mean().backward()
    
    
import os
os.environ["CUDA_LAUNCH_BLOCKING"] = "1"


test()
test()

> 1.6.0 0.7.0 10.2
> RuntimeError: CUDA error: an illegal memory access was encountered
```

Seems like error is input size dependent (no error with 110x110 input) and if we comment out `out.mean().backward()` code runs without any issue event with `timeit`.


**EDIT**: seems like problem is related to the number of image channels = 9. If tested with smaller number of channels (=8) there is no issue.

<summary>
Here is error stack
<details>

```
RuntimeError: CUDA error: an illegal memory access was encountered
Exception raised from launch_vectorized_kernel at /pytorch/aten/src/ATen/native/cuda/CUDALoops.cuh:146 (most recent call first):
frame #0: c10::Error::Error(c10::SourceLocation, std::string) + 0x42 (0x7f529c78f1e2 in /opt/conda/lib/python3.7/site-packages/torch/lib/libc10.so)
frame #1: void at::native::gpu_kernel_impl<__nv_hdl_wrapper_t<false, false, __nv_dl_tag<void (*)(at::TensorIterator&, c10::Scalar), &at::native::fill_kernel_cuda, 4u>, float (), float> >(at::TensorIterator&, __nv_hdl_wrapper_t<false, false, __nv_dl_tag<void (*)(at::TensorIterator&, c10::Scalar), &at::native::fill_kernel_cuda, 4u>, float (), float> const&) + 0x482 (0x7f52471be0d2 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cuda.so)
frame #2: void at::native::gpu_kernel<__nv_hdl_wrapper_t<false, false, __nv_dl_tag<void (*)(at::TensorIterator&, c10::Scalar), &at::native::fill_kernel_cuda, 4u>, float (), float> >(at::TensorIterator&, __nv_hdl_wrapper_t<false, false, __nv_dl_tag<void (*)(at::TensorIterator&, c10::Scalar), &at::native::fill_kernel_cuda, 4u>, float (), float> const&) + 0x11b (0x7f52471bf72b in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cuda.so)
frame #3: <unknown function> + 0x1f77748 (0x7f52471b6748 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cuda.so)
frame #4: at::native::fill_kernel_cuda(at::TensorIterator&, c10::Scalar) + 0x1a (0x7f52471b742a in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cuda.so)
frame #5: <unknown function> + 0xc89f10 (0x7f5280ac8f10 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #6: <unknown function> + 0x129342a (0x7f52810d242a in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #7: at::native::zero_(at::Tensor&) + 0x29 (0x7f5280ac5ce9 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #8: <unknown function> + 0xf24e4d (0x7f5246163e4d in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cuda.so)
frame #9: <unknown function> + 0x2edab20 (0x7f5282d19b20 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #10: at::native::zeros_like(at::Tensor const&, c10::TensorOptions const&, c10::optional<c10::MemoryFormat>) + 0x4e (0x7f5280cab6fe in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #11: <unknown function> + 0x128b434 (0x7f52810ca434 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #12: <unknown function> + 0x2eea7d1 (0x7f5282d297d1 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #13: <unknown function> + 0x12c3863 (0x7f5281102863 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #14: at::zeros_like(at::Tensor const&, c10::TensorOptions const&, c10::optional<c10::MemoryFormat>) + 0x101 (0x7f5281016dd1 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #15: <unknown function> + 0x998f4 (0x7f523cd5d8f4 in /opt/conda/lib/python3.7/site-packages/torchvision/_C.so)
frame #16: DeformConv2d_backward_cuda(at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::pair<int, int>, std::pair<int, int>, std::pair<int, int>, int, int) + 0x229 (0x7f523cd5e68c in /opt/conda/lib/python3.7/site-packages/torchvision/_C.so)
frame #17: DeformConv2d_backward(at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::pair<int, int> const&, std::pair<int, int> const&, std::pair<int, int> const&, int, int) + 0xc9 (0x7f523ccf4339 in /opt/conda/lib/python3.7/site-packages/torchvision/_C.so)
frame #18: DeformConv2dFunction::backward(torch::autograd::AutogradContext*, std::vector<at::Tensor, std::allocator<at::Tensor> >) + 0x533 (0x7f523cd05443 in /opt/conda/lib/python3.7/site-packages/torchvision/_C.so)
frame #19: torch::autograd::CppNode<DeformConv2dFunction>::apply(std::vector<at::Tensor, std::allocator<at::Tensor> >&&) + 0x154 (0x7f523cd102c4 in /opt/conda/lib/python3.7/site-packages/torchvision/_C.so)
frame #20: <unknown function> + 0x3375bb7 (0x7f52831b4bb7 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #21: torch::autograd::Engine::evaluate_function(std::shared_ptr<torch::autograd::GraphTask>&, torch::autograd::Node*, torch::autograd::InputBuffer&, std::shared_ptr<torch::autograd::ReadyQueue> const&) + 0x1400 (0x7f52831b0400 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #22: torch::autograd::Engine::thread_main(std::shared_ptr<torch::autograd::GraphTask> const&) + 0x451 (0x7f52831b0fa1 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #23: torch::autograd::Engine::thread_init(int, std::shared_ptr<torch::autograd::ReadyQueue> const&, bool) + 0x89 (0x7f52831a9119 in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_cpu.so)
frame #24: torch::autograd::python::PythonEngine::thread_init(int, std::shared_ptr<torch::autograd::ReadyQueue> const&, bool) + 0x4a (0x7f529d541dea in /opt/conda/lib/python3.7/site-packages/torch/lib/libtorch_python.so)
frame #25: <unknown function> + 0xc819d (0x7f52aefa419d in /opt/conda/lib/python3.7/site-packages/zmq/backend/cython/../../../../.././libstdc++.so.6)
frame #26: <unknown function> + 0x76db (0x7f52b1e746db in /lib/x86_64-linux-gnu/libpthread.so.0)
frame #27: clone + 0x3f (0x7f52b1b9d88f in /lib/x86_64-linux-gnu/libc.so.6)
```

</details>
</summary>



@vfdev-5 great that you managed to reproduce it.

Does this error also happen on the CPU btw?
@fmassa looks like CPU implementation is fine.
Ok, interesting. Maybe the issue is that we are using `int` to hold some pointers to the data in the CUDA kernel, and the image that we are passing overflows int32?
Here is the error when debugged with `cuda-gdb`:
```
> DeformConv2d_backward_cuda
>> deform_conv_backward_input_cuda
>>> compute_grad_input
>>>> deformable_im2col
>>>>> deformable_im2col_gpu_kernel

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x5585aae7e818

Thread 5 "python" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 18, block (339,0,0), thread (352,0,0), device 0, sm 0, warp 8, lane 0]
0x00005585aae7e968 in void deformable_col2im_gpu_kernel<float>(int, float const*, float const*, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, f
loat*)<<<(65535,1,1),(1024,1,1)>>> ()
```

Probably, this is due to requested number of kernels: `71280000` which is translated to number of blocks=69610 (71280000/1024) and truncated by 65535 max possible number of blocks. But we still would like to run `num_kernels` of 71280000 :

https://github.com/pytorch/vision/blob/baa592b215804927e28638f6a7f3318cbc411d49/torchvision/csrc/cuda/DeformConv_cuda.cu#L510-L520

Nice debugging! ~I would follow a similar pattern to [what we do in `im2col` in PyTorch](https://github.com/pytorch/pytorch/blob/c2511bdfa4bef89e6c328a161d3a4d3466261ff7/aten/src/ATen/native/cuda/im2col.cuh#L93-L95)~ Forget that, this is incorrect.

We will need to change a bit more the code to make it work I'm afraid. ccing @ngimel if she has any ideas on how to easily fix this.

Commit message: - Redefined kMaxGridNum as max grid according to current CUDA device
- Added test to check the code from issue and compared grads CPU/CUDA

Patch to review:
diff --git a/test/test_ops.py b/test/test_ops.py
index 564d5d54559..2b836371238 100644
--- a/test/test_ops.py
+++ b/test/test_ops.py
@@ -553,6 +553,35 @@ def script_func(x_, offset_, weight_, bias_, stride_, pad_, dilation_):
         gradcheck(lambda z, off, wei, bi: script_func(z, off, wei, bi, stride, padding, dilation),
                   (x, offset, weight, bias), nondet_tol=1e-5)
 
+        # Test from https://github.com/pytorch/vision/issues/2598
+        # Run on CUDA only
+        if "cuda" in device.type:
+            # compare grads computed on CUDA with grads computed on CPU
+            true_cpu_grads = None
+
+            init_weight = torch.randn(9, 9, 3, 3, requires_grad=True)
+            img = torch.randn(8, 9, 1000, 110)
+            offset = torch.rand(8, 2 * 3 * 3, 1000, 110)
+
+            if not contiguous:
+                img = img.permute(0, 1, 3, 2).contiguous().permute(0, 1, 3, 2)
+                offset = offset.permute(1, 3, 0, 2).contiguous().permute(2, 0, 3, 1)
+                weight = init_weight.permute(3, 2, 0, 1).contiguous().permute(2, 3, 1, 0)
+            else:
+                weight = init_weight
+
+            for d in ["cpu", "cuda"]:
+
+                out = ops.deform_conv2d(img.to(d), offset.to(d), weight.to(d), padding=1)
+                out.mean().backward()
+                if true_cpu_grads is None:
+                    true_cpu_grads = init_weight.grad
+                    self.assertTrue(true_cpu_grads is not None)
+                else:
+                    self.assertTrue(init_weight.grad is not None)
+                    res_grads = init_weight.grad.to("cpu")
+                    self.assertTrue(true_cpu_grads.allclose(res_grads))
+
 
 class FrozenBNTester(unittest.TestCase):
     def test_frozenbatchnorm2d_repr(self):
diff --git a/torchvision/csrc/cuda/DeformConv_cuda.cu b/torchvision/csrc/cuda/DeformConv_cuda.cu
index 3bf5145602e..49052f9024e 100644
--- a/torchvision/csrc/cuda/DeformConv_cuda.cu
+++ b/torchvision/csrc/cuda/DeformConv_cuda.cu
@@ -80,12 +80,11 @@
 
 using namespace at;
 
-const int CUDA_NUM_THREADS = 1024;
-const int kMaxGridNum = 65535;
-
+const unsigned int CUDA_NUM_THREADS = 1024;
+const unsigned int kMaxGridNum = at::cuda::getCurrentDeviceProperties()->maxGridSize[0];
 const int kMaxParallelImgs = 32;
 
-inline int GET_BLOCKS(const int N) {
+inline unsigned int GET_BLOCKS(const unsigned int N) {
   return std::min(kMaxGridNum, (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS);
 }
 


Code review comments:
diff_hunk: @@ -80,12 +80,11 @@
 
 using namespace at;
 
-const int CUDA_NUM_THREADS = 1024;
-const int kMaxGridNum = 65535;
-
+const unsigned int CUDA_NUM_THREADS = 1024;
+const unsigned int kMaxGridNum = at::cuda::getCurrentDeviceProperties()->maxGridSize[0];
path: torchvision/csrc/cuda/DeformConv_cuda.cu
text: @user1:
For my understanding, can we make this a `const` in the outer namespace? Doesn't it depend on the machine it's being run on, and thus could change from the machine it was compiled to the machine it's being run on?

@author:
Yes, I agree that this variable should be assigned at runtime. It is very unlikely, taking into account how implemented [`at::cuda::getCurrentDeviceProperties()`](https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/cuda/CUDAContext.cpp#L41), that `kMaxGridNum` is evaluated at compilation time.

@author:
Yes, you are right. I'll change it to dynamic evaluation