diff --git a/dpnp/tests/third_party/cupy/core_tests/test_cub_reduction.py b/dpnp/tests/third_party/cupy/core_tests/test_cub_reduction.py index 0bbc1296a3f..be68997b264 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_cub_reduction.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_cub_reduction.py @@ -52,8 +52,22 @@ def _test_can_use(self, i_shape, o_shape, r_axis, o_axis, order, expected): assert result is expected -@pytest.mark.parametrize("shape", [(2,), (2, 3), (2, 3, 4), (2, 3, 4, 5)]) -@pytest.mark.parametrize("order", ["C", "F"]) +_MIN_SIZE = cupy._core._cub_reduction._CUB_REDUCE_SIZE_THRESHOLD + + +@pytest.mark.parametrize( + "shape", + [ + (_MIN_SIZE,), + (_MIN_SIZE, _MIN_SIZE + 1), + (_MIN_SIZE, 3, _MIN_SIZE + 1), + (_MIN_SIZE, 3, 4, _MIN_SIZE + 1), + ], +) +@pytest.mark.parametrize( + "order", + ["C", "F"], +) class TestSimpleCubReductionKernelContiguity(CubReductionTestBase): @testing.for_contiguous_axes() @@ -139,15 +153,15 @@ def test_can_use_cub_oversize_input4(self): b = cupy.empty((), dtype=cupy.int8) assert self.can_use([a], [b], (1,), (0,)) is None - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp - # @pytest.mark.thread_unsafe( - # reason="AssertFunctionIsCalled and accelerate mutation.") + @pytest.mark.thread_unsafe( + reason="AssertFunctionIsCalled and accelerate mutation." + ) def test_can_use_accelerator_set_unset(self): # ensure we use CUB block reduction and not CUB device reduction old_routine_accelerators = _accelerator.get_routine_accelerators() _accelerator.set_routine_accelerators([]) - a = cupy.random.random((10, 10)) + a = cupy.random.random((10, _cub_reduction._CUB_REDUCE_SIZE_THRESHOLD)) # this is the only function we can mock; the rest is cdef'd func_name = "".join( ( diff --git a/dpnp/tests/third_party/cupy/core_tests/test_ndarray.py b/dpnp/tests/third_party/cupy/core_tests/test_ndarray.py index e17d09e2edd..26cd46221cd 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_ndarray.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_ndarray.py @@ -132,6 +132,55 @@ class UserNdarray(cupy.ndarray): b.custom_attr = 100 +@testing.parameterize( + *testing.product( + { + "np_order": ["C", "F"], + "np_pinned": [False], # no pinned memory + "pinned_alloc_fails": [False], # no pinned memory + "cp_setup": [ + ("C", False, (48, 16, 4)), + ("C", True, (16, 4)), + ("F", False, (4, 8, 24)), + ("F", True, (4, 8)), + ], + } + ) +) +class TestAsarray(unittest.TestCase): + def test_asarray(self): + cp_order, view, strides = self.cp_setup + shape = (2, 3, 4) + if self.np_pinned: + count = numpy.prod(shape) + dtype = numpy.float32() + pinned_ptr = cupy.cuda.alloc_pinned_memory(count * dtype.itemsize) + a_cpu = numpy.frombuffer( + pinned_ptr, dtype=dtype, count=count + ).reshape(shape) + else: + a_cpu = numpy.ndarray( + shape, dtype=numpy.float32, order=self.np_order + ) + a_cpu[...] = numpy.arange(a_cpu.size).reshape(a_cpu.shape) + if view: + a_cpu = a_cpu[:, 1, :] + try: + if self.pinned_alloc_fails: + cupy.cuda.set_pinned_memory_allocator(lambda _: None) + a = cupy.asarray(a_cpu, order=cp_order) + finally: + # None means "no pool", not "the default pool" + # cupy.cuda.set_pinned_memory_allocator( + # cupy.get_default_pinned_memory_pool().malloc + # ) + pass + assert a.flags.c_contiguous == (cp_order == "C") + assert a.flags.f_contiguous == (cp_order == "F") + assert a.strides == strides + testing.assert_array_equal(a_cpu, a) + + @testing.parameterize( *testing.product( { @@ -278,14 +327,16 @@ def test_copy_multi_device_with_stream(self): ) +@pytest.mark.filterwarnings( + # Shape setting is deprecated starting NumPy 2.5 + "ignore::DeprecationWarning" +) class TestNdarrayShape(unittest.TestCase): - @testing.with_requires("numpy>=2.5") @testing.numpy_cupy_array_equal() def test_shape_set(self, xp): arr = xp.ndarray((2, 3)) - with testing.assert_warns(DeprecationWarning): - arr.shape = (3, 2) + arr.shape = (3, 2) return xp.array(arr.shape) @pytest.mark.skip( @@ -298,15 +349,12 @@ def test_shape_set_infer(self, xp): arr.shape = (3, -1) return xp.array(arr.shape) - @testing.with_requires("numpy>=2.5") @testing.numpy_cupy_array_equal() def test_shape_set_int(self, xp): arr = xp.ndarray((2, 3)) - with testing.assert_warns(DeprecationWarning): - arr.shape = 6 + arr.shape = 6 return xp.array(arr.shape) - @pytest.mark.filterwarnings("ignore::DeprecationWarning") def test_shape_need_copy(self): # from cupy/cupy#5470 for xp in (numpy, cupy): @@ -563,6 +611,22 @@ def test_shape_mismatch(self): wrap_take(a, i, out=o) +@testing.parameterize( + {"shape": (3, 4, 5), "indices": (2, 3), "out_shape": (2, 3)}, + {"shape": (), "indices": (), "out_shape": ()}, +) +@pytest.mark.skip("no exception since NumPy 2.5") +class TestNdarrayTakeErrorTypeMismatch(unittest.TestCase): + + def test_output_type_mismatch(self): + for xp in (numpy, cupy): + a = testing.shaped_arange(self.shape, xp, numpy.int32) + i = testing.shaped_arange(self.indices, xp, numpy.int32) % 3 + o = testing.shaped_arange(self.out_shape, xp, numpy.float32) + with pytest.raises(TypeError): + wrap_take(a, i, out=o) + + @testing.parameterize( {"shape": (0,), "indices": (0,), "axis": None}, {"shape": (0,), "indices": (0, 1), "axis": None}, diff --git a/dpnp/tests/third_party/cupy/core_tests/test_ndarray_reduction.py b/dpnp/tests/third_party/cupy/core_tests/test_ndarray_reduction.py index b774cdff439..bf5cb8de56d 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_ndarray_reduction.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_ndarray_reduction.py @@ -6,6 +6,9 @@ import pytest import dpnp as cupy + +# import cupy._core._accelerator as _acc +# from cupy._core import _cub_reduction from dpnp.tests.third_party.cupy import testing @@ -34,11 +37,10 @@ def exclude_cutensor(cls): # pass # _acc.set_reduction_accelerators(red_acc) - # yield + yield # _acc.set_routine_accelerators(old_routine_accelerators) # _acc.set_reduction_accelerators(old_reduction_accelerators) - pass @testing.for_all_dtypes() @testing.numpy_cupy_allclose(contiguous_check=False) @@ -298,15 +300,7 @@ def _axes_for_shape(shape): "shape,axis", [ (shape, axis) - for shape in [ - (), - (0,), - (0, 2), - (2, 0), - (0, 2, 3), - (2, 0, 3), - (2, 3, 0), - ] + for shape in [(), (0,), (0, 2), (2, 0), (0, 2, 3), (2, 0, 3), (2, 3, 0)] for axis in _axes_for_shape(shape) ], ) @@ -331,13 +325,16 @@ def test_zero_size(self, xp, shape, axis, order, func): # This class compares CUB results against NumPy's. ("fallback" is CuPy's # original kernel, also tested here to reduce code duplication.) +# Non-empty shapes keep both the first and last axis >= 128 so the +# contiguous reduction stays on the CUB block-reduction path rather than +# the short-axis fallback. @pytest.mark.parametrize( "shape", [ - (10,), - (10, 20), - (10, 20, 30), - (10, 20, 30, 40), + (128,), + (128, 128), + (128, 2, 128), + (128, 2, 2, 128), # skip (2, 3, 0) because it would not hit the CUB code path (0,), (2, 0), 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 43002cba815..bb548ad810c 100644 --- a/dpnp/tests/third_party/cupy/core_tests/test_raw.py +++ b/dpnp/tests/third_party/cupy/core_tests/test_raw.py @@ -9,7 +9,6 @@ import sys import tempfile import threading -import unittest from unittest import mock import pytest @@ -20,6 +19,7 @@ # from cupy import _util # from cupy._core import _accelerator # from cupy.cuda import compiler +# from cupy.cuda import _compiler_cache # from cupy.cuda import memory # from cupy_backends.cuda.libs import nvrtc @@ -357,11 +357,11 @@ def use_temporary_cache_dir(): # Note uses mock, so not thread-safe (except at class/method level) # tempdir fixture could be used instead. - target1 = "cupy.cuda.compiler._kernel_cache_backend._cache_dir" + target1 = "cupy.cuda.compiler._kernel_cache_backend" target2 = "cupy.cuda.compiler._empty_file_preprocess_cache" temp_cache = {} with tempfile.TemporaryDirectory() as path: - with mock.patch(target1, path): + with mock.patch(target1, _compiler_cache.DiskKernelCacheBackend(path)): with mock.patch(target2, temp_cache): yield path @@ -401,14 +401,101 @@ def find_nvcc_ver(): return int(major) * 1000 + int(minor) * 10 -# TODO(leofang): Further refactor the test suite to avoid using unittest? -class _TestRawBase: +# Recent CCCL has made Jitify cold-launch very slow, see the discussion +# starting https://github.com/cupy/cupy/pull/8899#issuecomment-2613022424. +no_jitify_markers = [ + pytest.mark.filterwarnings( + "ignore:The jitify argument is deprecated:DeprecationWarning" + ) +] +jitify_markers = [ + testing.slow(), + pytest.mark.filterwarnings( + "ignore:jitify=True is deprecated:DeprecationWarning" + ), + pytest.mark.thread_unsafe( + reason="Jitify seems to have problems, skip as largely unmaintained." + ), +] + + +@pytest.mark.parametrize( + "backend,in_memory,clean_up,jitify", + [ + # First test NVRTC + pytest.param( + "nvrtc", + False, + False, + False, + marks=no_jitify_markers, + id="NVRTC (no-jitify)", + ), + # this run will read from in-memory cache + pytest.param( + "nvrtc", + True, + False, + False, + marks=no_jitify_markers, + id="NVRTC (in-memory cache, no-jitify)", + ), + # this run will force recompilation + pytest.param( + "nvrtc", + True, + True, + False, + marks=no_jitify_markers, + id="NVRTC (in-memory cache, re-compile, no-jitify)", + ), + # Finally, we test NVCC + pytest.param( + "nvcc", + False, + False, + False, + marks=no_jitify_markers, + id="NVCC (no-jitify)", + ), + # Below is the same set of NVRTC tests, with Jitify turned on. + # For tests that can already pass, it shouldn't matter whether + # Jitify is on or not, and the only side effect is to add overhead. + # It doesn't make sense to test NVCC + Jitify. + pytest.param( + "nvrtc", + False, + False, + True, + marks=jitify_markers, + id="NVRTC (jitify)", + ), + pytest.param( + "nvrtc", + True, + False, + True, + marks=jitify_markers, + id="NVRTC (in-memory cache, jitify)", + ), + pytest.param( + "nvrtc", + True, + True, + True, + marks=jitify_markers, + id="NVRTC (in-memory cache, re-compile, jitify)", + ), + ], +) +class TestRaw: _nvcc_ver = None _nvrtc_ver = None - def setUp(self): - if getattr(self, "clean_up", False): + @pytest.fixture(autouse=True) + def configure(self, backend, in_memory, clean_up, jitify): + if clean_up: if cupy.cuda.runtime.is_hip: # Clearing memo triggers recompiling kernels using name # expressions in other tests, e.g. dot and matmul, which @@ -418,43 +505,40 @@ def setUp(self): self.dev = cupy.cuda.runtime.getDevice() assert self.dev != 1 - self.jitify = getattr(self, "jitify", False) - if cupy.cuda.runtime.is_hip and self.jitify: + if cupy.cuda.runtime.is_hip and jitify: pytest.skip("Jitify does not support ROCm/HIP") - self.temporary_cache_dir_context = use_temporary_cache_dir() - self.in_memory_context = compile_in_memory(self.in_memory) - self.cache_dir = self.temporary_cache_dir_context.__enter__() - self.in_memory_context.__enter__() - - self.kern = cupy.RawKernel( - _test_source1, "test_sum", backend=self.backend, jitify=self.jitify + kern = cupy.RawKernel( + _test_source1, "test_sum", backend=backend, jitify=jitify ) - self.mod2 = cupy.RawModule( - code=_test_source2, backend=self.backend, jitify=self.jitify + mod2 = cupy.RawModule( + code=_test_source2, backend=backend, jitify=jitify ) - self.mod3 = cupy.RawModule( + mod3 = cupy.RawModule( code=_test_source3, options=("-DPRECISION=2",), - backend=self.backend, - jitify=self.jitify, + backend=backend, + jitify=jitify, ) - def tearDown(self): - if ( - self.in_memory - and _accelerator.ACCELERATOR_CUB - not in _accelerator.get_reduction_accelerators() + with ( + compile_in_memory(in_memory), + use_temporary_cache_dir() as cache_dir, ): - # should not write any file to the cache dir, but the CUB reduction - # kernel uses nvcc, with which I/O cannot be avoided - files = os.listdir(self.cache_dir) - for f in files: - # only test_load_cubin_*.cu files should be present - assert re.match(r"test_load_cubin_(\d+)\.cu", f) - - self.in_memory_context.__exit__(*sys.exc_info()) - self.temporary_cache_dir_context.__exit__(*sys.exc_info()) + yield kern, mod2, mod3, cache_dir + + if ( + in_memory + and _accelerator.ACCELERATOR_CUB + not in _accelerator.get_reduction_accelerators() + ): + # should not write any file to the cache dir, + # but the CUB reduction + # kernel uses nvcc, with which I/O cannot be avoided + files = os.listdir(cache_dir) + for f in files: + # only test_load_cubin_*.cu files should be present + assert re.match(r"test_load_cubin_(\d+)\.cu", f) def _helper(self, kernel, dtype): N = 10 @@ -464,12 +548,14 @@ def _helper(self, kernel, dtype): kernel((N,), (N,), (x1, x2, y, N**2)) return x1, x2, y - def test_basic(self): - x1, x2, y = self._helper(self.kern, cupy.float32) + def test_basic(self, configure): + kern, _, _, _ = configure + x1, x2, y = self._helper(kern, cupy.float32) assert cupy.allclose(y, x1 + x2) - def test_kernel_attributes(self): - attrs = self.kern.attributes + def test_kernel_attributes(self, configure): + kern, _, _, _ = configure + attrs = kern.attributes for attribute in [ "binary_version", "cache_mode_ca", @@ -485,12 +571,13 @@ def test_kernel_attributes(self): assert attribute in attrs # TODO(leofang): investigate why this fails on ROCm 3.5.0 if not cupy.cuda.runtime.is_hip: - assert self.kern.num_regs > 0 - assert self.kern.max_threads_per_block > 0 - assert self.kern.shared_size_bytes == 0 + assert kern.num_regs > 0 + assert kern.max_threads_per_block > 0 + assert kern.shared_size_bytes == 0 - def test_module(self): - module = self.mod2 + def test_module(self, configure): + _, mod2, _, _ = configure + module = mod2 ker_sum = module.get_function("test_sum") ker_times = module.get_function("test_multiply") @@ -500,8 +587,9 @@ def test_module(self): x1, x2, y = self._helper(ker_times, cupy.float32) assert cupy.allclose(y, x1 * x2) - def test_compiler_flag(self): - module = self.mod3 + def test_compiler_flag(self, configure): + _, _, mod3, _ = configure + module = mod3 ker_sum = module.get_function("test_sum") ker_times = module.get_function("test_multiply") @@ -511,11 +599,11 @@ def test_compiler_flag(self): x1, x2, y = self._helper(ker_times, cupy.float64) assert cupy.allclose(y, x1 * x2) - def test_invalid_compiler_flag(self): - if cupy.cuda.runtime.is_hip and self.backend == "nvrtc": - self.skipTest("hiprtc does not handle #error macro properly") + def test_invalid_compiler_flag(self, backend, jitify): + if cupy.cuda.runtime.is_hip and backend == "nvrtc": + pytest.skip("hiprtc does not handle #error macro properly") - if self.jitify: + if jitify: ex_type = cupy.cuda.compiler.JitifyException else: ex_type = cupy.cuda.compiler.CompileException @@ -524,12 +612,12 @@ def test_invalid_compiler_flag(self): mod = cupy.RawModule( code=_test_source3, options=("-DPRECISION=3",), - backend=self.backend, - jitify=self.jitify, + backend=backend, + jitify=jitify, ) mod.get_function("test_sum") # enforce compilation - if not self.jitify: + if not jitify: assert "precision not supported" in str(ex.value) def _find_nvcc_ver(self): @@ -561,7 +649,7 @@ def _check_ptx_loadable(self, compiler: str): if driver_ver < compiler_ver: raise pytest.skip() - def _generate_file(self, ext: str): + def _generate_file(self, cache_dir, ext: str): # generate cubin/ptx by calling nvcc/hipcc if not cupy.cuda.runtime.is_hip: @@ -578,8 +666,8 @@ def _generate_file(self, ext: str): # split() is needed because nvcc could come from the env var NVCC cmd = cc.split() thread_id = threading.get_ident() - source = f"{self.cache_dir}/test_load_cubin_{thread_id}.cu" - file_path = self.cache_dir + f"test_load_cubin_{thread_id}" + source = f"{cache_dir}/test_load_cubin_{thread_id}.cu" + file_path = cache_dir + f"test_load_cubin_{thread_id}" with open(source, "w") as f: f.write(code) if not cupy.cuda.runtime.is_hip: @@ -596,54 +684,61 @@ def _generate_file(self, ext: str): flag = "--genco" cmd += [arch, flag, source, "-o", file_path] cc = "nvcc" if not cupy.cuda.runtime.is_hip else "hipcc" - compiler._run_cc(cmd, self.cache_dir, cc) + compiler._run_cc(cmd, cache_dir, cc) return file_path - @unittest.skipIf(cupy.cuda.runtime.is_hip, "HIP uses hsaco, not cubin") - def test_load_cubin(self): + @pytest.mark.skipif( + cupy.cuda.runtime.is_hip, reason="HIP uses hsaco, not cubin" + ) + def test_load_cubin(self, configure, backend): + _, _, _, cache_dir = configure # generate cubin in the temp dir - file_path = self._generate_file("cubin") + file_path = self._generate_file(cache_dir, "cubin") # load cubin and test the kernel - mod = cupy.RawModule(path=file_path, backend=self.backend) + mod = cupy.RawModule(path=file_path, backend=backend) ker = mod.get_function("test_div") x1, x2, y = self._helper(ker, cupy.float32) assert cupy.allclose(y, x1 / (x2 + 1.0)) - @unittest.skipIf(cupy.cuda.runtime.is_hip, "HIP uses hsaco, not ptx") - def test_load_ptx(self): + @pytest.mark.skipif( + cupy.cuda.runtime.is_hip, reason="HIP uses hsaco, not ptx" + ) + def test_load_ptx(self, configure, backend): + _, _, _, cache_dir = configure # use nvcc to generate ptx in the temp dir self._check_ptx_loadable("nvcc") - file_path = self._generate_file("ptx") + file_path = self._generate_file(cache_dir, "ptx") # load ptx and test the kernel - mod = cupy.RawModule(path=file_path, backend=self.backend) + mod = cupy.RawModule(path=file_path, backend=backend) ker = mod.get_function("test_div") x1, x2, y = self._helper(ker, cupy.float32) assert cupy.allclose(y, x1 / (x2 + 1.0)) - @unittest.skipIf( - not cupy.cuda.runtime.is_hip, "CUDA uses cubin/ptx, not hsaco" + @pytest.mark.skipif( + not cupy.cuda.runtime.is_hip, reason="CUDA uses cubin/ptx, not hsaco" ) - def test_load_hsaco(self): + def test_load_hsaco(self, configure, backend): + _, _, _, cache_dir = configure # generate hsaco in the temp dir - file_path = self._generate_file("hsaco") + file_path = self._generate_file(cache_dir, "hsaco") # load cubin and test the kernel - mod = cupy.RawModule(path=file_path, backend=self.backend) + mod = cupy.RawModule(path=file_path, backend=backend) ker = mod.get_function("test_div") x1, x2, y = self._helper(ker, cupy.float32) assert cupy.allclose(y, x1 / (x2 + 1.0)) - def test_module_load_failure(self): + def test_module_load_failure(self, backend): # in principle this test is better done in test_driver.py, but # this error is more likely to appear when using RawModule, so # let us do it here with pytest.raises(cupy.cuda.driver.CUDADriverError) as ex: mod = cupy.RawModule( path=os.path.expanduser("~/this_does_not_exist.cubin"), - backend=self.backend, + backend=backend, ) mod.get_function("nonexisting_kernel") # enforce loading assert "CUDA_ERROR_FILE_NOT_FOUND" in str( @@ -660,30 +755,31 @@ def test_module_both_code_and_path(self): with pytest.raises(TypeError): cupy.RawModule(code=_test_source1, path="test.cubin") - def test_get_function_failure(self): + def test_get_function_failure(self, configure): + _, mod2, _, _ = configure # in principle this test is better done in test_driver.py, but # this error is more likely to appear when using RawModule, so # let us do it here with pytest.raises(cupy.cuda.driver.CUDADriverError) as ex: - self.mod2.get_function("no_such_kernel") + mod2.get_function("no_such_kernel") assert "CUDA_ERROR_NOT_FOUND" in str( ex.value ) or "hipErrorNotFound" in str( # for CUDA ex.value ) # for HIP - @unittest.skipIf( + @pytest.mark.skipif( cupy.cuda.runtime.is_hip, - "ROCm/HIP does not support dynamic parallelism", + reason="ROCm/HIP does not support dynamic parallelism", ) - def test_dynamical_parallelism(self): + def test_dynamical_parallelism(self, backend, jitify): self._check_ptx_loadable("nvrtc") ker = cupy.RawKernel( _test_source4, "test_kernel", options=("-dc",), - backend=self.backend, - jitify=self.jitify, + backend=backend, + jitify=jitify, ) N = 169 inner_chunk = 13 @@ -691,21 +787,18 @@ def test_dynamical_parallelism(self): ker((1,), (N // inner_chunk,), (x, N, inner_chunk)) assert (x == 1.0).all() - def test_dynamical_parallelism_compile_failure(self): + def test_dynamical_parallelism_compile_failure(self, backend, jitify): # no option for separate compilation is given should cause an error ker = cupy.RawKernel( - _test_source4, - "test_kernel", - backend=self.backend, - jitify=self.jitify, + _test_source4, "test_kernel", backend=backend, jitify=jitify ) N = 10 inner_chunk = 2 x = cupy.zeros((N,), dtype=cupy.float32) use_ptx = os.environ.get("CUPY_COMPILE_WITH_PTX", False) - if self.jitify: + if jitify: error = cupy.cuda.compiler.JitifyException - elif self.backend == "nvrtc" and ( + elif backend == "nvrtc" and ( use_ptx or ( cupy.cuda.driver._is_cuda_python() @@ -724,17 +817,18 @@ def test_dynamical_parallelism_compile_failure(self): with pytest.raises(error): ker((1,), (N // inner_chunk,), (x, N, inner_chunk)) - @unittest.skipIf( - cupy.cuda.runtime.is_hip, "HIP code should not use cuFloatComplex" + @pytest.mark.skipif( + cupy.cuda.runtime.is_hip, + reason="HIP code should not use cuFloatComplex", ) - def test_cuFloatComplex(self): + def test_cuFloatComplex(self, jitify): N = 100 block = 32 grid = (N + block - 1) // block dtype = cupy.complex64 mod = cupy.RawModule( - code=_test_cuComplex, translate_cucomplex=True, jitify=self.jitify + code=_test_cuComplex, translate_cucomplex=True, jitify=jitify ) a = cupy.random.random((N,)) + 1j * cupy.random.random((N,)) a = a.astype(dtype) @@ -789,17 +883,18 @@ def test_cuFloatComplex(self): ker((grid,), (block,), (a, b, out)) assert (out == a + b).all() - @unittest.skipIf( - cupy.cuda.runtime.is_hip, "HIP code should not use cuDoubleComplex" + @pytest.mark.skipif( + cupy.cuda.runtime.is_hip, + reason="HIP code should not use cuDoubleComplex", ) - def test_cuDoubleComplex(self): + def test_cuDoubleComplex(self, jitify): N = 100 block = 32 grid = (N + block - 1) // block dtype = cupy.complex128 mod = cupy.RawModule( - code=_test_cuComplex, translate_cucomplex=True, jitify=self.jitify + code=_test_cuComplex, translate_cucomplex=True, jitify=jitify ) a = cupy.random.random((N,)) + 1j * cupy.random.random((N,)) a = a.astype(dtype) @@ -860,9 +955,9 @@ def test_cuDoubleComplex(self): assert (out == a + b).all() @pytest.mark.thread_unsafe(reason="mutates global in RawModule") - def test_const_memory(self): + def test_const_memory(self, backend, jitify): mod = cupy.RawModule( - code=test_const_mem, backend=self.backend, jitify=self.jitify + code=test_const_mem, backend=backend, jitify=jitify ) ker = mod.get_function("multiply_by_const") mem_ptr = mod.get_global("some_array") @@ -873,15 +968,15 @@ def test_const_memory(self): ker((1,), (100,), (output_arr, cupy.int32(100))) assert (data == output_arr).all() - def test_template_specialization(self): - if self.backend == "nvcc": - self.skipTest("nvcc does not support template specialization") + def test_template_specialization(self, backend, clean_up, jitify): + if backend == "nvcc": + pytest.skip(reason="nvcc does not support template specialization") # TODO(leofang): investigate why hiprtc generates a wrong source code # when the same code is compiled and discarded. It seems hiprtc has # an internal cache that conflicts with the 2nd compilation attempt. - if cupy.cuda.runtime.is_hip and hasattr(self, "clean_up"): - self.skipTest("skip a potential hiprtc bug") + if cupy.cuda.runtime.is_hip and clean_up: + pytest.skip(reason="skip a potential hiprtc bug") # compile code if cupy.cuda.runtime.is_hip: @@ -904,7 +999,7 @@ def test_template_specialization(self): mod = cupy.RawModule( code=test_cxx_template, name_expressions=name_expressions, - jitify=self.jitify, + jitify=jitify, ) dtypes = (cupy.int32, cupy.float32, cupy.complex128, cupy.float64) @@ -928,22 +1023,22 @@ def test_template_specialization(self): # check results assert cupy.allclose(in_arr, out_arr) - def test_template_failure(self): + def test_template_failure(self, backend, jitify): name_expressions = ["my_sqrt"] # 1. nvcc is disabled for this feature - if self.backend == "nvcc": + if backend == "nvcc": with pytest.raises(ValueError) as e: cupy.RawModule( code=test_cxx_template, - backend=self.backend, + backend=backend, name_expressions=name_expressions, ) assert "nvrtc" in str(e.value) return # the rest of tests do not apply to nvcc # 2. compile code without specializations - mod = cupy.RawModule(code=test_cxx_template, jitify=self.jitify) + mod = cupy.RawModule(code=test_cxx_template, jitify=jitify) # ...try to get a specialized kernel match = ( "named symbol not found" @@ -957,7 +1052,7 @@ def test_template_failure(self): mod = cupy.RawModule( code=test_cxx_template, name_expressions=name_expressions, - jitify=self.jitify, + jitify=jitify, ) if cupy.cuda.runtime.is_hip: msg = "hipErrorNotFound" @@ -966,10 +1061,8 @@ def test_template_failure(self): with pytest.raises(cupy.cuda.driver.CUDADriverError, match=msg): mod.get_function("my_sqrt") - def test_raw_pointer(self): - mod = cupy.RawModule( - code=test_cast, backend=self.backend, jitify=self.jitify - ) + def test_raw_pointer(self, backend, jitify): + mod = cupy.RawModule(code=test_cast, backend=backend, jitify=jitify) ker = mod.get_function("my_func") a = cupy.ones((100,), dtype=cupy.float64) @@ -982,21 +1075,23 @@ def test_raw_pointer(self): assert (a == b).all() @testing.multi_gpu(2) - def test_context_switch_RawKernel(self): + def test_context_switch_RawKernel(self, configure): + kern, _, _, _ = configure # run test_basic() on another device # we need to launch it once to force compiling - x1, x2, y = self._helper(self.kern, cupy.float32) + x1, x2, y = self._helper(kern, cupy.float32) with cupy.cuda.Device(1): - x1, x2, y = self._helper(self.kern, cupy.float32) + x1, x2, y = self._helper(kern, cupy.float32) assert cupy.allclose(y, x1 + x2) @testing.multi_gpu(2) - def test_context_switch_RawModule1(self): + def test_context_switch_RawModule1(self, configure): + _, mod2, _, _ = configure # run test_module() on another device # in this test, re-compiling happens at 2nd get_function() - module = self.mod2 + module = mod2 with cupy.cuda.Device(0): module.get_function("test_sum") @@ -1006,10 +1101,11 @@ def test_context_switch_RawModule1(self): assert cupy.allclose(y, x1 + x2) @testing.multi_gpu(2) - def test_context_switch_RawModule2(self): + def test_context_switch_RawModule2(self, configure): + _, mod2, _, _ = configure # run test_module() on another device # in this test, re-compiling happens at kernel launch - module = self.mod2 + module = mod2 with cupy.cuda.Device(0): ker_sum = module.get_function("test_sum") @@ -1018,7 +1114,8 @@ def test_context_switch_RawModule2(self): assert cupy.allclose(y, x1 + x2) @testing.multi_gpu(2) - def test_context_switch_RawModule3(self): + def test_context_switch_RawModule3(self, configure, backend): + _, _, _, cache_dir = configure # run test_load_cubin() on another device # generate cubin in the temp dir and load it on device 0 @@ -1028,8 +1125,8 @@ def test_context_switch_RawModule3(self): raise pytest.skip() with device0: - file_path = self._generate_file("cubin") - mod = cupy.RawModule(path=file_path, backend=self.backend) + file_path = self._generate_file(cache_dir, "cubin") + mod = cupy.RawModule(path=file_path, backend=backend) mod.get_function("test_div") # in this test, reloading happens at 2nd get_function() @@ -1039,7 +1136,8 @@ def test_context_switch_RawModule3(self): assert cupy.allclose(y, x1 / (x2 + 1.0)) @testing.multi_gpu(2) - def test_context_switch_RawModule4(self): + def test_context_switch_RawModule4(self, configure, backend): + _, _, _, cache_dir = configure # run test_load_cubin() on another device # generate cubin in the temp dir and load it on device 0 @@ -1049,8 +1147,8 @@ def test_context_switch_RawModule4(self): raise pytest.skip() with device0: - file_path = self._generate_file("cubin") - mod = cupy.RawModule(path=file_path, backend=self.backend) + file_path = self._generate_file(cache_dir, "cubin") + mod = cupy.RawModule(path=file_path, backend=backend) ker = mod.get_function("test_div") # in this test, reloading happens at kernel launch @@ -1059,11 +1157,11 @@ def test_context_switch_RawModule4(self): assert cupy.allclose(y, x1 / (x2 + 1.0)) @testing.multi_gpu(2) - def test_context_switch_RawModule5(self): + def test_context_switch_RawModule5(self, backend, jitify): # run test_template_specialization() on another device # in this test, re-compiling happens at get_function() - if self.backend == "nvcc": - self.skipTest("nvcc does not support template specialization") + if backend == "nvcc": + pytest.skip(reason="nvcc does not support template specialization") # compile code name_expressions = ["my_sqrt"] @@ -1072,7 +1170,7 @@ def test_context_switch_RawModule5(self): mod = cupy.RawModule( code=test_cxx_template, name_expressions=name_expressions, - jitify=self.jitify, + jitify=jitify, ) # get specialized kernels @@ -1094,11 +1192,11 @@ def test_context_switch_RawModule5(self): assert cupy.allclose(in_arr, out_arr) @testing.multi_gpu(2) - def test_context_switch_RawModule6(self): + def test_context_switch_RawModule6(self, backend, jitify): # run test_template_specialization() on another device # in this test, re-compiling happens at kernel launch - if self.backend == "nvcc": - self.skipTest("nvcc does not support template specialization") + if backend == "nvcc": + pytest.skip(reason="nvcc does not support template specialization") # compile code name_expressions = ["my_sqrt"] @@ -1107,7 +1205,7 @@ def test_context_switch_RawModule6(self): mod = cupy.RawModule( code=test_cxx_template, name_expressions=name_expressions, - jitify=self.jitify, + jitify=jitify, ) # get specialized kernels @@ -1125,17 +1223,17 @@ def test_context_switch_RawModule6(self): # check results assert cupy.allclose(in_arr, out_arr) - @unittest.skipUnless( - not cupy.cuda.runtime.is_hip, "only CUDA raises warning" + @pytest.mark.skipif( + cupy.cuda.runtime.is_hip, reason="only CUDA raises warning" ) @pytest.mark.thread_unsafe(reason="mutates global cache directory") - def test_compile_kernel(self): + def test_compile_kernel(self, backend, jitify): kern = cupy.RawKernel( _test_compile_src, "test_op", options=("-DOP=+",), - backend=self.backend, - jitify=self.jitify, + backend=backend, + jitify=jitify, ) log = io.StringIO() with use_temporary_cache_dir(): @@ -1144,16 +1242,16 @@ def test_compile_kernel(self): x1, x2, y = self._helper(kern, cupy.float32) assert cupy.allclose(y, x1 + x2) - @unittest.skipUnless( - not cupy.cuda.runtime.is_hip, "only CUDA raises warning" + @pytest.mark.skipif( + cupy.cuda.runtime.is_hip, reason="only CUDA raises warning" ) @pytest.mark.thread_unsafe(reason="mutates global cache directory") - def test_compile_module(self): + def test_compile_module(self, backend, jitify): module = cupy.RawModule( code=_test_compile_src, - backend=self.backend, + backend=backend, options=("-DOP=+",), - jitify=self.jitify, + jitify=jitify, ) log = io.StringIO() with use_temporary_cache_dir(): @@ -1164,42 +1262,6 @@ def test_compile_module(self): assert cupy.allclose(y, x1 + x2) -@testing.parameterize( - # First test NVRTC - {"backend": "nvrtc", "in_memory": False}, - # this run will read from in-memory cache - {"backend": "nvrtc", "in_memory": True}, - # this run will force recompilation - {"backend": "nvrtc", "in_memory": True, "clean_up": True}, - # Finally, we test NVCC - {"backend": "nvcc", "in_memory": False}, -) -@pytest.mark.filterwarnings("ignore:.*jitify=False:DeprecationWarning") -class TestRaw(_TestRawBase, unittest.TestCase): - pass - - -# Recent CCCL has made Jitify cold-launch very slow, see the discussion -# starting https://github.com/cupy/cupy/pull/8899#issuecomment-2613022424. -# TODO(leofang): Further refactor the test suite? -@testing.parameterize( - # Below is the same set of NVRTC tests, with Jitify turned on. For tests - # that can already pass, it shouldn't matter whether Jitify is on or not, - # and the only side effect is to add overhead. It doesn't make sense to - # test NVCC + Jitify. - {"backend": "nvrtc", "in_memory": False, "jitify": True}, - {"backend": "nvrtc", "in_memory": True, "jitify": True}, - {"backend": "nvrtc", "in_memory": True, "clean_up": True, "jitify": True}, -) -@testing.slow -@pytest.mark.thread_unsafe( - reason="Jitify seems to have problems, skip as largely unmaintained." -) -@pytest.mark.filterwarnings("ignore:jitify=True:DeprecationWarning") -class TestRawWithJitify(_TestRawBase, unittest.TestCase): - pass - - _test_grid_sync = r""" #include @@ -1220,29 +1282,24 @@ class TestRawWithJitify(_TestRawBase, unittest.TestCase): """ -@testing.parameterize( - *testing.product( - { - "n": [10, 100, 1000], - "block": [64, 256], - } - ) -) -@unittest.skipIf( +@pytest.mark.parametrize("n", [10, 100, 1000]) +@pytest.mark.parametrize("block", [64, 256]) +@pytest.mark.skipif( cupy.cuda.runtime.is_hip or find_nvcc_ver() >= 12020, - "fp16 header compatibility issue, see cupy#8412 (Skip on HIP)", + reason="fp16 header compatibility issue, see cupy#8412 (Skip on HIP)", ) -@unittest.skipUnless( - 9000 <= cupy.cuda.runtime.runtimeGetVersion(), "Requires CUDA 9.x or later" +@pytest.mark.skipif( + 9000 > cupy.cuda.runtime.runtimeGetVersion(), + reason="Requires CUDA 9.x or later", ) -@unittest.skipUnless( - 60 <= int(cupy.cuda.device.get_compute_capability()), - "Requires compute capability 6.0 or later", +@pytest.mark.skipif( + 60 > int(cupy.cuda.device.get_compute_capability()), + reason="Requires compute capability 6.0 or later", ) -class TestRawGridSync(unittest.TestCase): +class TestRawGridSync: + @pytest.mark.thread_unsafe(reason="mutates global cache directory") - def test_grid_sync_rawkernel(self): - n = self.n + def test_grid_sync_rawkernel(self, n, block): with use_temporary_cache_dir(): kern_grid_sync = cupy.RawKernel( _test_grid_sync, @@ -1253,14 +1310,12 @@ def test_grid_sync_rawkernel(self): x1 = cupy.arange(n**2, dtype="float32").reshape(n, n) x2 = cupy.ones((n, n), dtype="float32") y = cupy.zeros((n, n), dtype="float32") - block = self.block grid = (n * n + block - 1) // block kern_grid_sync((grid,), (block,), (x1, x2, y, n**2)) assert cupy.allclose(y, x1 + x2) @pytest.mark.thread_unsafe(reason="mutates global cache directory") - def test_grid_sync_rawmodule(self): - n = self.n + def test_grid_sync_rawmodule(self, n, block): with use_temporary_cache_dir(): mod_grid_sync = cupy.RawModule( code=_test_grid_sync, @@ -1271,7 +1326,6 @@ def test_grid_sync_rawmodule(self): x2 = cupy.ones((n, n), dtype="float32") y = cupy.zeros((n, n), dtype="float32") kern = mod_grid_sync.get_function("test_grid_sync") - block = self.block grid = (n * n + block - 1) // block kern((grid,), (block,), (x1, x2, y, n**2)) assert cupy.allclose(y, x1 + x2) @@ -1303,82 +1357,83 @@ def test_grid_sync_rawmodule(self): # Pickling/unpickling a RawModule should always success, whereas # pickling/unpickling a RawKernel would fail if we don't enforce # recompiling after unpickling it. -@testing.parameterize( - *testing.product( - { - "compile": (False, True), - "raw": ("ker", "mod", "mod_ker"), - } - ) +@pytest.mark.parametrize( + "compile", + [ + pytest.param(False, id="nocompile"), + pytest.param(True, id="compile"), + ], ) -@unittest.skipUnless( - 60 <= int(cupy.cuda.device.get_compute_capability()), - "Requires compute capability 6.0 or later", +@pytest.mark.parametrize("raw", ["ker", "mod", "mod_ker"]) +@pytest.mark.skipif( + 60 > int(cupy.cuda.device.get_compute_capability()), + reason="Requires compute capability 6.0 or later", ) -@unittest.skipIf( - cupy.cuda.runtime.is_hip, "HIP does not support enable_cooperative_groups" +@pytest.mark.skipif( + cupy.cuda.runtime.is_hip, + reason="HIP does not support enable_cooperative_groups", ) -class TestRawPicklable(unittest.TestCase): - def setUp(self): - self.temporary_dir_context = use_temporary_cache_dir() - self.temp_dir = self.temporary_dir_context.__enter__() - +class TestRawPicklable: + @pytest.fixture(autouse=True) + def configure(self, compile, raw): # test if kw-only arguments are properly handled or not - if self.raw == "ker": - self.ker = cupy.RawKernel( + ker, mod = None, None + if raw == "ker": + ker = cupy.RawKernel( _test_source1, "test_sum", backend="nvcc", enable_cooperative_groups=True, ) else: - self.mod = cupy.RawModule( + mod = cupy.RawModule( code=_test_source1, backend="nvcc", enable_cooperative_groups=True, ) - def tearDown(self): - self.temporary_dir_context.__exit__(*sys.exc_info()) + with use_temporary_cache_dir() as temp_dir: + yield compile, raw, ker, mod, temp_dir - def _helper(self): + def _helper(self, raw, ker, mod): N = 10 x1 = cupy.arange(N**2, dtype=cupy.float32).reshape(N, N) x2 = cupy.ones((N, N), dtype=cupy.float32) y = cupy.zeros((N, N), dtype=cupy.float32) - if self.raw == "ker": - ker = self.ker + + if raw == "ker": + impl = ker else: - ker = self.mod.get_function("test_sum") - ker((N,), (N,), (x1, x2, y, N**2)) + impl = mod.get_function("test_sum") + impl((N,), (N,), (x1, x2, y, N**2)) assert cupy.allclose(x1 + x2, y) - def test_raw_picklable(self): + def test_raw_picklable(self, configure): + compile, raw, ker, mod, temp_dir = configure # force compiling before pickling - if self.compile: - self._helper() + if compile: + self._helper(raw, ker, mod) - if self.raw == "ker": + if raw == "ker": # pickle the RawKernel - obj = self.ker - elif self.raw == "mod": + obj = ker + elif raw == "mod": # pickle the RawModule - obj = self.mod - elif self.raw == "mod_ker": + obj = mod + elif raw == "mod_ker": # pickle the RawKernel fetched from the RawModule - obj = self.mod.get_function("test_sum") - with open(self.temp_dir + "/raw.pkl", "wb") as f: + obj = mod.get_function("test_sum") + with open(temp_dir + "/raw.pkl", "wb") as f: pickle.dump(obj, f) # dump test script to temp dir - with open(self.temp_dir + "/TestRawPicklable.py", "w") as f: + with open(temp_dir + "/TestRawPicklable.py", "w") as f: f.write(_test_script) - test_args = ["test_sum"] if self.raw == "mod" else [] + test_args = ["test_sum"] if raw == "mod" else [] # run another process to check the pickle s = subprocess.run( - [sys.executable, "TestRawPicklable.py"] + test_args, - cwd=self.temp_dir, + [sys.executable, "TestRawPicklable.py"] + test_args, cwd=temp_dir ) s.check_returncode() # raise if unsuccessful @@ -1400,20 +1455,28 @@ def test_raw_picklable(self): # Recent CCCL has made Jitify cold-launch very slow, see the discussion # starting https://github.com/cupy/cupy/pull/8899#issuecomment-2613022424. -# TODO(leofang): Further refactor the test suite? -class _TestRawJitify: - def setUp(self): - self.temporary_dir_context = use_temporary_cache_dir() - self.temp_dir = self.temporary_dir_context.__enter__() - - def tearDown(self): - self.temporary_dir_context.__exit__(*sys.exc_info()) +@testing.slow +@pytest.mark.parametrize( + "jitify", + [ + pytest.param(False, marks=no_jitify_markers, id="no-jitify"), + pytest.param(True, marks=jitify_markers, id="jitify"), + ], +) +@pytest.mark.skipif( + cupy.cuda.runtime.is_hip, reason="Jitify does not support ROCm/HIP" +) +class TestRawJitify: + @pytest.fixture(autouse=True) + def configure(self): + with use_temporary_cache_dir() as temp_dir: + yield temp_dir - def _helper(self, header, options=()): + def _helper(self, jitify, header, options=()): code = header code += _test_source1 mod1 = cupy.RawModule( - code=code, backend="nvrtc", options=options, jitify=self.jitify + code=code, backend="nvrtc", options=options, jitify=jitify ) N = 10 @@ -1424,10 +1487,10 @@ def _helper(self, header, options=()): ker((N,), (N,), (x1, x2, y, N**2)) assert cupy.allclose(x1 + x2, y) - def _helper2(self, type_str): + def _helper2(self, jitify, type_str): mod2 = cupy.RawModule( code=std_code, - jitify=self.jitify, + jitify=jitify, name_expressions=("shift<%s>" % type_str,), ) ker = mod2.get_function("shift<%s>" % type_str) @@ -1437,7 +1500,7 @@ def _helper2(self, type_str): ker((1,), (N,), (a, N)) assert cupy.allclose(a, b + 100) - def test_jitify1(self): + def test_jitify1(self, jitify): # simply prepend an unused header hdr = "#include \n" # Starting CUDA 12.2, fp16/bf16 headers are intertwined, but due to @@ -1447,91 +1510,76 @@ def test_jitify1(self): options = ("-DCUB_DISABLE_BF16_SUPPORT",) # Compiling CUB headers now works with or without Jitify. - self._helper(hdr, options) + self._helper(jitify, hdr, options) - def test_jitify2(self): + def test_jitify2(self, jitify): # NVRTC cannot compile any code involving std - if self.jitify: + if jitify: # Jitify will make it work - self._helper2("int") + self._helper2(jitify, "int") else: with pytest.raises(cupy.cuda.compiler.CompileException) as ex: - self._helper2("int") + self._helper2(jitify, "int") assert "cannot open source file" in str(ex.value) - def test_jitify3(self): + def test_jitify3(self, jitify): # We supply a type impossible to specialize. Jitify is still able to # locate the headers, but when it comes to the actual compilation, # NVRTC fails (raising the same exception) with different error # messages. ex_type = cupy.cuda.compiler.CompileException with pytest.raises(ex_type) as ex: - self._helper2("float") - if self.jitify: + self._helper2(jitify, "float") + if jitify: assert "Error in parsing name expression" in str(ex.value) else: assert "cannot open source file" in str(ex.value) - def test_jitify4(self): + def test_jitify4(self, jitify): # ensure JitifyException is raised with a broken code code = r""" __global__ void i_am_broken() { """ - if self.jitify: + if jitify: ex_type = cupy.cuda.compiler.JitifyException else: ex_type = cupy.cuda.compiler.CompileException with pytest.raises(ex_type): - mod = cupy.RawModule(code=code, jitify=self.jitify) + mod = cupy.RawModule(code=code, jitify=jitify) ker = mod.get_function("i_am_broken") # if Jitify could redirect its output, we would be able to check # the error log here as well (NVIDIA/jitify#79) - def test_jitify5(self): + def test_jitify5(self, configure, jitify): + temp_dir = configure # If including a header that does not exist, Jitify would attempt to # comment it out and proceed. If this header is actually unused, then # everything would run just fine. hdr = "I_INCLUDE_SOMETHING.h" - with open(self.temp_dir + "/" + hdr, "w") as f: + with open(temp_dir + "/" + hdr, "w") as f: dummy = "#include \n" f.write(dummy) hdr = '#include "' + hdr + '"\n' - if self.jitify: + if jitify: # Jitify would print a warning "[jitify] File not found" to stdout, # but as mentioned above and elsewhere, we can't capture it. - self._helper(hdr, options=("-I" + self.temp_dir,)) + self._helper(jitify, hdr, options=("-I" + temp_dir,)) else: with pytest.raises(cupy.cuda.compiler.CompileException) as ex: - self._helper(hdr, options=("-I" + self.temp_dir,)) + self._helper(jitify, hdr, options=(f"-I{temp_dir}",)) assert "cannot open source file" in str(ex.value) -@unittest.skipIf(cupy.cuda.runtime.is_hip, "Jitify does not support ROCm/HIP") -@testing.slow -@pytest.mark.filterwarnings("ignore:.*jitify=False:DeprecationWarning") -class TestRawJitifyNoJitify(_TestRawJitify, unittest.TestCase): - jitify = False - - -@unittest.skipIf(cupy.cuda.runtime.is_hip, "Jitify does not support ROCm/HIP") -@testing.slow -@pytest.mark.thread_unsafe( - reason="Jitify seems to have problems, skip as largely unmaintained." -) -@pytest.mark.filterwarnings("ignore:jitify=True:DeprecationWarning") -class TestRawJitifyJitify(_TestRawJitify, unittest.TestCase): - jitify = True - - @pytest.mark.parametrize( - "jitify,match", - [(True, ".*"), (False, "Avoid passing.*jitify=False")], + "jitify,match", [(True, ".*"), (False, "Avoid passing.*jitify=False")] +) +@pytest.mark.skipif( + cupy.cuda.runtime.is_hip, reason="Jitify does not support ROCm/HIP" ) -@unittest.skipIf(cupy.cuda.runtime.is_hip, "Jitify does not support ROCm/HIP") @testing.slow @pytest.mark.thread_unsafe(reason="uses temporary cache dir") @use_temporary_cache_dir() @@ -1546,5 +1594,4 @@ def test_jitify_deprecation_warning(jitify, match): # Not technically part of the rawkernel, but test warning in compile here: with pytest.warns(DeprecationWarning, match=match): - # compiler is not imported in dpnp (module is skipped) compiler.compile_using_nvrtc("", options=(), jitify=jitify) diff --git a/dpnp/tests/third_party/cupy/math_tests/test_sumprod.py b/dpnp/tests/third_party/cupy/math_tests/test_sumprod.py index 7fc5a3b80d2..d743e14775c 100644 --- a/dpnp/tests/third_party/cupy/math_tests/test_sumprod.py +++ b/dpnp/tests/third_party/cupy/math_tests/test_sumprod.py @@ -1,3 +1,5 @@ +from __future__ import annotations + import math from itertools import product as iproduct @@ -5,6 +7,10 @@ import pytest import dpnp as cupy + +# import cupy._core._accelerator as _acc +# import cupy.cuda.cutensor +# from cupy._core import _cub_reduction from dpnp.exceptions import AxisError from dpnp.tests.helper import ( has_support_aspect16, @@ -14,13 +20,12 @@ class TestSumprod: - @pytest.fixture(autouse=True) def tearDown(self): + yield # Free huge memory for slow test # cupy.get_default_memory_pool().free_all_blocks() # cupy.get_default_pinned_memory_pool().free_all_blocks() - pass @testing.for_all_dtypes() @testing.numpy_cupy_allclose() @@ -41,7 +46,7 @@ def test_external_sum_all(self, xp, dtype): return xp.sum(a) @testing.for_all_dtypes() - @testing.numpy_cupy_allclose(rtol=1e-06) + @testing.numpy_cupy_allclose(rtol=1e-6) def test_sum_all2(self, xp, dtype): a = testing.shaped_arange((20, 30, 40), xp, dtype) return a.sum() @@ -53,7 +58,7 @@ def test_sum_all_transposed(self, xp, dtype): return a.sum() @testing.for_all_dtypes() - @testing.numpy_cupy_allclose(rtol=1e-06) + @testing.numpy_cupy_allclose(rtol=1e-6) def test_sum_all_transposed2(self, xp, dtype): a = testing.shaped_arange((20, 30, 40), xp, dtype).transpose(2, 0, 1) return a.sum() @@ -66,7 +71,6 @@ def test_sum_axis(self, xp, dtype): @testing.slow @testing.numpy_cupy_allclose() - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="too large allocations") def test_sum_axis_huge(self, xp): a = testing.shaped_random((2048, 1, 1024), xp, "b") @@ -208,12 +212,32 @@ def test_prod_dtype(self, xp, src_dtype, dst_dtype): return a.prod(dtype=dst_dtype) -# This class compares CUB results against NumPy's +# This class compares CUB results against NumPy's. +# Use _min_cub to make sure that the CUB path is used on these files +# _MIN_CUB = _cub_reduction._CUB_REDUCE_SIZE_THRESHOLD +_MIN_CUB = 0 + + @pytest.mark.parametrize( - "shape", [(10,), (10, 20), (10, 20, 30), (10, 20, 30, 40)] + "shape", + [ + (_MIN_CUB,), + (_MIN_CUB, _MIN_CUB), + (_MIN_CUB, 2, _MIN_CUB), + (_MIN_CUB, 2, 2, _MIN_CUB), + ], ) -@pytest.mark.parametrize("order", ["C", "F"]) -@pytest.mark.parametrize("backend", ["device", "block"]) +@pytest.mark.parametrize( + "order", + ["C", "F"], +) +@pytest.mark.parametrize( + "backend", + ["device", "block"], +) +# @pytest.mark.skipif( +# not cupy.cuda.cub.available, reason="The CUB routine is not enabled" +# ) @pytest.mark.skip("_cub_reduction is not supported") class TestCubReduction: @@ -231,12 +255,11 @@ def setUp(self, backend): _acc.set_routine_accelerators(old_routine_accelerators) _acc.set_reduction_accelerators(old_reduction_accelerators) - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp - # @pytest.mark.thread_unsafe(reason="unsafe AssertFunctionIsCalled.") @testing.for_contiguous_axes() # sum supports less dtypes; don't test float16 as it's not as accurate? @testing.for_dtypes("qQfdFD") @testing.numpy_cupy_allclose(rtol=1e-5) + # @pytest.mark.thread_unsafe(reason="unsafe AssertFunctionIsCalled.") def test_cub_sum(self, xp, dtype, axis, shape, order, backend): a = testing.shaped_random(shape, xp, dtype) if order in ("c", "C"): @@ -284,7 +307,6 @@ def test_cub_sum_empty_axis(self, xp, dtype, shape, order, backend): a = xp.asfortranarray(a) return a.sum(axis=()) - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="unsafe AssertFunctionIsCalled.") @testing.for_contiguous_axes() # prod supports less dtypes; don't test float16 as it's not as accurate? @@ -328,7 +350,6 @@ def test_cub_prod(self, xp, dtype, axis, shape, order, backend): # TODO(leofang): test axis after support is added # don't test float16 as it's not as accurate? - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="unsafe AssertFunctionIsCalled.") @testing.for_dtypes("bhilBHILfdFD") @testing.numpy_cupy_allclose(rtol=1e-4) @@ -355,7 +376,6 @@ def test_cub_cumsum(self, xp, dtype, shape, order, backend): # TODO(leofang): test axis after support is added # don't test float16 as it's not as accurate? - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="unsafe AssertFunctionIsCalled.") @testing.for_dtypes("bhilBHILfdFD") @testing.numpy_cupy_allclose(rtol=1e-4) @@ -383,7 +403,7 @@ def test_cub_cumprod(self, xp, dtype, shape, order, backend): return self._mitigate_cumprod(xp, dtype, result) def _mitigate_cumprod(self, xp, dtype, result): - # for testing cumprod against complex arrays, the catch is CuPy may + # for testing cumprod against complex arrays, the got you is CuPy may # produce only Inf at the position where NumPy starts to give NaN. So, # an error would be raised during assert_allclose where the positions # of NaNs are examined. Since this is both algorithm and architecture @@ -398,10 +418,9 @@ def _mitigate_cumprod(self, xp, dtype, result): INT32_MAX = numpy.iinfo(numpy.int32).max -# CUB is not supported by dpnp; the original skipif on cupy.cuda.cub.available -# cannot be evaluated (dpnp has no cupy.cuda), so skip unconditionally. # @pytest.mark.skipif( -# not cupy.cuda.cub.available, reason="The CUB routine is not enabled") +# not cupy.cuda.cub.available, reason="The CUB routine is not enabled" +# ) @pytest.mark.skip("CUB reduction is not supported") @testing.slow class TestReductionSizeOverInt32Max: @@ -458,8 +477,7 @@ def test_reduce(self, shape, axis, dtype, part): a.max(axis=axis), cupy.full(s.shape, 3, dtype=dtype) ) testing.assert_array_equal( - a.argmin(axis), - cupy.full(s.shape, a.shape[axis] - 1), + a.argmin(axis), cupy.full(s.shape, a.shape[axis] - 1) ) else: if axis is None: @@ -514,9 +532,17 @@ def test_cumprod_size_over_int32_max(self, dtype): # This class compares cuTENSOR results against NumPy's @pytest.mark.parametrize( - "shape", [(10,), (10, 20), (10, 20, 30), (10, 20, 30, 40)] + "shape", + [(10,), (10, 20), (10, 20, 30), (10, 20, 30, 40)], ) -@pytest.mark.parametrize("order", ["C", "F"]) +@pytest.mark.parametrize( + "order", + ["C", "F"], +) +# @pytest.mark.skipif( +# not cupy.cuda.cutensor.available, +# reason="The cuTENSOR routine is not enabled", +# ) @pytest.mark.skip("cutensor is not supported") class TestCuTensorReduction: @@ -528,6 +554,7 @@ def setup(cls): yield cupy._core.set_routine_accelerators(old_accelerators) + # @pytest.mark.thread_unsafe(reason="unsafe AssertFunctionIsCalled.") @testing.for_contiguous_axes() # sum supports less dtypes; don't test float16 as it's not as accurate? @testing.for_dtypes("qQfdFD") @@ -868,12 +895,15 @@ def test_ndarray_cumprod_2dim_with_axis(self, xp, dtype): @testing.slow def test_cumprod_huge_array(self): size = 2**32 - a = cupy.ones(size, dtype="b") + # Free huge memory for slow test + cupy.get_default_memory_pool().free_all_blocks() + a = cupy.ones(size, "b") result = cupy.cumprod(a, dtype="b") del a assert (result == 1).all() # Free huge memory for slow test del result + cupy.get_default_memory_pool().free_all_blocks() @testing.for_all_dtypes() def test_invalid_axis_lower1(self, dtype): @@ -1327,3 +1357,130 @@ def test_trapz_1dim_with_x_and_dx(self, xp, dtype): a = testing.shaped_arange((5,), xp, dtype) x = testing.shaped_arange((5,), xp, dtype) return xp.trapezoid(a, x=x, dx=0.1) + + +@testing.with_requires("numpy>=2.1") +@pytest.mark.parametrize("func", ["cumulative_sum", "cumulative_prod"]) +class TestCumulativeSumProd: + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_1d(self, xp, dtype, func): + a = testing.shaped_arange((5,), xp, dtype) + return getattr(xp, func)(a) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6, contiguous_check=False) + def test_axis(self, xp, dtype, func): + a = testing.shaped_arange((3, 4, 5), xp, dtype) + return getattr(xp, func)(a, axis=1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_negative_axis(self, xp, dtype, func): + a = testing.shaped_arange((3, 4, 5), xp, dtype) + return getattr(xp, func)(a, axis=-1) + + @testing.for_all_dtypes(no_bool=True) + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_dtype(self, xp, dtype, func): + a = testing.shaped_arange((5,), xp, numpy.int16) + return getattr(xp, func)(a, dtype=dtype) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_1d_include_initial(self, xp, dtype, func): + a = testing.shaped_arange((5,), xp, dtype) + return getattr(xp, func)(a, include_initial=True) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6, contiguous_check=False) + def test_axis_include_initial(self, xp, dtype, func): + a = testing.shaped_arange((3, 4, 5), xp, dtype) + return getattr(xp, func)(a, axis=1, include_initial=True) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_negative_axis_include_initial(self, xp, dtype, func): + a = testing.shaped_arange((3, 4, 5), xp, dtype) + return getattr(xp, func)(a, axis=-1, include_initial=True) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_out(self, xp, dtype, func): + a = testing.shaped_arange((3, 4, 5), xp, dtype) + out = xp.zeros((3, 4, 5), dtype=dtype) + getattr(xp, func)(a, axis=1, out=out) + return out + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_out_include_initial(self, xp, dtype, func): + a = testing.shaped_arange((3, 4, 5), xp, dtype) + out = xp.zeros((3, 5, 5), dtype=dtype) + getattr(xp, func)(a, axis=1, out=out, include_initial=True) + return out + + @testing.for_all_dtypes() + def test_axis_none_requires_1d(self, dtype, func): + for xp in (numpy, cupy): + a = testing.shaped_arange((3, 4), xp, dtype) + with pytest.raises(ValueError): + getattr(xp, func)(a) + + @testing.for_all_dtypes() + def test_invalid_axis(self, dtype, func): + for xp in (numpy, cupy): + a = testing.shaped_arange((3, 4), xp, dtype) + with pytest.raises(AxisError): + getattr(xp, func)(a, axis=3) + + def test_out_shape_mismatch(self, func): + a = testing.shaped_arange((3, 4), cupy, numpy.float32) + out = cupy.zeros((3, 5), dtype=numpy.float32) + with pytest.raises(ValueError): + getattr(cupy, func)(a, axis=1, out=out) + + @pytest.mark.skip("implementation-defined acc to Python array API") + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_0d(self, xp, dtype, func): + # numpy accepts 0-D input; cupy matches via atleast_1d. + a = xp.asarray(3, dtype=dtype) + return getattr(xp, func)(a) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_0d_include_initial(self, xp, dtype, func): + a = xp.asarray(3, dtype=dtype) + return getattr(xp, func)(a, include_initial=True) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_empty_axis(self, xp, dtype, func): + a = xp.empty((3, 0, 4), dtype=dtype) + return getattr(xp, func)(a, axis=1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_empty_axis_include_initial(self, xp, dtype, func): + a = xp.empty((3, 0, 4), dtype=dtype) + return getattr(xp, func)(a, axis=1, include_initial=True) + + @pytest.mark.parametrize("in_dtype", [numpy.int8, numpy.uint16]) + @testing.numpy_cupy_array_equal() + def test_narrow_int_promotion_include_initial(self, xp, in_dtype, func): + # Narrow integer inputs must promote to int64 / uint64 (matches + # numpy's default-platform-integer rule) even when include_initial + # takes the internal-allocation path. + a = testing.shaped_arange((5,), xp, in_dtype) + return getattr(xp, func)(a, include_initial=True) + + @testing.for_all_dtypes(no_bool=True) + @testing.numpy_cupy_allclose(rtol=1e-6) + def test_out_dtype_cast(self, xp, dtype, func): + # out has a different dtype than x -- result must be cast. + a = testing.shaped_arange((3, 4, 5), xp, numpy.int16) + out = xp.zeros((3, 4, 5), dtype=dtype) + getattr(xp, func)(a, axis=1, out=out) + return out diff --git a/dpnp/tests/third_party/cupy/random_tests/test_generator_api.py b/dpnp/tests/third_party/cupy/random_tests/test_generator_api.py index 1e03d0a4246..4f7a5d9783b 100644 --- a/dpnp/tests/third_party/cupy/random_tests/test_generator_api.py +++ b/dpnp/tests/third_party/cupy/random_tests/test_generator_api.py @@ -1,3 +1,5 @@ +from __future__ import annotations + import threading import unittest @@ -231,6 +233,50 @@ def test_integers_ks_large2(self): self.check_ks(0.05)(2**40, size=2000) +@testing.parameterize( + *[ + # Check values around, 2**31 < high <= 2**32 (used to error). + {"high": 2**31, "endpoint": False}, + {"high": 2**31 + 1, "endpoint": False}, + {"high": 2**32, "endpoint": False}, # full uint32 path + {"high": 2**31, "endpoint": True}, + ] +) +@testing.fix_random() +class TestIntegersLargeBound(GeneratorTestCase): + target_method = "integers" + + def test_integers_large_bound(self): + out = self.generate( + 0, self.high, size=10000, dtype=numpy.int64, endpoint=self.endpoint + ) + assert out.dtype == numpy.int64 + assert (0 <= out).all() + if self.endpoint: + assert (out <= self.high).all() + else: + assert (out < self.high).all() + + +@testing.parameterize( + *[ + # The spans are 1.5x a power of two, so the mask is one bit wider than + # the span and the rejection loop is exercised. The 32-bit span also + # exceeds 2**31, the range that used to overflow the kernel parameter. + {"low": -(2**30), "high": 2**31 - 1}, + {"low": -(2**61), "high": 2**62 - 1}, + ] +) +@testing.with_requires("numpy>=1.17.0") +@testing.fix_random() +class TestIntegersLargeBoundKS(GeneratorTestCase): + target_method = "integers" + + @_condition.repeat_with_success_at_least(10, 3) + def test_integers_large_bound_ks(self): + self.check_ks(0.05)(self.low, self.high, size=2000) + + @testing.with_requires("numpy>=1.17.0") @testing.fix_random() class TestRandom(InvalidOutsMixin, GeneratorTestCase): @@ -376,7 +422,6 @@ class TestDrichlet(common_distributions.Dirichlet, GeneratorTestCase): @testing.slow class TestLarge: - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="allocates large memory") def test_large(self): gen = random.Generator(random.XORWOW(1234)) 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 1e41e885301..98925845cec 100644 --- a/dpnp/tests/third_party/cupy/sorting_tests/test_search.py +++ b/dpnp/tests/third_party/cupy/sorting_tests/test_search.py @@ -1,7 +1,12 @@ +from __future__ import annotations + import numpy import pytest import dpnp as cupy + +# import cupy._core._accelerator as _acc +# from cupy._core import _cub_reduction from dpnp.tests.helper import has_support_aspect64 from dpnp.tests.third_party.cupy import testing @@ -83,10 +88,9 @@ def test_argmax_zero_size_axis1(self, xp, dtype): return a.argmax(axis=1) @testing.slow - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="allocation too large.") def test_argmax_int32_overflow(self): - a = testing.shaped_arange((2**32 + 1,), cupy, numpy.float64) + a = cupy.arange(2**32 + 1, dtype=cupy.float64) assert a.argmax().item() == 2**32 @testing.for_all_dtypes(no_complex=True) @@ -164,10 +168,9 @@ def test_argmin_zero_size_axis1(self, xp, dtype): return a.argmin(axis=1) @testing.slow - # thread_unsafe marker requires pytest-run-parallel, not used by dpnp # @pytest.mark.thread_unsafe(reason="allocation too large.") def test_argmin_int32_overflow(self): - a = testing.shaped_arange((2**32 + 1,), cupy, numpy.float64) + a = cupy.arange(2**32 + 1, dtype=cupy.float64) cupy.negative(a, out=a) assert a.argmin().item() == 2**32 @@ -184,13 +187,17 @@ def _skip_cuda90(dtype): @testing.parameterize( *testing.product( { - "shape": [(10,), (10, 20), (10, 20, 30), (10, 20, 30, 40)], + # Keep the contiguous reduction axis (last for C, first for F) >= 128 so + # the CUB block-reduction path is used rather than the short-axis fallback. + "shape": [(128,), (128, 128), (128, 2, 128), (128, 2, 2, 128)], "order_and_axis": (("C", -1), ("C", None), ("F", 0), ("F", None)), "backend": ("device", "block"), } ) ) -# thread_unsafe marker requires pytest-run-parallel, not used by dpnp +# @pytest.mark.skipif( +# not cupy.cuda.cub.available, reason="The CUB routine is not enabled" +# ) # @pytest.mark.thread_unsafe(reason="unsafe setUp and counts function calls.") @pytest.mark.skip("The CUB routine is not enabled") class TestCubReduction: @@ -716,7 +723,9 @@ def test_searchsorted(self, xp, dtype): x = testing.shaped_arange(self.shape, xp, dtype) bins = xp.array(self.bins) y = xp.searchsorted(bins, x, side=self.side) - return (y,) + # python scalar for `v` + y1 = xp.searchsorted(bins, 42.5, side=self.side) + return y, y1 @testing.for_all_dtypes(no_bool=True) @testing.numpy_cupy_array_equal() @@ -727,6 +736,15 @@ def test_ndarray_searchsorted(self, xp, dtype): return (y,) +class TestSearchSortedScalar: + @pytest.mark.parametrize("side", ["left", "right"]) + @pytest.mark.parametrize("v", [-1.0, 0.0, 1.0, 1.5, 2.0, 4.0, 4.5]) + @testing.numpy_cupy_array_equal() + def test_searchsorted_scalar(self, xp, v, side): + y = xp.searchsorted(xp.arange(3), v) + return (y,) + + @testing.parameterize({"side": "left"}, {"side": "right"}) class TestSearchSortedNanInf: