From f1fa1e34e69ae73ae298548cff9e2a25affc8f6e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 10 Oct 2026 14:52:52 +0200 Subject: [PATCH 1/5] Sync third-party CuPy tests for kernel argument handling Add new tests for ufunc, reduction, elementwise kernel, copyto and concatenate argument handling. Tests relying on unsupported features (ufunc interface, CUDA array interface, ReductionKernel, ElementwiseKernel) are skipped; expected error messages are adapted to dpnp where only the wording differs. --- .../cupy/core_tests/test_ndarray_ufunc.py | 22 ++++++++++ .../cupy/core_tests/test_reduction.py | 42 +++++++++++++++++++ .../cupy/core_tests/test_userkernel.py | 40 ++++++++++++++++++ .../cupy/manipulation_tests/test_basic.py | 6 +++ .../cupy/manipulation_tests/test_join.py | 20 +++++++++ 5 files changed, 130 insertions(+) diff --git a/dpnp/tests/third_party/cupy/core_tests/test_ndarray_ufunc.py b/dpnp/tests/third_party/cupy/core_tests/test_ndarray_ufunc.py index b2edad0ddb2b..8c02d3d0df98 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_ndarray_ufunc.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_ndarray_ufunc.py @@ -116,6 +116,18 @@ def test_subclass_unary_op(self): outb = numpy.sin(b) testing.assert_allclose(outa, outb) + @pytest.mark.parametrize("rows", [0, 2]) + def test_subclass_broadcast_template(self, rows): + class Array(cupy.ndarray): + def __array_finalize__(self, parent): + self.parent = parent + + a = cupy.ones((1, 3)).view(Array) + result = cupy.add(a, cupy.ones((rows, 3))) + assert isinstance(result, Array) + assert result.parent is a + testing.assert_array_equal(result, numpy.full((rows, 3), 2.0)) + def test_subclass_binary_op(self): a0 = cupy.array([0, 1, 2]).view(C) a0.info = 1 @@ -185,6 +197,16 @@ def test_ufunc_reduceat(self, xp): class TestUfunc: + def test_broadcast_error_includes_output(self): + a = cupy.empty((0, 3)) + out = cupy.empty((2, 3)) + with pytest.raises(ValueError) as error: + cupy.add(a, 1, out=out) + assert str(error.value) == ( + "The shape of input and output arrays are inconsistent. " + "Expected output shape is (0, 3), got (2, 3)" + ) + @pytest.mark.parametrize( "ufunc", [ diff --git a/dpnp/tests/third_party/cupy/core_tests/test_reduction.py b/dpnp/tests/third_party/cupy/core_tests/test_reduction.py index 660471a3e189..084831f83800 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_reduction.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_reduction.py @@ -65,6 +65,27 @@ def get_sum_func(self): class TestSimpleReductionFunction( unittest.TestCase, SimpleReductionFunctionTestBase ): + def test_overlapping_output(self): + kernel = _core.create_reduction_func( + "overlapping_sum", ("q->q",), ("in0", "a + b", "out0 = a", None), 0 + ) + a = cupy.arange(64 * 64, dtype="int64").reshape(64, 64) + expected = a.get() + numpy.sum(expected, axis=1, out=expected[0]) + out = a[0] + result = kernel(a, axis=1, out=out) + assert result is out + testing.assert_array_equal(a, expected) + + def test_subclass_output(self): + class Array(cupy.ndarray): + pass + + a = cupy.arange(6, dtype="int8").view(Array) + result = self.get_sum_func()(a) + assert type(result) is cupy.ndarray + testing.assert_array_equal(result, 15) + def test_shape1(self): for i in range(1, 10): self.check_int8_sum((2**i,)) @@ -170,6 +191,27 @@ def get_sum_func(self): class TestReductionKernel(ReductionKernelTestBase, unittest.TestCase): + def test_explicit_none_output(self): + a = cupy.arange(6, dtype="float32").reshape(2, 3) + result = self.get_sum_func()(a, None, axis=1) + testing.assert_array_equal(result, a.get().sum(axis=1)) + + def test_supplied_output_shape(self): + kernel = self.get_sum_func() + for shape in ((2, 1), (2, 3), (0, 3)): + a = cupy.arange(numpy.prod(shape), dtype="float32").reshape(shape) + out = cupy.empty(shape[0], dtype="float32") + result = kernel(a, axis=1, out=out) + assert result is out + testing.assert_array_equal(result, a.get().sum(axis=1)) + + def test_scalar_input_without_shape(self): + kernel = cupy.ReductionKernel( + "float32 x", "float32 y", "x", "a + b", "y = a", "0", "scalar_sum" + ) + with pytest.raises(ValueError, match="Loop size is undecided"): + kernel(1) + def test_shape1(self): for i in range(1, 10): self.check_int8_sum((2**i,)) diff --git a/dpnp/tests/third_party/cupy/core_tests/test_userkernel.py b/dpnp/tests/third_party/cupy/core_tests/test_userkernel.py index aa839f61e9bc..b006fc6c671c 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_userkernel.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_userkernel.py @@ -18,6 +18,17 @@ class TestUserkernel(unittest.TestCase): + def test_explicit_none_output(self): + kernel = cupy.ElementwiseKernel( + "T x", "T y, T z", "y = x + 1; z = x + 2", "optional_outputs" + ) + a = cupy.arange(6, dtype="float32") + out = cupy.empty_like(a) + y, z = kernel(a, None, out) + assert z is out + testing.assert_array_equal(y, a.get() + 1) + testing.assert_array_equal(z, a.get() + 2) + def test_manual_indexing(self, n=100): in1 = cupy.random.uniform(-1, 1, n).astype(cupy.float32) in2 = cupy.random.uniform(-1, 1, n).astype(cupy.float32) @@ -43,6 +54,23 @@ def test_manual_indexing(self, n=100): testing.assert_array_equal(out1, out2) + def test_broadcast_error_ignores_raw_shape(self): + kernel = cupy.ElementwiseKernel( + "raw float32 raw_x, float32 x, float32 y", + "float32 z", + "z = x + y", + "raw_broadcast_error", + ) + raw = cupy.empty(17, dtype="float32") + x = cupy.empty(2, dtype="float32") + y = cupy.empty(3, dtype="float32") + with pytest.raises(ValueError) as error: + kernel(raw, x, y) + assert str(error.value) == ( + "operands could not be broadcast together with shapes " + "() (2,) (3,)" + ) + def test_python_scalar(self): for typ in (int, float, bool): dtype = numpy.dtype(typ).type @@ -143,6 +171,18 @@ def test_all_raws(self): with self.raises_size_required(): kernel2(self.arr1, self.arr2) + def test_all_scalars(self): + kernel = self.create_kernel((False,), (False,)) + with self.raises_size_required(): + kernel(1) + assert kernel(1, size=2).shape == (2,) + + def test_no_arguments(self): + kernel = self.create_kernel((), (False,)) + with self.raises_size_required(): + kernel() + assert kernel(size=2).shape == (2,) + def test_all_nonraws(self): # All arrays are not raw -> size not allowed kernel1 = self.create_kernel((False, False), (False,)) diff --git a/dpnp/tests/third_party/cupy/manipulation_tests/test_basic.py b/dpnp/tests/third_party/cupy/manipulation_tests/test_basic.py index 1d790eeb7597..14b4bef5a7c1 100644 --- a/dpnp/tests/third_party/cupy/manipulation_tests/test_basic.py +++ b/dpnp/tests/third_party/cupy/manipulation_tests/test_basic.py @@ -96,6 +96,12 @@ def test_copyto_where(self, xp, dtype): xp.copyto(a, b, where=c) return a + @testing.numpy_cupy_array_equal() + def test_copyto_overlapping_where(self, xp): + a = xp.ones((1024, 1024), dtype=bool) + xp.copyto(a, False, where=a[::-1, ::-1]) + return a + @testing.for_all_dtypes() @testing.numpy_cupy_array_equal() def test_copyto_where_squeeze_broadcast(self, xp, dtype): diff --git a/dpnp/tests/third_party/cupy/manipulation_tests/test_join.py b/dpnp/tests/third_party/cupy/manipulation_tests/test_join.py index 5020313febd8..c10d1bc18d96 100644 --- a/dpnp/tests/third_party/cupy/manipulation_tests/test_join.py +++ b/dpnp/tests/third_party/cupy/manipulation_tests/test_join.py @@ -15,10 +15,30 @@ from dpnp.tests.helper import has_support_aspect64 from dpnp.tests.third_party.cupy import testing from dpnp.tests.third_party.cupy.testing._helper import skip_if_after_baseline +from dpnp.tests.third_party.cupy.testing._protocol_helpers import ( + DummyObjectWithCudaArrayInterface, + DummyObjectWithCuPyGetNDArray, +) class TestJoin: + @pytest.mark.skip("CUDA array interface is not supported") + @pytest.mark.parametrize( + "wrapper", + [DummyObjectWithCuPyGetNDArray, DummyObjectWithCudaArrayInterface], + ) + def test_concatenate_array_protocols(self, wrapper): + a = cupy.arange(3) + result = cupy.concatenate((a, wrapper(a))) + testing.assert_array_equal(result, [0, 1, 2, 0, 1, 2]) + + def test_concatenate_invalid_input(self): + with pytest.raises( + TypeError, match="An array must be any of supported type" + ): + cupy.concatenate((cupy.ones(2), 1.0)) + @testing.for_all_dtypes(name="dtype1") @testing.for_all_dtypes(name="dtype2") @testing.numpy_cupy_array_equal(type_check=has_support_aspect64()) From 9669e3e082463991a9e83c3bed29d975fce48aa2 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 10 Oct 2026 15:09:43 +0200 Subject: [PATCH 2/5] Skip large-memory third-party CuPy tests when allocation fails Skip the argmax/argmin int32 overflow tests when the large array cannot be allocated, rework the huge FFT allocation test to check free memory upfront and clean up afterwards, and raise the timeout for the jitify tests. --- .../third_party/cupy/core_tests/test_raw.py | 1 + .../third_party/cupy/fft_tests/test_fft.py | 35 ++++++++++++------- .../cupy/sorting_tests/test_search.py | 11 ++++-- 3 files changed, 32 insertions(+), 15 deletions(-) diff --git a/dpnp/tests/third_party/cupy/core_tests/test_raw.py b/dpnp/tests/third_party/cupy/core_tests/test_raw.py index bb548ad810c6..259772a418cc 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_raw.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_raw.py @@ -410,6 +410,7 @@ def find_nvcc_ver(): ] jitify_markers = [ testing.slow(), + pytest.mark.timeout(1200), # Cold header initialization can take minutes. pytest.mark.filterwarnings( "ignore:jitify=True is deprecated:DeprecationWarning" ), diff --git a/dpnp/tests/third_party/cupy/fft_tests/test_fft.py b/dpnp/tests/third_party/cupy/fft_tests/test_fft.py index 3c0f9b3ad52a..acb4867595e1 100644 --- a/dpnp/tests/third_party/cupy/fft_tests/test_fft.py +++ b/dpnp/tests/third_party/cupy/fft_tests/test_fft.py @@ -358,20 +358,29 @@ def test_default_fft_func(self, enable_nd): class TestFftAllocate: # @pytest.mark.thread_unsafe(reason="does large allocations") def test_fft_allocate(self): - # Check CuFFTError is not raised when the GPU memory is enough. - # See https://github.com/cupy/cupy/issues/1063 - # TODO(mizuno): Simplify "a" after memory compaction is implemented. - a = [] - for i in range(10): - a.append(cupy.empty(100000000)) - del a - b = cupy.empty(100000007, dtype=cupy.float32) - cupy.fft.fft(b) - # Free huge memory for slow test - del b - cupy.get_default_memory_pool().free_all_blocks() - # Clean up FFT plan cache + # Check that a huge FFT succeeds when GPU memory is sufficient. + # TODO(mizuno): Revisit this workaround after memory compaction is + # implemented. + pool = cupy.get_default_memory_pool() + n = 100000007 cupy.fft.config.clear_plan_cache() + pool.free_all_blocks() + free_memory, _ = cupy.cuda.Device().mem_info + # Float32 input, complex64 input/output, and cuFFT's worst-case + # workspace of eight complex64 arrays need n*(4 + 2*8 + 8*8) bytes + # (~8.4 GB). Round up to 10 GiB for planning allocations and alignment. + # https://docs.nvidia.com/cuda/cufft/index.html#fourier-transform-setup + required_memory = 10 * 1024**3 + if free_memory < required_memory: + pytest.skip("Not enough GPU memory for cuFFT planning") + + b = cupy.empty(n, dtype=cupy.float32) + try: + cupy.fft.fft(b) + finally: + del b + cupy.fft.config.clear_plan_cache() + pool.free_all_blocks() @testing.with_requires("numpy>=2.0") diff --git a/dpnp/tests/third_party/cupy/sorting_tests/test_search.py b/dpnp/tests/third_party/cupy/sorting_tests/test_search.py index 11618225c163..07ef3a2bc896 100644 --- a/dpnp/tests/third_party/cupy/sorting_tests/test_search.py +++ b/dpnp/tests/third_party/cupy/sorting_tests/test_search.py @@ -7,6 +7,7 @@ # import cupy._core._accelerator as _acc # from cupy._core import _cub_reduction +from dpnp.exceptions import USMAllocationError from dpnp.tests.helper import has_support_aspect64 from dpnp.tests.third_party.cupy import testing @@ -90,7 +91,10 @@ def test_argmax_zero_size_axis1(self, xp, dtype): @testing.slow # @pytest.mark.thread_unsafe(reason="allocation too large.") def test_argmax_int32_overflow(self): - a = cupy.arange(2**32 + 1, dtype=cupy.default_float_type()) + try: + a = cupy.arange(2**32 + 1, dtype=cupy.default_float_type()) + except USMAllocationError: + pytest.skip("Not enough GPU memory for the large array") assert a.argmax().item() == 2**32 @testing.for_all_dtypes(no_complex=True) @@ -170,7 +174,10 @@ def test_argmin_zero_size_axis1(self, xp, dtype): @testing.slow # @pytest.mark.thread_unsafe(reason="allocation too large.") def test_argmin_int32_overflow(self): - a = cupy.arange(2**32 + 1, dtype=cupy.default_float_type()) + try: + a = cupy.arange(2**32 + 1, dtype=cupy.default_float_type()) + except USMAllocationError: + pytest.skip("Not enough GPU memory for the large array") cupy.negative(a, out=a) assert a.argmin().item() == 2**32 From 0907bdc82ede412465cfc1a478cc613ed2a7c946 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 10 Oct 2026 15:14:19 +0200 Subject: [PATCH 3/5] Add Windows thread-safety helper to third-party CuPy dlpack tests Add the thread_unsafe_on_windows helper and keep its use on test_conversion_device_to_cpu commented out, since the thread_unsafe marker is not registered in dpnp and the test is already skipped. --- dpnp/tests/third_party/cupy/core_tests/test_dlpack.py | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/dpnp/tests/third_party/cupy/core_tests/test_dlpack.py b/dpnp/tests/third_party/cupy/core_tests/test_dlpack.py index e8fac152f758..7c1fff434ed5 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_dlpack.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_dlpack.py @@ -1,6 +1,7 @@ from __future__ import annotations import ctypes +import sys import dpctl import numpy @@ -17,6 +18,15 @@ ] +def thread_unsafe_on_windows(func): + """Serialize CPU access to managed memory on Windows.""" + if sys.platform == "win32": + return pytest.mark.thread_unsafe( + reason="CPU access to managed memory overlaps GPU work" + )(func) + return func + + # TODO: to roll back the changes once the issue with CUDA support is resolved for random def _gen_array(dtype, alloc_q=None): if cupy.issubdtype(dtype, numpy.unsignedinteger): @@ -237,6 +247,7 @@ def test_conversion_bad_device(self): arr.__dlpack__(dl_device=(9, 0), max_version=(1, 0)) @pytest.mark.skip("numpy doesn't support kDLOneAPI device type") + # @thread_unsafe_on_windows def test_conversion_device_to_cpu(self): # NOTE: This defaults to the old unversioned, which is needed for # NumPy 1.x support. From 47a60cb56200c519ec5ca543229c9966aaedd640 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 12 Oct 2026 19:16:52 +0200 Subject: [PATCH 4/5] Comment out timeout marker in third-party CuPy raw tests pytest-timeout is not part of the dpnp test dependencies. --- dpnp/tests/third_party/cupy/core_tests/test_raw.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpnp/tests/third_party/cupy/core_tests/test_raw.py b/dpnp/tests/third_party/cupy/core_tests/test_raw.py index 259772a418cc..1561a534d9b9 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_raw.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_raw.py @@ -410,7 +410,7 @@ def find_nvcc_ver(): ] jitify_markers = [ testing.slow(), - pytest.mark.timeout(1200), # Cold header initialization can take minutes. + # pytest.mark.timeout(1200), # Cold header initialization can take minutes. pytest.mark.filterwarnings( "ignore:jitify=True is deprecated:DeprecationWarning" ), From 178d3bfdd959796384a52239bef3699f0d4493fc Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 12 Oct 2026 19:39:30 +0200 Subject: [PATCH 5/5] Sync third-party CuPy quantile tests for outer indexing Parametrize test_quantile_out over C and F output order and check the supplied out array, and add tests for singleton non-reduced dimensions and NaN in the next slice. TestQuantileMethods remains skipped since dpnp.quantile() is not implemented yet. --- .../cupy/statistics_tests/test_order.py | 36 ++++++++++++++++--- 1 file changed, 32 insertions(+), 4 deletions(-) diff --git a/dpnp/tests/third_party/cupy/statistics_tests/test_order.py b/dpnp/tests/third_party/cupy/statistics_tests/test_order.py index 97d464fa63f2..074bb3d2baca 100644 --- a/dpnp/tests/third_party/cupy/statistics_tests/test_order.py +++ b/dpnp/tests/third_party/cupy/statistics_tests/test_order.py @@ -242,13 +242,15 @@ def test_quantile_keepdims(self, xp, dtype, method): q = testing.shaped_random((5,), xp, scale=1) return xp.quantile(a, q, axis=None, keepdims=True, method=method) + @pytest.mark.parametrize("order", ["C", "F"]) @testing.for_float_dtypes(no_float16=True) # NumPy raises error on int8 - @testing.numpy_cupy_allclose(rtol=1e-6) - def test_quantile_out(self, xp, dtype, method): + @testing.numpy_cupy_allclose(rtol=1e-6, contiguous_check=False) + def test_quantile_out(self, xp, dtype, method, order): a = testing.shaped_random((10, 2, 3, 2), xp, dtype) q = testing.shaped_random((5,), xp, dtype=dtype, scale=1) - out = testing.shaped_random((5, 10, 2, 3), xp, dtype) - return xp.quantile(a, q, axis=-1, method=method, out=out) + out = xp.empty((5, 10, 2, 3), dtype=dtype, order=order) + result = xp.quantile(a, q, axis=-1, method=method, out=out) + return result, out @testing.for_float_dtypes(no_float16=True) @testing.numpy_cupy_allclose(rtol=1e-6) @@ -285,6 +287,32 @@ def test_quantile_axis_and_keepdims(self, xp, dtype, method): q = testing.shaped_random((5,), xp, scale=1) return xp.quantile(a, q, axis=0, keepdims=True, method=method) + @pytest.mark.parametrize( + "shape, axis", + [ + ((1, 10), 1), + ((10, 1), 0), + ((1, 1, 10), -1), + ((1, 10, 1), 1), + ], + ) + @testing.for_all_dtypes(no_float16=True, no_bool=True, no_complex=True) + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_quantile_singleton_dimensions( + self, xp, dtype, method, shape, axis + ): + # See gh-9508 + # Every dimension that is not reduced has length 1 + a = testing.shaped_random(shape, xp, dtype) + q = testing.shaped_random((5,), xp, scale=1) + return xp.quantile(a, q, axis=axis, method=method) + + @testing.for_float_dtypes(no_float16=True) + @testing.numpy_cupy_allclose() + def test_quantile_nan_in_next_slice(self, xp, dtype, method): + a = xp.array([[1.0, 2.0], [xp.nan, xp.nan]], dtype=dtype) + return xp.quantile(a, 1.0, axis=1, method=method) + class TestOrder: