From 9034fca9a2e08aae1ee2546d925f49c4dc3566f8 Mon Sep 17 00:00:00 2001 From: gpinkert Date: Thu, 19 Feb 2026 03:44:48 +0000 Subject: [PATCH 1/6] fix(hip): use 64-bit mask for warp shuffle/vote intrinsics ROCm 7+ requires a 64-bit mask type for __shfl_*_sync and __any_sync. Passing a 32-bit literal (0xffffffff) causes a compile-time static assertion failure in HIPRTC. - Consolidate mask constants into cupy/_core/_kernel.pyx - Derive the mask value from the device warp size on HIP - Append ULL suffix to the C literal on HIP to satisfy the 64-bit type requirement; on older ROCm the mask is stripped by macros so 64-bit is harmless - Guard the __shfl_*_sync compatibility macros behind HIP_VERSION < 70000000 so they do not conflict with native declarations - Cast the user-supplied mask to unsigned long long in the JIT rawkernel path when targeting ROCm 7+ - Query actual device warp size via _get_warpsize() for the JIT default shuffle width instead of assuming 64 Closes #9742 --- cupy/_core/__init__.py | 3 +++ cupy/_core/_kernel.pyx | 20 ++++++++++++++++++++ cupy/_core/_routines_indexing.pyx | 11 ++++++----- cupy/_core/_routines_math.pyx | 9 +++++---- cupy/_core/_routines_sorting.pyx | 11 +++++++---- cupyx/jit/_builtin_funcs.py | 12 ++++++++---- tests/cupyx_tests/jit_tests/test_raw.py | 3 ++- 7 files changed, 51 insertions(+), 18 deletions(-) diff --git a/cupy/_core/__init__.py b/cupy/_core/__init__.py index e4ce673c78d..68edc48052e 100644 --- a/cupy/_core/__init__.py +++ b/cupy/_core/__init__.py @@ -21,6 +21,9 @@ from cupy._core._kernel import ElementwiseKernel # NOQA from cupy._core._kernel import ufunc # NOQA from cupy._core._kernel import _get_warpsize # NOQA +from cupy._core._kernel import _full_mask # NOQA +from cupy._core._kernel import _full_mask_hex # NOQA +from cupy._core._kernel import _is_hip_7_plus # NOQA from cupy._core._reduction import create_reduction_func # NOQA from cupy._core._reduction import ReductionKernel # NOQA from cupy._core._routines_binary import bitwise_and # NOQA diff --git a/cupy/_core/_kernel.pyx b/cupy/_core/_kernel.pyx index 058b1be0669..ee6059396da 100644 --- a/cupy/_core/_kernel.pyx +++ b/cupy/_core/_kernel.pyx @@ -28,6 +28,7 @@ from cupy._core.core cimport compile_with_cache from cupy._core.core cimport _ndarray_base from cupy._core cimport internal from cupy_backends.cuda.api cimport runtime +from cupy_backends.cuda.api import driver as _driver try: from cupy_backends.cuda.libs import cutensor as cuda_cutensor @@ -50,6 +51,25 @@ def _get_warpsize(): return runtime.getDeviceProperties(device_id)['warpSize'] +# ROCm 7+ requires a 64-bit mask type for __shfl_*_sync / __any_sync. +_is_hip_7_plus = (runtime._is_hip_environment + and _driver.get_build_version() >= 7_00_00000) + +# Full-lane mask: value is warp-size-dependent on HIP, always 32-bit on CUDA. +# On HIP the mask *type* must be 64-bit (ROCm 7+ enforces via static_assert), +# so _full_mask_hex() appends 'ULL' to the C literal. +if runtime._is_hip_environment: + _full_mask = (1 << _get_warpsize()) - 1 +else: + _full_mask = 0xffffffff + + +def _full_mask_hex(): + if runtime._is_hip_environment: + return hex(_full_mask) + 'ULL' + return hex(_full_mask) + + cdef str _get_simple_elementwise_kernel_code( tuple params_, tuple arginfos, str operation, str name, _TypeMap type_map, str preamble, str loop_prep='', str after_loop=''): diff --git a/cupy/_core/_routines_indexing.pyx b/cupy/_core/_routines_indexing.pyx index 921a603fa5e..f9b9320e6fe 100644 --- a/cupy/_core/_routines_indexing.pyx +++ b/cupy/_core/_routines_indexing.pyx @@ -6,7 +6,7 @@ import numpy import cupy import cupy._core.core as core from cupy.exceptions import AxisError -from cupy._core._kernel import ElementwiseKernel, _get_warpsize +from cupy._core._kernel import ElementwiseKernel, _get_warpsize, _full_mask_hex from cupy._core._ufuncs import elementwise_copy from libcpp cimport vector @@ -490,7 +490,7 @@ def _nonzero_kernel_incomplete_scan(block_size, warp_size=32): S x = 0; if (i < a.size()) x = a[i]; for (int j = 1; j < ${warp_size}; j *= 2) { - S tmp = __shfl_up_sync(0xffffffff, x, j, ${warp_size}); + S tmp = __shfl_up_sync(${full_mask}, x, j, ${warp_size}); if (lane_id - j >= 0) x += tmp; } if (lane_id == ${warp_size} - 1) smem[warp_id] = x; @@ -499,7 +499,7 @@ def _nonzero_kernel_incomplete_scan(block_size, warp_size=32): S y = 0; if (lane_id < n_warp) y = smem[lane_id]; for (int j = 1; j < n_warp; j *= 2) { - S tmp = __shfl_up_sync(0xffffffff, y, j, ${warp_size}); + S tmp = __shfl_up_sync(${full_mask}, y, j, ${warp_size}); if (lane_id - j >= 0) y += tmp; } int block_id = i / ${block_size}; @@ -510,7 +510,7 @@ def _nonzero_kernel_incomplete_scan(block_size, warp_size=32): } __syncthreads(); x += smem[warp_id]; - S x0 = __shfl_up_sync(0xffffffff, x, 1, ${warp_size}); + S x0 = __shfl_up_sync(${full_mask}, x, 1, ${warp_size}); if (lane_id == 0) { x0 = smem[warp_id]; } @@ -523,7 +523,8 @@ def _nonzero_kernel_incomplete_scan(block_size, warp_size=32): j = j_next; } } - """).substitute(block_size=block_size, warp_size=warp_size) + """).substitute(block_size=block_size, warp_size=warp_size, + full_mask=_full_mask_hex()) return cupy.ElementwiseKernel(in_params, out_params, loop_body, 'cupy_nonzero_kernel_incomplete_scan', loop_prep=loop_prep) diff --git a/cupy/_core/_routines_math.pyx b/cupy/_core/_routines_math.pyx index b8b4e9f3ffc..05437666fdb 100644 --- a/cupy/_core/_routines_math.pyx +++ b/cupy/_core/_routines_math.pyx @@ -5,7 +5,7 @@ import numpy import cupy from cupy._core._reduction import create_reduction_func -from cupy._core._kernel import create_ufunc, _get_warpsize +from cupy._core._kernel import create_ufunc, _get_warpsize, _full_mask_hex from cupy._core._scalar import get_typename from cupy._core._ufuncs import elementwise_copy import cupy._core.core as core @@ -193,7 +193,7 @@ def _cupy_bsum_shfl(op, chunk_size, warp_size=32): if (2*i < a.size()) x = a[2*i]; if (2*i + 1 < a.size()) x ${op}= a[2*i + 1]; for (int j = 1; j < ${warp_size}; j *= 2) { - x ${op}= __shfl_xor_sync(0xffffffff, x, j, ${warp_size}); + x ${op}= __shfl_xor_sync(${full_mask}, x, j, ${warp_size}); } if (lane_id == 0) smem[warp_id] = x; __syncthreads(); @@ -201,13 +201,14 @@ def _cupy_bsum_shfl(op, chunk_size, warp_size=32): x = ${identity}; if (lane_id < n_warp) x = smem[lane_id]; for (int j = 1; j < n_warp; j *= 2) { - x ${op}= __shfl_xor_sync(0xffffffff, x, j, ${warp_size}); + x ${op}= __shfl_xor_sync(${full_mask}, x, j, ${warp_size}); } int block_id = i / ${block_size}; if (lane_id == 0) b[block_id] = x; } """).substitute(block_size=block_size, warp_size=warp_size, - op=_op_char[op], identity=_identity[op]) + op=_op_char[op], identity=_identity[op], + full_mask=_full_mask_hex()) return cupy.ElementwiseKernel(in_params, out_params, loop_body, 'cupy_bsum_shfl', loop_prep=loop_prep) diff --git a/cupy/_core/_routines_sorting.pyx b/cupy/_core/_routines_sorting.pyx index 0918d22466f..2e2351cf367 100644 --- a/cupy/_core/_routines_sorting.pyx +++ b/cupy/_core/_routines_sorting.pyx @@ -6,6 +6,7 @@ import cupy from cupy.exceptions import AxisError from cupy._core._scalar import get_typename as _get_typename from cupy._core._ufuncs import elementwise_copy +from cupy._core._kernel import _full_mask_hex import cupy._core.core as core from cupy import _util from cupy.cuda import thrust @@ -387,7 +388,7 @@ def _partition_kernel(dtype): // If at least one thread in the warp has found t values that // can be selected, we update the first k elements. #if __CUDACC_VER_MAJOR__ >= 9 - if (__any_sync(0xffffffff, x >= t)) { + if (__any_sync(${full_mask}, x >= t)) { #else if (__any(x >= t)) { #endif @@ -419,7 +420,8 @@ def _partition_kernel(dtype): } } ''').substitute(name=name, merge_kernel=merge_kernel, dtype=dtype, - type_headers=type_headers) + type_headers=type_headers, + full_mask=_full_mask_hex()) module = compile_with_cache(source) return module.get_function(name), module.get_function(merge_kernel) @@ -513,7 +515,7 @@ def _argpartition_kernel(dtype): // If at least one thread in the warp has found t values that // can be selected, we update the first k elements. #if __CUDACC_VER_MAJOR__ >= 9 - if (__any_sync(0xffffffff, x >= t)) { + if (__any_sync(${full_mask}, x >= t)) { #else if (__any(x >= t)) { #endif @@ -546,6 +548,7 @@ def _argpartition_kernel(dtype): } } ''').substitute(name=name, merge_kernel=merge_kernel, dtype=dtype, - type_headers=type_headers) + type_headers=type_headers, + full_mask=_full_mask_hex()) module = compile_with_cache(source) return module.get_function(name), module.get_function(merge_kernel) diff --git a/cupyx/jit/_builtin_funcs.py b/cupyx/jit/_builtin_funcs.py index 83f8d52656d..f59a35b47a8 100644 --- a/cupyx/jit/_builtin_funcs.py +++ b/cupyx/jit/_builtin_funcs.py @@ -379,9 +379,11 @@ def call(self, env, mask, var, val_id, *, width=None): mask = mask.obj except Exception: raise TypeError('mask must be an integer') - if runtime.is_hip: + + _hip_7_plus = cupy._core._is_hip_7_plus + if runtime.is_hip and not _hip_7_plus: warnings.warn(f'mask {mask} is ignored on HIP', RuntimeWarning) - elif not (0x0 <= mask <= 0xffffffff): + if not (0x0 <= mask <= cupy._core._full_mask): raise ValueError('mask is out of range') # val_id refers to "delta" for shfl_{up, down}, "srcLane" for shfl, and @@ -398,12 +400,14 @@ def call(self, env, mask, var, val_id, *, width=None): if width.obj not in (2, 4, 8, 16, 32): raise ValueError('width needs to be power of 2') else: - width = Constant(64) if runtime.is_hip else Constant(32) + width = Constant(cupy._core._get_warpsize()) width = _compile._astype_scalar( width, _cuda_types.int32, 'same_kind', env) width = Data.init(width, env) - code = f'{name}({hex(mask)}, {var.code}, {val_id.code}' + mask_str = (f'(unsigned long long){hex(mask)}' + if _hip_7_plus else hex(mask)) + code = f'{name}({mask_str}, {var.code}, {val_id.code}' code += f', {width.code})' return Data(code, ctype) diff --git a/tests/cupyx_tests/jit_tests/test_raw.py b/tests/cupyx_tests/jit_tests/test_raw.py index e438e345db0..6ec8aab323e 100644 --- a/tests/cupyx_tests/jit_tests/test_raw.py +++ b/tests/cupyx_tests/jit_tests/test_raw.py @@ -654,10 +654,11 @@ def test_shfl_down(self, dtype): N = 5 # __shfl_down() on HIP does not seem to have the same behavior... block = cupy._core._get_warpsize() + full_mask = (1 << block) - 1 @jit.rawkernel() def f(a): - value = jit.shfl_down_sync(0xffffffff, a[jit.threadIdx.x], N) + value = jit.shfl_down_sync(full_mask, a[jit.threadIdx.x], N) a[jit.threadIdx.x] = value a = cupy.arange(block, dtype=dtype) From 2bc717582ee4500108ee86c7d8cfb305a97d01c6 Mon Sep 17 00:00:00 2001 From: gpinkert Date: Thu, 19 Feb 2026 05:39:22 +0000 Subject: [PATCH 3/6] fix(hip): use correct enum types in TextureDesc Cython struct The Cython declaration of cudaTextureDesc used plain 'int' for the addressMode, filterMode, and readMode fields. On HIP these are distinct enum types (TextureAddressMode, TextureFilterMode, TextureReadMode), causing type mismatch warnings and potential undefined behavior. Use the proper enum types to match the C header. --- cupy_backends/cuda/api/_runtime_typedef.pxi | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cupy_backends/cuda/api/_runtime_typedef.pxi b/cupy_backends/cuda/api/_runtime_typedef.pxi index 257cd6c1c67..6337afa3f8b 100644 --- a/cupy_backends/cuda/api/_runtime_typedef.pxi +++ b/cupy_backends/cuda/api/_runtime_typedef.pxi @@ -94,8 +94,8 @@ cdef extern from *: ctypedef struct TextureDesc 'cudaTextureDesc': TextureAddressMode addressMode[3] - int filterMode - int readMode + TextureFilterMode filterMode + TextureReadMode readMode int sRGB float borderColor[4] int normalizedCoords From 26bfcf4d97e4a49ff4dc32f2278eae490e8d6717 Mon Sep 17 00:00:00 2001 From: gpinkert Date: Wed, 1 Apr 2026 18:11:18 +0000 Subject: [PATCH 4/6] =?UTF-8?q?fix(hip):=20extend=20shfl=5Fsync=20workarou?= =?UTF-8?q?nd=20macros=20to=20cover=20ROCm=206.2=E2=80=936.4?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The previous guard (HIP_VERSION < 60200000) disabled CuPy's fallback macros starting at ROCm 6.2, assuming native __shfl_*_sync functions were available. However, ROCm 6.2–6.4 only provide them behind an opt-in flag (HIP_ENABLE_WARP_SYNC_BUILTINS) that is not set by default, leaving __shfl_*_sync undeclared. Replace the single version check with per-era logic: - ROCm < 6.2: always use workaround macros (no native builtins) - ROCm 6.2–6.4: use macros unless HIP_ENABLE_WARP_SYNC_BUILTINS - ROCm 7.0+: skip macros unless HIP_DISABLE_WARP_SYNC_BUILTINS --- cupy/_core/include/cupy/hip_workaround.cuh | 44 +++++++++++++++++----- 1 file changed, 34 insertions(+), 10 deletions(-) diff --git a/cupy/_core/include/cupy/hip_workaround.cuh b/cupy/_core/include/cupy/hip_workaround.cuh index c99c10ffba9..dff769104ca 100644 --- a/cupy/_core/include/cupy/hip_workaround.cuh +++ b/cupy/_core/include/cupy/hip_workaround.cuh @@ -3,17 +3,41 @@ #ifdef __HIP_DEVICE_COMPILE__ -// As per the comment below, use workaround conditionally: -// https://github.com/ROCm/clr/blob/68147fe9b20a72aa43e7898bdd9ba39bca4afd14/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h#L25 -#if (HIP_VERSION < 60200000) || defined(HIP_DISABLE_WARP_SYNC_BUILTINS) +// Determine whether native __shfl_*_sync functions are available. +// (defined in /opt/rocm/include/hip/amd_detail/amd_warp_sync_functions.h) +// +// The native builtins require a 64-bit mask and were introduced in ROCm 6.2 +// behind an opt-in macro: +// ROCm < 6.2 — not available at all +// ROCm 6.2–6.4 — available only if HIP_ENABLE_WARP_SYNC_BUILTINS is defined +// ROCm 7.0+ — available by default, disabled if HIP_DISABLE_WARP_SYNC_BUILTINS is defined +// +// When the native builtins are not available, we define compatibility macros +// that rewrite __shfl_*_sync(mask, ...) to __shfl_*(...), stripping the mask. +// This is safe because HIP wavefronts execute in lock-step. +#if !defined(HIP_VERSION) || HIP_VERSION < 60200000 + // ROCm < 6.2: no native builtins + #define CUPY_HIP_SHFL_WORKAROUND +#elif HIP_VERSION < 70000000 + // ROCm 6.2–6.4: native builtins only if user opted in + #if !defined(HIP_ENABLE_WARP_SYNC_BUILTINS) + #define CUPY_HIP_SHFL_WORKAROUND + #endif +#else + // ROCm 7.0+: native builtins unless user opted out + #if defined(HIP_DISABLE_WARP_SYNC_BUILTINS) + #define CUPY_HIP_SHFL_WORKAROUND + #endif +#endif -// ignore mask -#define __shfl_sync(mask, ...) __shfl(__VA_ARGS__) -#define __shfl_up_sync(mask, ...) __shfl_up(__VA_ARGS__) -#define __shfl_down_sync(mask, ...) __shfl_down(__VA_ARGS__) -#define __shfl_xor_sync(mask, ...) __shfl_xor(__VA_ARGS__) - -#endif // (HIP_VERSION < 60200000) || defined(HIP_DISABLE_WARP_SYNC_BUILTINS) +#ifdef CUPY_HIP_SHFL_WORKAROUND + // ignore mask + #define __shfl_sync(mask, ...) __shfl(__VA_ARGS__) + #define __shfl_up_sync(mask, ...) __shfl_up(__VA_ARGS__) + #define __shfl_down_sync(mask, ...) __shfl_down(__VA_ARGS__) + #define __shfl_xor_sync(mask, ...) __shfl_xor(__VA_ARGS__) + #undef CUPY_HIP_SHFL_WORKAROUND +#endif // In ROCm, threads in a warp march in lock-step, so we don't need to // synchronize the threads. But it doesn't guarantee the memory order, From 32bddf8a57686cf9d78fe0cd9cb077df2390c4e8 Mon Sep 17 00:00:00 2001 From: gpinkert Date: Sat, 18 Apr 2026 00:08:45 +0000 Subject: [PATCH 5/6] fixup! fix(hip): use 64-bit mask for warp shuffle/vote intrinsics --- cupy/_core/_kernel.pxd | 2 ++ cupy/_core/_kernel.pyx | 2 +- cupy/_core/_routines_indexing.pyx | 3 ++- cupy/_core/_routines_math.pyx | 3 ++- cupy/_core/_routines_sorting.pyx | 2 +- 5 files changed, 8 insertions(+), 4 deletions(-) diff --git a/cupy/_core/_kernel.pxd b/cupy/_core/_kernel.pxd index a5d445d5e38..2f4548a05e6 100644 --- a/cupy/_core/_kernel.pxd +++ b/cupy/_core/_kernel.pxd @@ -150,6 +150,8 @@ cpdef create_ufunc(name, ops, routine=*, preamble=*, doc=*, default_casting=*, loop_prep=*, out_ops=*, cutensor_op=*, scatter_op=*) +cpdef str _full_mask_hex() + cdef tuple _get_arginfos(list args) cdef str _get_kernel_params(tuple params, tuple arginfos, type_headers=*) diff --git a/cupy/_core/_kernel.pyx b/cupy/_core/_kernel.pyx index ee6059396da..e13b1dd90a4 100644 --- a/cupy/_core/_kernel.pyx +++ b/cupy/_core/_kernel.pyx @@ -64,7 +64,7 @@ else: _full_mask = 0xffffffff -def _full_mask_hex(): +cpdef str _full_mask_hex(): if runtime._is_hip_environment: return hex(_full_mask) + 'ULL' return hex(_full_mask) diff --git a/cupy/_core/_routines_indexing.pyx b/cupy/_core/_routines_indexing.pyx index f9b9320e6fe..30272324928 100644 --- a/cupy/_core/_routines_indexing.pyx +++ b/cupy/_core/_routines_indexing.pyx @@ -6,7 +6,8 @@ import numpy import cupy import cupy._core.core as core from cupy.exceptions import AxisError -from cupy._core._kernel import ElementwiseKernel, _get_warpsize, _full_mask_hex +from cupy._core._kernel import ElementwiseKernel, _get_warpsize +from cupy._core._kernel cimport _full_mask_hex from cupy._core._ufuncs import elementwise_copy from libcpp cimport vector diff --git a/cupy/_core/_routines_math.pyx b/cupy/_core/_routines_math.pyx index 05437666fdb..79df4c2daae 100644 --- a/cupy/_core/_routines_math.pyx +++ b/cupy/_core/_routines_math.pyx @@ -5,7 +5,8 @@ import numpy import cupy from cupy._core._reduction import create_reduction_func -from cupy._core._kernel import create_ufunc, _get_warpsize, _full_mask_hex +from cupy._core._kernel import create_ufunc, _get_warpsize +from cupy._core._kernel cimport _full_mask_hex from cupy._core._scalar import get_typename from cupy._core._ufuncs import elementwise_copy import cupy._core.core as core diff --git a/cupy/_core/_routines_sorting.pyx b/cupy/_core/_routines_sorting.pyx index 2e2351cf367..e92f1afcd7d 100644 --- a/cupy/_core/_routines_sorting.pyx +++ b/cupy/_core/_routines_sorting.pyx @@ -6,7 +6,7 @@ import cupy from cupy.exceptions import AxisError from cupy._core._scalar import get_typename as _get_typename from cupy._core._ufuncs import elementwise_copy -from cupy._core._kernel import _full_mask_hex +from cupy._core._kernel cimport _full_mask_hex import cupy._core.core as core from cupy import _util from cupy.cuda import thrust From e783ef49359e1180b9f49d465085f7ac633992af Mon Sep 17 00:00:00 2001 From: gpinkert Date: Sat, 18 Apr 2026 00:13:32 +0000 Subject: [PATCH 6/6] fixup! fix(hip): use 64-bit mask for warp shuffle/vote intrinsics --- cupy/_core/_kernel.pyx | 23 ++++++++++++++--------- cupyx/jit/_builtin_funcs.py | 2 +- tests/cupyx_tests/jit_tests/test_raw.py | 2 +- 3 files changed, 16 insertions(+), 11 deletions(-) diff --git a/cupy/_core/_kernel.pyx b/cupy/_core/_kernel.pyx index e13b1dd90a4..f05a6d1ecf1 100644 --- a/cupy/_core/_kernel.pyx +++ b/cupy/_core/_kernel.pyx @@ -55,19 +55,24 @@ def _get_warpsize(): _is_hip_7_plus = (runtime._is_hip_environment and _driver.get_build_version() >= 7_00_00000) -# Full-lane mask: value is warp-size-dependent on HIP, always 32-bit on CUDA. -# On HIP the mask *type* must be 64-bit (ROCm 7+ enforces via static_assert), -# so _full_mask_hex() appends 'ULL' to the C literal. -if runtime._is_hip_environment: - _full_mask = (1 << _get_warpsize()) - 1 -else: - _full_mask = 0xffffffff + +# Full-lane mask: warp-size-dependent on HIP, always 32-bit on CUDA. +# Computed lazily (per device) so that importing CuPy does not require a +# functional ROCm/CUDA runtime. +@_util.memoize(for_each_device=True) +def _full_mask(): + if runtime._is_hip_environment: + return (1 << _get_warpsize()) - 1 + return 0xffffffff +# On HIP the mask *type* must be 64-bit (ROCm 7+ enforces via static_assert), +# so the C literal is suffixed with 'ULL'. cpdef str _full_mask_hex(): + cdef str s = hex(_full_mask()) if runtime._is_hip_environment: - return hex(_full_mask) + 'ULL' - return hex(_full_mask) + return s + 'ULL' + return s cdef str _get_simple_elementwise_kernel_code( diff --git a/cupyx/jit/_builtin_funcs.py b/cupyx/jit/_builtin_funcs.py index f59a35b47a8..aa60e7341b1 100644 --- a/cupyx/jit/_builtin_funcs.py +++ b/cupyx/jit/_builtin_funcs.py @@ -383,7 +383,7 @@ def call(self, env, mask, var, val_id, *, width=None): _hip_7_plus = cupy._core._is_hip_7_plus if runtime.is_hip and not _hip_7_plus: warnings.warn(f'mask {mask} is ignored on HIP', RuntimeWarning) - if not (0x0 <= mask <= cupy._core._full_mask): + if not (0x0 <= mask <= cupy._core._full_mask()): raise ValueError('mask is out of range') # val_id refers to "delta" for shfl_{up, down}, "srcLane" for shfl, and diff --git a/tests/cupyx_tests/jit_tests/test_raw.py b/tests/cupyx_tests/jit_tests/test_raw.py index 6ec8aab323e..6e30b5c6dc4 100644 --- a/tests/cupyx_tests/jit_tests/test_raw.py +++ b/tests/cupyx_tests/jit_tests/test_raw.py @@ -654,7 +654,7 @@ def test_shfl_down(self, dtype): N = 5 # __shfl_down() on HIP does not seem to have the same behavior... block = cupy._core._get_warpsize() - full_mask = (1 << block) - 1 + full_mask = cupy._core._full_mask() @jit.rawkernel() def f(a):