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. 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_raw.py b/dpnp/tests/third_party/cupy/core_tests/test_raw.py index bb548ad810c6..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,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/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/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/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()) 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 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: