Local maxima detection with cupy and friends

Hi python / CUDA folks,

I’m searching for a (local) maxima finder for 3D data that can make use of GPUs. Ideally it should be using dask-image, cupy or any other compatible sibling. I found cupy.argmax and cusignal.peak_finding.peak_finding. argrelmax but both have an axis parameter suggesting that these only work in one dimension…

And hint / link is appreciated! Is there by chance any spot-detection + region growing tutorial for CUDA based python libraries by chance?

Thanks!
Robert

1 Like

Hi Robert,

cupyimg has a CUDA implementation of peak_local_max:

2 Likes

Thanks @VolkerH , that was a great hint!

I just tried and receive an error. In short:

res = peak_local_max(blurred, min_distance=1)

results in

catastrophic error: cannot open source file "type_traits"

1 catastrophic error detected in the compilation of [...]_2.cubin.cu.
Compilation terminated.

I’ve had another issue of similar kind today

At the moment I run on CUDA 10.2, cupy 0.8.5, Windows 10, NVidia 2080 Ti.

The issue is not of high priority. I was just exploring what’s possible with CUDA. I may continue that exploration later when things are more mature. But thanks for the hint anyway!

Cheers,
Robert

Full code:

import cupy
import numpy as np
from cupyx.scipy.ndimage import gaussian_filter as cupy_gaussian_filter
image = np.random.random((200,200,200))

from cupyimg.skimage.feature.peak import peak_local_max

cuda_image = cupy.asarray(image)
cuda_blurred = cupy_gaussian_filter(cuda_image, sigma=20)

res = peak_local_max(cuda_blurred, min_distance=1)

result1 = cupy.asnumpy(res)

Full error message:

---------------------------------------------------------------------------
NVRTCError                                Traceback (most recent call last)
C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\cuda\compiler.py in compile(self, options, log_stream)
    548                     nvrtc.addAddNameExpression(self.ptr, ker)
--> 549             nvrtc.compileProgram(self.ptr, options)
    550             mapping = None

cupy_backends\cuda\libs\nvrtc.pyx in cupy_backends.cuda.libs.nvrtc.compileProgram()

cupy_backends\cuda\libs\nvrtc.pyx in cupy_backends.cuda.libs.nvrtc.compileProgram()

cupy_backends\cuda\libs\nvrtc.pyx in cupy_backends.cuda.libs.nvrtc.check_status()

NVRTCError: NVRTC_ERROR_COMPILATION (6)

During handling of the above exception, another exception occurred:

CompileException                          Traceback (most recent call last)
<ipython-input-1-859328117964> in <module>
      9 cuda_blurred = cupy_gaussian_filter(cuda_image, sigma=20)
     10 
---> 11 res = peak_local_max(cuda_blurred, min_distance=1)
     12 
     13 result1 = cupy.asnumpy(res)

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\skimage\_shared\utils.py in fixed_func(*args, **kwargs)
    118                 # warn that arg_name is deprecated
    119                 warnings.warn(warning_msg, FutureWarning, stacklevel=2)
--> 120             return func(*args, **kwargs)
    121 
    122         return fixed_func

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\skimage\feature\peak.py in peak_local_max(image, min_distance, threshold_abs, threshold_rel, exclude_border, indices, num_peaks, footprint, labels, num_peaks_per_label, p_norm)
    273     if labels is None:
    274         # Non maximum filter
--> 275         mask = _get_peak_mask(image, footprint, threshold)
    276 
    277         mask = _exclude_border(mask, border_width)

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\skimage\feature\peak.py in _get_peak_mask(image, footprint, threshold, mask)
     42         return image > threshold
     43 
---> 44     image_max = ndi.maximum_filter(image, footprint=footprint, mode="constant")
     45 
     46     out = image == image_max

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\scipy\ndimage\filters.py in maximum_filter(input, size, footprint, output, mode, cval, origin)
   1366     .. seealso:: :func:`scipy.ndimage.maximum_filter`
   1367     """
-> 1368     return _min_or_max_filter(
   1369         input, size, footprint, None, output, mode, cval, origin, "max"
   1370     )

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\scipy\ndimage\filters.py in _min_or_max_filter(input, size, ftprnt, structure, output, mode, cval, origin, func)
   1386         # Seperable filter, run as a series of 1D filters
   1387         fltr = minimum_filter1d if func == "min" else maximum_filter1d
-> 1388         return _filters_core._run_1d_filters(
   1389             [fltr if size > 1 else None for size in sizes],
   1390             input,

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\scipy\ndimage\_filters_core.py in _run_1d_filters(filters, input, args, output, mode, cval, origin, dtype_mode)
    101         if fltr is None:
    102             continue
--> 103         fltr(input, arg, axis, output, mode, cval, origin)
    104         input, output = output, temp if first else input
    105         first = False

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\scipy\ndimage\filters.py in maximum_filter1d(input, size, axis, output, mode, cval, origin)
   1473     .. seealso:: :func:`scipy.ndimage.maximum_filter1d`
   1474     """
-> 1475     return _min_or_max_1d(input, size, axis, output, mode, cval, origin, "max")
   1476 
   1477 

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\scipy\ndimage\filters.py in _min_or_max_1d(input, size, axis, output, mode, cval, origin, func)
   1503         has_weights=False,
   1504     )
-> 1505     return _filters_core._call_kernel(
   1506         kernel, input, None, output, weights_dtype=bool
   1507     )

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupyimg\scipy\ndimage\_filters_core.py in _call_kernel(kernel, input, weights, output, structure, weights_dtype, structure_dtype)
    150         output, temp = _util._get_output(output.dtype, input), output
    151     args.append(output)
--> 152     kernel(*args)
    153     if needs_temp:
    154         temp[...] = output[...]

cupy\core\_kernel.pyx in cupy.core._kernel.ElementwiseKernel.__call__()

cupy\core\_kernel.pyx in cupy.core._kernel.ElementwiseKernel._get_elementwise_kernel()

cupy\_util.pyx in cupy._util.memoize.decorator.ret()

cupy\core\_kernel.pyx in cupy.core._kernel._get_elementwise_kernel()

cupy\core\_kernel.pyx in cupy.core._kernel._get_simple_elementwise_kernel()

cupy\core\_kernel.pyx in cupy.core._kernel._get_simple_elementwise_kernel()

cupy\core\core.pyx in cupy.core.core.compile_with_cache()

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\cuda\compiler.py in compile_with_cache(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream)
    368             name_expressions, log_stream, cache_in_memory)
    369     else:
--> 370         return _compile_with_cache_cuda(
    371             source, options, arch, cache_dir, extra_source, backend,
    372             enable_cooperative_groups, name_expressions, log_stream,

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\cuda\compiler.py in _compile_with_cache_cuda(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, cache_in_memory)
    433     if backend == 'nvrtc':
    434         cu_name = '' if cache_in_memory else name + '.cu'
--> 435         ptx, mapping = compile_using_nvrtc(
    436             source, options, arch, cu_name, name_expressions,
    437             log_stream, cache_in_memory)

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\cuda\compiler.py in compile_using_nvrtc(source, options, arch, filename, name_expressions, log_stream, cache_in_memory)
    209                 cu_file.write(source)
    210 
--> 211             return _compile(source, options, cu_path,
    212                             name_expressions, log_stream)
    213     else:

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\cuda\compiler.py in _compile(source, options, cu_path, name_expressions, log_stream)
    193                              name_expressions=name_expressions)
    194         try:
--> 195             ptx, mapping = prog.compile(options, log_stream)
    196         except CompileException as e:
    197             dump = _get_bool_env_variable(

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\cuda\compiler.py in compile(self, options, log_stream)
    558         except nvrtc.NVRTCError:
    559             log = nvrtc.getProgramLog(self.ptr)
--> 560             raise CompileException(log, self.src, self.name, options,
    561                                    'nvrtc' if not runtime.is_hip else 'hiprtc')
    562 

CompileException: C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\core\include\cupy/complex/complex.h(94): warning: __host__ annotation is ignored on a function("complex") that is explicitly defaulted on its first declaration

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\core\include\cupy/complex/complex.h(94): warning: __device__ annotation is ignored on a function("complex") that is explicitly defaulted on its first declaration

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\core\include\cupy/complex/complex.h(101): warning: __host__ annotation is ignored on a function("complex") that is explicitly defaulted on its first declaration

C:\Programs\miniconda3\envs\beetlesafari\lib\site-packages\cupy\core\include\cupy/complex/complex.h(101): warning: __device__ annotation is ignored on a function("complex") that is explicitly defaulted on its first declaration

C:\Users\ROBERT~1\AppData\Local\Temp\tmpsxqh350g\7dc0f5bf0f642b710c0b8a179916eb98_2.cubin.cu(10): catastrophic error: cannot open source file "type_traits"

1 catastrophic error detected in the compilation of "C:\Users\ROBERT~1\AppData\Local\Temp\tmpsxqh350g\7dc0f5bf0f642b710c0b8a179916eb98_2.cubin.cu".
Compilation terminated.

Yes, getting a working version of cupyimg is still a bit tricky as it and all its dependencies are under constant developement.
I had very similar problems until a few days ago.
You may want to refer to this issue:

For me, some of the recent changes that Greg pushed together with cupy build 9.0.0b3 from conda-forge made it work (on Ubuntu/conda). He also announced that this will be released more publicly soon, probably also with easier installation.

1 Like

that definitely will no longer work (see issue above)

1 Like

Yeah, I will wait for that. The list of supported functions looks amazing and good things take a while. Thanks again for the support!

Yes, I am patiently waiting for the cupy/numpy compatible API to clEsperanto.

1 Like

I was considering to work on API-compatibility with scikit-image as soon as its version 1.0 is out.

1 Like