From 6e3fd316cda41d34cff7922993aa981a15435020 Mon Sep 17 00:00:00 2001 From: Gregory Lee Date: Wed, 25 Sep 2024 23:58:01 -0400 Subject: [PATCH] update vendored CUDA includes to match CuPy >= 13.3 (#781) CuPy 13.3 removed dependence on jitify (see: https://github.com/cupy/cupy/pull/8473) Make the same change in our vendored code (but only if CuPy >= 13.3). Authors: - Gregory Lee (https://github.com/grlee77) Approvers: - https://github.com/jakirkham URL: https://github.com/rapidsai/cucim/pull/781 --- .../_vendored/_ndimage_filters_core.py | 22 +++++++++++++++++-- .../skimage/filters/_separable_filtering.py | 6 ++++- 2 files changed, 25 insertions(+), 3 deletions(-) diff --git a/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters_core.py b/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters_core.py index 788d34497..be5348da4 100644 --- a/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters_core.py +++ b/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters_core.py @@ -3,12 +3,15 @@ import cupy import numpy +from packaging.version import parse from cucim.skimage._vendored import ( _internal as internal, _ndimage_util as _util, ) +CUPY_GTE_13_3_0 = parse(cupy.__version__) >= parse("13.3.0") + def _origins_to_offsets(origins, w_shape): return tuple(x // 2 + o for x, o in zip(w_shape, origins)) @@ -182,13 +185,24 @@ def _call_kernel( return output -_ndimage_includes = r""" +if CUPY_GTE_13_3_0: + _includes = r""" +#include // provide std:: coverage +""" +else: + _includes = r""" #include // let Jitify handle this +""" + +_ndimage_includes = ( + _includes + + r""" #include template<> struct std::is_floating_point : std::true_type {}; template<> struct std::is_signed : std::true_type {}; """ +) _ndimage_CAST_FUNCTION = """ @@ -352,7 +366,11 @@ def _generate_nd_kernel( if has_mask: name += "_with_mask" preamble = _ndimage_includes + _ndimage_CAST_FUNCTION + preamble - options += ("--std=c++11", "-DCUPY_USE_JITIFY") + + if CUPY_GTE_13_3_0: + options += ("--std=c++11",) + else: + options += ("--std=c++11", "-DCUPY_USE_JITIFY") return cupy.ElementwiseKernel( in_params, out_params, diff --git a/python/cucim/src/cucim/skimage/filters/_separable_filtering.py b/python/cucim/src/cucim/skimage/filters/_separable_filtering.py index 31effe313..7f2da7cf0 100644 --- a/python/cucim/src/cucim/skimage/filters/_separable_filtering.py +++ b/python/cucim/src/cucim/skimage/filters/_separable_filtering.py @@ -5,6 +5,7 @@ from cucim.skimage._vendored import _ndimage_util as util from cucim.skimage._vendored._internal import _normalize_axis_index from cucim.skimage._vendored._ndimage_filters_core import ( + CUPY_GTE_13_3_0, _ndimage_CAST_FUNCTION, _ndimage_includes, ) @@ -917,7 +918,10 @@ def _get_separable_conv_kernel( patch_per_block=patch_per_block, flip_kernel=flip_kernel, ) - options = ("--std=c++11", "-DCUPY_USE_JITIFY") + if CUPY_GTE_13_3_0: + options = ("--std=c++11",) + else: + options = ("--std=c++11", "-DCUPY_USE_JITIFY") m = cp.RawModule(code=code, options=options) return m.get_function(func_name), block, patch_per_block