antonwolfy · Oct 13, 2026 · Oct 10, 2026 · Oct 10, 2026 · Oct 10, 2026 · Oct 10, 2026
@@ -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.

@@ -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",
         [

@@ -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"
     ),

@@ -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,))

@@ -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,))

@@ -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")

@@ -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):

@@ -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())

@@ -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
 

@@ -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: