Repository: rapidsai/cucim (language: python)

Issue:
[BUG] CUDA_ERROR_ILLEGAL_ADDRESS with `morphology.isotropic_closing`
**Describe the bug**
I accidentally landed on a specific combination of conditions which causes illegal address access, resulting in a crash. I believe this has something to do with the time taken to write to memory, whilst another operation is happening. However, the timing required is incredibly inconsistent: sometimes it will happen repeatedly, then start working for a while, and go back to crashing.

**Steps/Code to reproduce bug**
Snippet 1 (MVP of my actual code):
```py
import cupy as cp
from cucim.skimage import morphology
x = cp.random.rand(1000, 1000)  # some large array simulating a 2D image
y = x.copy()
y = y >= 0.1
morphology.isotropic_closing(y, 3)
```
- all at once (in a function): crashes 5/5 times
- wait 10 seconds between `y = y >= 0.1` and `morphology.isotropic_closing(y, 3)`: crashes 5/5 times

Snippet 2:
```py
import cupy as cp
from cucim.skimage import morphology
x = cp.random.rand(1000, 1000)  # some large array simulating a 2D image
y = x.copy()
z = y >= 0.1
morphology.isotropic_closing(z, 3)
```
- doesn't crash

Snippet 3:
```py
import cupy as cp
from cucim.skimage import morphology
x = cp.random.rand(1000, 1000)  # some large array simulating a 2D image
y = x >= 0.1
morphology.isotropic_closing(y, 3)
```
- all at once (in a function): crashes 5/5 times
- wait 10 seconds between `y = x >= 0.1` and `morphology.isotropic_closing(y, 3)`: crashes 3/5 times

**Expected behavior**
All of the above snippets should run fine.

**Environment details (please complete the following information):**
 - Environment location: HPC Cluster
 - Method of cuCIM install: conda, stable v23.02 and nightly v23.04a
 - GPU: Tesla V100-SXM2-32GB
 - CUDA: 11.4


**Additional context**
- This ONLY happens to `morphology.isotropic_closing`, `morphology.isotropic_opening` works fine
- `morphology.binary_closing` (i.e. `morphology.binary_closing(y, morphology.disk(3))`) works fine
-  When digging through the GitHub issues, I also found https://github.com/rapidsai/cucim/pull/421#issuecomment-1330140735, not sure if relevant.


Discussion / hints:
Thanks for reporting the issue. I saw this type of failure inconsistently on CI earlier this week (albeit only for aarch64, not x86_64). I was unable to reproduce it running the test suite locally. I will try to look into it again and see where the issue could be occurring.
I was unable to reproduce with size (1000, 1000) even if running many times, but it did fail when running with size (4000, 4000) for me. The traceback pointed to a specific kernel within the 2D parallel banding algorithm used for the distance transform:
```py
File ~/Dropbox/Quansight/public/NVIDIA/cucim/python/cucim/src/cucim/core/operations/morphology/_pba_2d.py:416, in _pba_2d(arr, sampling, return_distances, return_indices, block_params, check_warp_size, float64_distances, distances, indices)
    409 kernelPropagateInterband(
    410     grid,
    411     block,
    412     (input_arr, margin, size, bandSize1),
    413 )
    414 # kernelUpdateVertical stores output into an intermediate array of
    415 # transposed shape
--> 416 kernelUpdateVertical(
    417     grid,
    418     block,
    419     (input_arr, margin, output, size, bandSize1),
    420 )
    422 # phase 2
    423 block = (block_size, 1, 1)
File ~/src/public/cupy/cupy/_core/raw.pyx:89, in cupy._core.raw.RawKernel.__call__()

File ~/src/public/cupy/cupy/cuda/function.pyx:224, in cupy.cuda.function.Function.__call__()

File ~/src/public/cupy/cupy/cuda/function.pyx:206, in cupy.cuda.function._launch()

File ~/src/public/cupy/cupy_backends/cuda/api/driver.pyx:253, in cupy_backends.cuda.api.driver.launchKernel()

File ~/src/public/cupy/cupy_backends/cuda/api/driver.pyx:60, in cupy_backends.cuda.api.driver.check_status()

CUDADriverError: CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered

```

I would guess the difference in size at which it occurs relates to grid/block size settings and differences between GPU models. It seems perhaps the default setting of these parameters is not working as expected in some cases.

Actually, that specific traceback above may be a false alarm because it was not the same on a repeated run. I can also confirm that I only seem to encounter the issue for `isotropic_closing` and not the other isotropic functions.


Fix patch:
diff --git a/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h b/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h
index 74e5f0255..41fc6698e 100644
--- a/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h
+++ b/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h
@@ -585,7 +585,7 @@ __global__ void kernelColor(pixel_int2_t *input, pixel_int2_t *output, int size)
     last2 = input[TOID(tx, lasty, size)];
 
     if (last2.x == MARKER) {
-        lasty = last2.y;
+        lasty = max(last2.y, 0);
         last2 = input[TOID(tx, lasty, size)];
     }
 
@@ -648,7 +648,7 @@ __global__ void kernelColorWithSpacing(pixel_int2_t *input, pixel_int2_t *output
     last2 = input[TOID(tx, lasty, size)];
 
     if (last2.x == MARKER) {
-        lasty = last2.y;
+        lasty = max(last2.y, 0);
         last2 = input[TOID(tx, lasty, size)];
     }
 


Test patch:
diff --git a/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py b/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py
index edd67bc7e..d0ecdb219 100644
--- a/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py
+++ b/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py
@@ -1,5 +1,3 @@
-import platform
-
 import cupy as cp
 import numpy as np
 import pytest
@@ -12,12 +10,7 @@
 img = color.rgb2gray(cp.asarray(data.astronaut()))
 bw_img = img > 100 / 255.
 
-# TODO: Some tests fail unexpectedly on ARM.
-ON_AARCH64 = platform.machine() == "aarch64"
-ON_AARCH64_REASON = "TODO: Test fails unexpectedly on ARM."
-
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_non_square_image():
     isotropic_res = morphology.isotropic_erosion(bw_img[:100, :200], 3)
     binary_res = img_as_bool(morphology.binary_erosion(
@@ -25,7 +18,6 @@ def test_non_square_image():
     assert_array_equal(isotropic_res, binary_res)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_isotropic_erosion():
     isotropic_res = morphology.isotropic_erosion(bw_img, 3)
     binary_res = img_as_bool(
@@ -52,7 +44,6 @@ def _disk_with_spacing(
     return cp.asarray((X ** 2 + Y ** 2) <= radius ** 2, dtype=dtype)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_isotropic_erosion_spacing():
     isotropic_res = morphology.isotropic_dilation(bw_img, 6, spacing=(1, 2))
     binary_res = img_as_bool(
@@ -63,7 +54,6 @@ def test_isotropic_erosion_spacing():
     assert_array_equal(isotropic_res, binary_res)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_isotropic_dilation():
     isotropic_res = morphology.isotropic_dilation(bw_img, 3)
     binary_res = img_as_bool(
@@ -72,7 +62,6 @@ def test_isotropic_dilation():
     assert_array_equal(isotropic_res, binary_res)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_isotropic_closing():
     isotropic_res = morphology.isotropic_closing(bw_img, 3)
     binary_res = img_as_bool(
@@ -81,7 +70,6 @@ def test_isotropic_closing():
     assert_array_equal(isotropic_res, binary_res)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_isotropic_opening():
     isotropic_res = morphology.isotropic_opening(bw_img, 3)
     binary_res = img_as_bool(
@@ -90,7 +78,6 @@ def test_isotropic_opening():
     assert_array_equal(isotropic_res, binary_res)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 def test_footprint_overflow():
     img = cp.zeros((20, 20), dtype=bool)
     img[2:19, 2:19] = True
@@ -101,7 +88,6 @@ def test_footprint_overflow():
     assert_array_equal(isotropic_res, binary_res)
 
 
-@pytest.mark.xfail(ON_AARCH64, reason=ON_AARCH64_REASON)
 @pytest.mark.parametrize('out_dtype', [bool, cp.uint8, cp.int32])
 def test_out_argument(out_dtype):
     for func in (morphology.isotropic_erosion, morphology.isotropic_dilation,