From ae1bc56664a607ebae19b071a569657bbbfc7554 Mon Sep 17 00:00:00 2001 From: Max Date: Sat, 10 Oct 2026 20:36:24 +0200 Subject: [PATCH 1/5] AMD support via hipify --- .gitlab-ci.yml | 37 +++++++++++++ docs/source/kernels/cuda-headers.md | 15 +++++ src/cunumpy/_cuda_kernel.py | 42 ++++++++++++-- src/cunumpy/_device.py | 32 +++++++++++ src/cunumpy/cuda/__init__.py | 2 + src/cunumpy/cuda/include/cunumpy/reduce.cuh | 6 ++ src/cunumpy/cuda/include/cunumpy/scan.cuh | 1 + src/cunumpy/kernel_testing.py | 61 +++++++++++++++++++-- tests/unit/test_cuda_collectives.py | 3 +- 9 files changed, 187 insertions(+), 12 deletions(-) diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index 1e91db3..0f98d3b 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -1,10 +1,15 @@ variables: CUDA_IMAGE: "gitlab-registry.mpcdf.mpg.de/mpcdf/ci-module-image/nvhpcsdk_26:2026" + ROCM_IMAGE: "gitlab-registry.mpcdf.mpg.de/mpcdf/ci-module-image/gcc_15-openmpi_5_0-rocm_7_1:2026" PIP_DISABLE_PIP_VERSION_CHECK: "1" CUNUMPY_BACKEND: "cupy" CUNUMPY_REQUIRE_CUDA: "1" # empty: the newest CuPy; gpu_tests_min_cupy pins the oldest supported one CUPY_PIN: "" + # MI210 (MPCDF-GPU-05) is CDNA2, offload architecture gfx90a; also covers + # MI200/MI250(X). Change this (and the gpu-amd-gfx* tag below) for another + # AMD GPU generation. + HCC_AMDGPU_TARGET: "gfx90a" stages: - test @@ -26,12 +31,44 @@ stages: - export PATH="$HOME/.local/bin:$PATH" - python3 -c "import cupy; cupy.show_config()" +# CuPy has no prebuilt ROCm wheel for recent ROCm (the old cupy-rocm-* +# packages stop at ROCm 4.3), so this builds it from source against the +# module-provided ROCm. HCC_AMDGPU_TARGET above must match the runner's GPU. +.gpu_job_amd: + stage: test + image: ${ROCM_IMAGE} + tags: + - gpu-amd + - gpu-amd-gfx90a + before_script: + - module load gcc/15 + - module load openmpi/5.0 + - module load rocm/7.1 + - module load python-waterboa/2025.06 + - rocm-smi + - export ROCM_HOME="${ROCM_PATH:-/opt/rocm}" + - export CUPY_INSTALL_USE_HIP=1 + - python3 -m pip install --user "cupy${CUPY_PIN}" + - python3 -m pip install --user -e ".[test]" + - export PATH="$HOME/.local/bin:$PATH" + - python3 -c "import cupy; cupy.show_config()" + - python3 -c "import cupy; assert cupy.cuda.runtime.is_hip, 'expected a HIP build of CuPy'" + gpu_tests: extends: .gpu_job script: - python3 -m pytest -q tests - python3 testing/cupyx_names.py +# cunumpy/reduce.cuh and cunumpy/scan.cuh are not yet ported to HIP: tests that +# need them use `requires_warp_shuffle`, which skips (not fails) here even +# though CUNUMPY_REQUIRE_CUDA is set, see cunumpy.kernel_testing. +gpu_tests_amd: + extends: .gpu_job_amd + script: + - python3 -m pytest -q tests + - python3 testing/cupyx_names.py + # The oldest supported CuPy (cunumpy.MIN_CUPY_VERSION): the whole suite, and the # fake cupyx.scipy must list exactly its names (_fake_cupyx.NAMES). gpu_tests_min_cupy: diff --git a/docs/source/kernels/cuda-headers.md b/docs/source/kernels/cuda-headers.md index 1d2eea3..66e39b5 100644 --- a/docs/source/kernels/cuda-headers.md +++ b/docs/source/kernels/cuda-headers.md @@ -19,6 +19,13 @@ in [Writing CUDA kernels](cuda-kernel.md)). | `cunumpy/random.cuh` | counter-based random numbers, equal to `xp.rng.philox_*` | | `cunumpy/morton.cuh` | Morton (Z-order) keys, equal to `xp.algorithms.morton_*` | +**AMD/ROCm.** CuPy's own API is the same on a ROCm build, so `xp.set_backend("cupy")` +and the device/memory/stream helpers of `xp.cuda` work unchanged; check +`xp.cuda.is_hip()` to tell the two apart where it matters. `CudaKernel` +compiles plain CUDA C through `cupy.RawKernel`, which HIPRTC accepts for +most kernels; the one shipped header that is not yet portable is +`cunumpy/reduce.cuh` (and `cunumpy/scan.cuh`, which includes it), see below. + ## `cunumpy/index.cuh` ```c @@ -113,6 +120,14 @@ gets the result. default requires all 32 lanes. These rules follow [NVIDIA's warp intrinsic constraints](https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html). +**Not ported to HIP/ROCm.** This header hardcodes a 32-lane warp and CUDA's +`_sync` shuffle intrinsics; AMD wavefronts are commonly 64 lanes wide (CDNA: +MI100/MI200/MI300), and HIP's shuffles have no mask argument. Rather than +compile silently-wrong reductions, `CudaKernel.compile()` raises +`NotImplementedError` for a kernel that includes `cunumpy/reduce.cuh` (or +`cunumpy/scan.cuh`, which includes it) when the active CuPy build targets +HIP (`xp.cuda.is_hip()` is True). + ```c extern "C" __global__ void kinetic_energy(const double* v, long long n, double mass, double* energy) { diff --git a/src/cunumpy/_cuda_kernel.py b/src/cunumpy/_cuda_kernel.py index 2b95972..9394585 100644 --- a/src/cunumpy/_cuda_kernel.py +++ b/src/cunumpy/_cuda_kernel.py @@ -2183,13 +2183,15 @@ def included_headers(self) -> tuple[Path, ...]: ) def compile_options(self) -> tuple[str, ...]: - """Return the NVRTC options a compilation now would use. + """Return the compiler options a compilation now would use. These are :attr:`options`, ``-I`` for cunumpy's header directory, - :data:`~cunumpy.cuda.DEBUG_OPTIONS` if :meth:`debug_active`, and, if - the source includes headers, ``-DCUNUMPY_INCLUDE_HASH=0x`` of - their contents (:func:`~cunumpy.cuda.include_hash`). CuPy keys its - kernel cache on the options, so a changed header means a recompile. + :data:`~cunumpy.cuda.DEBUG_OPTIONS` if :meth:`debug_active` (minus + ``-lineinfo`` on a HIP/ROCm build, an NVRTC-only flag HIPRTC does not + accept), and, if the source includes headers, + ``-DCUNUMPY_INCLUDE_HASH=0x`` of their contents + (:func:`~cunumpy.cuda.include_hash`). CuPy keys its kernel cache on + the options, so a changed header means a recompile. Returns ------- @@ -2202,7 +2204,12 @@ def compile_options(self) -> tuple[str, ...]: if cunumpy_include not in options: options += (cunumpy_include,) if self.debug_active(): - options += tuple(o for o in DEBUG_OPTIONS if o not in options) + from cunumpy._device import is_hip + + debug_options = DEBUG_OPTIONS + if is_hip(): + debug_options = tuple(o for o in debug_options if o != "-lineinfo") + options += tuple(o for o in debug_options if o not in options) headers = self.included_headers if headers: options += (f"-DCUNUMPY_INCLUDE_HASH=0x{include_hash(headers)}",) @@ -2287,6 +2294,13 @@ def compile(self, *, log_stream: Any = None) -> Any: ------ RuntimeError If CuPy or a GPU is not available. + NotImplementedError + If the source includes ``cunumpy/reduce.cuh`` or + ``cunumpy/scan.cuh`` and the active CuPy build targets HIP/ROCm: + their warp-level shuffles assume a 32-lane warp and CUDA's + ``_sync`` intrinsics, neither of which hold on AMD wavefronts + (commonly 64 lanes on CDNA GPUs), so compiling them as-is would + silently reduce or scan the wrong set of threads. """ from cunumpy.xp import cupy_available @@ -2297,8 +2311,24 @@ def compile(self, *, log_stream: Any = None) -> Any: ) import cupy as cp + from cunumpy._device import is_hip + device = _current_device() if device not in self._compiled: + if is_hip(): + unported = {"reduce.cuh", "scan.cuh"} & { + h.name for h in self.included_headers + } + if unported: + names = ", ".join(sorted(f"cunumpy/{n}" for n in unported)) + raise NotImplementedError( + f"CUDA kernel {self.expression!r} includes {names}, which are " + "not yet ported to HIP/ROCm: their warp-level shuffles assume " + "a 32-lane warp and CUDA's `_sync` intrinsics, neither of " + "which hold on AMD wavefronts (commonly 64 lanes on CDNA " + "GPUs). Compiling them as-is would silently reduce/scan the " + "wrong set of threads rather than fail loudly.", + ) options = self.compile_options() if self._template_args is None: raw = cp.RawKernel( diff --git a/src/cunumpy/_device.py b/src/cunumpy/_device.py index 77c72cb..915f546 100644 --- a/src/cunumpy/_device.py +++ b/src/cunumpy/_device.py @@ -64,6 +64,38 @@ def device_count() -> int: return 0 +def is_hip() -> bool: + """Return whether the active CuPy build targets AMD ROCm/HIP. + + CuPy's own API (``cp.cuda.Device``, ``memory_info``, streams, ...) is the + same on a ROCm build; this only distinguishes the GPU vendor for the + parts that differ, such as :meth:`~cunumpy.kernels.CudaKernel.compile_options` + (NVRTC-only flags) and the warp-level primitives of + ``cunumpy/reduce.cuh`` and ``cunumpy/scan.cuh`` (not yet ported to HIP, + see :doc:`/kernels/cuda-headers`). + + Returns + ------- + bool + True if CuPy is available and reports a HIP runtime, False for a + CUDA build or when CuPy/a GPU is unavailable. + + Examples + -------- + >>> xp.cuda.is_hip() # on a CUDA build, or without CuPy + False + """ + if not cupy_available(): + return False + + import cupy as cp + + try: + return bool(cp.cuda.runtime.is_hip) + except Exception: # noqa: BLE001 - tolerate any driver/runtime failure + return False + + def set_device_for_rank(rank: int, devices_per_node: int | None = None) -> int: """Select device ``rank % devices_per_node`` for an MPI rank. diff --git a/src/cunumpy/cuda/__init__.py b/src/cunumpy/cuda/__init__.py index 2ec9d24..8281dc7 100644 --- a/src/cunumpy/cuda/__init__.py +++ b/src/cunumpy/cuda/__init__.py @@ -35,6 +35,7 @@ device_count, free_memory, get_cuda_debug, + is_hip, max_shared_memory_per_block, memory_info, pin_memory, @@ -69,6 +70,7 @@ "free_memory", "get_cuda_debug", "include_hash", + "is_hip", "max_shared_memory_per_block", "memory_info", "parse_cuda_signature", diff --git a/src/cunumpy/cuda/include/cunumpy/reduce.cuh b/src/cunumpy/cuda/include/cunumpy/reduce.cuh index ddcc422..76ec3d6 100644 --- a/src/cunumpy/cuda/include/cunumpy/reduce.cuh +++ b/src/cunumpy/cuda/include/cunumpy/reduce.cuh @@ -27,6 +27,12 @@ // // The warp size is 32, as on all NVIDIA GPUs. // +// Not yet ported to HIP/ROCm: AMD wavefronts are commonly 64 lanes (CDNA: +// MI100/MI200/MI300) rather than 32, and HIP's shuffle intrinsics have no +// `_sync`/mask variant. `CudaKernel.compile()` raises `NotImplementedError` +// for a kernel that includes this header on a HIP build, instead of silently +// reducing/scanning the wrong set of threads. +// // The header directory is added to every CudaKernel's NVRTC options. #ifndef CUNUMPY_REDUCE_CUH diff --git a/src/cunumpy/cuda/include/cunumpy/scan.cuh b/src/cunumpy/cuda/include/cunumpy/scan.cuh index 8ad852c..ea5fb35 100644 --- a/src/cunumpy/cuda/include/cunumpy/scan.cuh +++ b/src/cunumpy/cuda/include/cunumpy/scan.cuh @@ -4,6 +4,7 @@ // every thread must call a block function, including padding threads (value 0). // Supports CUDA shuffle arithmetic types, arbitrary masks and partial warps. // Block scans use 32 shared values per type; calls must be collective. +// Not yet ported to HIP/ROCm, see cunumpy/reduce.cuh. #ifndef CUNUMPY_SCAN_CUH #define CUNUMPY_SCAN_CUH diff --git a/src/cunumpy/kernel_testing.py b/src/cunumpy/kernel_testing.py index 1ba0706..24b0a96 100644 --- a/src/cunumpy/kernel_testing.py +++ b/src/cunumpy/kernel_testing.py @@ -19,9 +19,14 @@ ``requires_device_backend`` A ``skipif`` marker for tests that need a GPU or the fake CuPy, see :func:`device_backend_available`. - - With ``CUNUMPY_REQUIRE_CUDA=1`` (:func:`cuda_required`) they fail - instead of skipping. + ``requires_warp_shuffle`` + A ``skipif`` marker for tests that compile ``cunumpy/reduce.cuh`` or + ``cunumpy/scan.cuh``, see :func:`warp_shuffle_available`. + + With ``CUNUMPY_REQUIRE_CUDA=1`` (:func:`cuda_required`) ``requires_cupy`` + and ``requires_device_backend`` fail instead of skipping; + ``requires_warp_shuffle`` always skips on HIP, a known gap rather than a + broken GPU/driver. * Test ``__device__`` helpers: :func:`device_function_kernel`. * Run CUDA kernels on the CPU, without a GPU (from ``cunumpy._emulation``): :func:`emulate_cuda_kernel`, :func:`emulated_launches`, @@ -100,7 +105,9 @@ def test_parity(kernel): "parity_cases", "requires_cupy", # noqa: F822 "requires_device_backend", # noqa: F822 + "requires_warp_shuffle", # noqa: F822 "run_in_fake_cupy_subprocess", + "warp_shuffle_available", ] SKIP_REASON = "CuPy/GPU not available" @@ -108,6 +115,9 @@ def test_parity(kernel): "neither a GPU nor the fake CuPy (CUNUMPY_FAKE_CUPY=1) is available" ) FAKE_SKIP_REASON = "the fake CuPy cannot run CUDA kernels" +WARP_HIP_SKIP_REASON = ( + "cunumpy/reduce.cuh and cunumpy/scan.cuh are not yet ported to HIP/ROCm" +) def fake_cupy_active() -> bool: @@ -203,6 +213,36 @@ def cuda_gate() -> bool: return False +def warp_shuffle_available() -> bool: + """Whether CUDA kernels using ``cunumpy/reduce.cuh`` or ``cunumpy/scan.cuh`` can run here. + + False on CuPy's HIP/ROCm build, even with a working GPU: their warp + shuffles assume a 32-lane warp and CUDA's ``_sync`` mask intrinsics, + which are not yet ported to HIP (see + :meth:`~cunumpy.kernels.CudaKernel.compile`). The marker + ``requires_warp_shuffle`` skips tests where this is false; unlike + ``requires_cupy``, it always skips on HIP (even under + ``CUNUMPY_REQUIRE_CUDA``), since that is an intentional, known gap and + not a broken GPU/driver. + + Returns + ------- + bool + True if CUDA kernels can be launched and the backend is not HIP. + + Examples + -------- + >>> from cunumpy.kernel_testing import warp_shuffle_available + >>> warp_shuffle_available() # doctest: +SKIP + True + """ + if not _can_launch(): + return False + from cunumpy._device import is_hip + + return not is_hip() + + def device_backend_available() -> bool: """Tell whether a CuPy-backend program can run here: a GPU, or the fake CuPy. @@ -436,6 +476,12 @@ def _build_lazy() -> None: requires_device_backend = pytest.mark.skipif( not device_backend_available(), reason=DEVICE_SKIP_REASON ) + # Always skips on HIP (never fails under CUNUMPY_REQUIRE_CUDA): the gap is a + # known, intentional one, not a broken GPU/driver that CI should catch. + requires_warp_shuffle = pytest.mark.skipif( + not warp_shuffle_available(), + reason=SKIP_REASON if not _can_launch() else WARP_HIP_SKIP_REASON, + ) backends = ["numpy", pytest.param("cupy", marks=requires_cupy)] @pytest.fixture(params=backends) @@ -461,13 +507,20 @@ def backend(request): _LAZY.update( requires_cupy=requires_cupy, requires_device_backend=requires_device_backend, + requires_warp_shuffle=requires_warp_shuffle, BACKENDS=backends, backend=backend, ) def __getattr__(name: str) -> Any: - if name in ("requires_cupy", "requires_device_backend", "BACKENDS", "backend"): + if name in ( + "requires_cupy", + "requires_device_backend", + "requires_warp_shuffle", + "BACKENDS", + "backend", + ): if not _LAZY: _build_lazy() return _LAZY[name] diff --git a/tests/unit/test_cuda_collectives.py b/tests/unit/test_cuda_collectives.py index 899cefb..8a251cf 100644 --- a/tests/unit/test_cuda_collectives.py +++ b/tests/unit/test_cuda_collectives.py @@ -7,10 +7,9 @@ import cunumpy as xp from cunumpy.kernel_testing import emulate_cuda_kernel, emulation_compiler +from cunumpy.kernel_testing import requires_warp_shuffle as requires_cuda from cunumpy.kernels import CudaKernel -requires_cuda = pytest.mark.skipif(not xp.cupy_available(), reason="requires CUDA") - TYPED_SOURCE = r""" #include template From 06b64be91b04b067b328dd72fe49505b5d269f75 Mon Sep 17 00:00:00 2001 From: Max Date: Sat, 10 Oct 2026 20:46:37 +0200 Subject: [PATCH 2/5] Updated job names --- .gitlab-ci.yml | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index 0f98d3b..1d852e4 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -14,7 +14,7 @@ variables: stages: - test -.gpu_job: +.gpu_job_nvidia: stage: test image: ${CUDA_IMAGE} tags: @@ -54,8 +54,8 @@ stages: - python3 -c "import cupy; cupy.show_config()" - python3 -c "import cupy; assert cupy.cuda.runtime.is_hip, 'expected a HIP build of CuPy'" -gpu_tests: - extends: .gpu_job +gpu_tests_nvidia: + extends: .gpu_job_nvidia script: - python3 -m pytest -q tests - python3 testing/cupyx_names.py @@ -71,8 +71,8 @@ gpu_tests_amd: # The oldest supported CuPy (cunumpy.MIN_CUPY_VERSION): the whole suite, and the # fake cupyx.scipy must list exactly its names (_fake_cupyx.NAMES). -gpu_tests_min_cupy: - extends: .gpu_job +gpu_tests_min_cupy_nvidia: + extends: .gpu_job_nvidia variables: CUPY_PIN: "==14.0.0" script: @@ -80,8 +80,8 @@ gpu_tests_min_cupy: - python3 -m pytest -q tests - python3 testing/cupyx_names.py --exact -gpu_sanitizers: - extends: .gpu_job +gpu_sanitizers_nvidia: + extends: .gpu_job_nvidia timeout: 45 minutes parallel: matrix: From 07f719a9534b9a9f92612c83df7d3c97cfd3abad Mon Sep 17 00:00:00 2001 From: Max Date: Sat, 10 Oct 2026 21:17:59 +0200 Subject: [PATCH 3/5] Remove _trap() on HIP --- .../cuda/include/cunumpy/array_view.cuh | 12 +++++- tests/unit/test_cuda_kernel.py | 37 ++++++++++++++----- 2 files changed, 39 insertions(+), 10 deletions(-) diff --git a/src/cunumpy/cuda/include/cunumpy/array_view.cuh b/src/cunumpy/cuda/include/cunumpy/array_view.cuh index 2a299ae..17bc2f1 100644 --- a/src/cunumpy/cuda/include/cunumpy/array_view.cuh +++ b/src/cunumpy/cuda/include/cunumpy/array_view.cuh @@ -42,13 +42,23 @@ #define CUNUMPY_ARRAY_VIEW_CUH #ifdef CUNUMPY_BOUNDS_CHECK +// __trap() is a CUDA device builtin, found unchanged by NVRTC and (deliberately +// shadowable) by the C++ compiler used for CPU emulation (cunumpy._emulation +// defines its own host __trap()). HIPRTC has no such builtin; only a real HIP +// device compile (never the CPU emulation path, which never defines these +// macros) uses __builtin_trap() instead. +#if defined(__HIP_DEVICE_COMPILE__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#define CUNUMPY_TRAP() __builtin_trap() +#else +#define CUNUMPY_TRAP() __trap() +#endif #define CUNUMPY_CHECK_INDEX(index, axis, extent) \ do { \ if ((index) < 0 || (index) >= (extent)) { \ printf("cunumpy: index %lld is out of bounds for axis %d with size " \ "%lld\n", \ (long long)(index), (int)(axis), (long long)(extent)); \ - __trap(); \ + CUNUMPY_TRAP(); \ } \ } while (0) #else diff --git a/tests/unit/test_cuda_kernel.py b/tests/unit/test_cuda_kernel.py index 84e3850..39d74a5 100644 --- a/tests/unit/test_cuda_kernel.py +++ b/tests/unit/test_cuda_kernel.py @@ -64,6 +64,14 @@ def _user_options(kernel): ) +def _debug_options(): + """``DEBUG_OPTIONS`` as actually added by ``compile_options()``: no ``-lineinfo`` + on a HIP/ROCm build, an NVRTC-only flag HIPRTC does not accept.""" + if xp.cuda.is_hip(): + return tuple(o for o in xp.cuda.DEBUG_OPTIONS if o != "-lineinfo") + return xp.cuda.DEBUG_OPTIONS + + class FakeDeviceArray: """Enough of a CuPy array for the argument checks: dtype, interface, address. @@ -101,6 +109,14 @@ def _skip_without_cupy(): pytest.skip("CuPy not installed or not functional") +def _skip_without_warp_shuffle(): + """Like `_skip_without_cupy`, also skipping on HIP: reduce.cuh/scan.cuh are + not yet ported (see `CudaKernel.compile`).""" + _skip_without_cupy() + if xp.cuda.is_hip(): + pytest.skip("cunumpy/reduce.cuh is not yet ported to HIP/ROCm") + + # --------------------------------------------------------------------------- # signature parsing # --------------------------------------------------------------------------- @@ -1478,11 +1494,7 @@ def test_compile_options_follow_the_global_setting(debug_off): with xp.cuda.cuda_debug(): # decided at call time: the kernel created before is affected assert kernel.debug_active() is True - assert _user_options(kernel) == ( - "-std=c++17", - "-lineinfo", - "-DCUNUMPY_BOUNDS_CHECK", - ) + assert _user_options(kernel) == ("-std=c++17", *_debug_options()) assert _user_options(kernel) == ("-std=c++17",) assert kernel.options == ("-std=c++17",) # the given options are unchanged @@ -1500,18 +1512,18 @@ def test_explicit_debug_overrides_the_global_setting(debug_off): assert on.debug is True and off.debug is False assert on.debug_active() is True assert off.debug_active() is False - assert set(xp.cuda.DEBUG_OPTIONS) <= set(on.compile_options()) + assert set(_debug_options()) <= set(on.compile_options()) with xp.cuda.cuda_debug(): assert off.debug_active() is False assert _user_options(off) == () - assert _user_options(on) == xp.cuda.DEBUG_OPTIONS + assert _user_options(on) == _debug_options() def test_debug_options_include_dirs_and_from_file(tmp_path, debug_off): (tmp_path / "axpy_cuda.cu").write_text(AXPY) kernel = CudaKernel.from_file(tmp_path / "axpy_cuda.cu", debug=True) - assert _user_options(kernel) == (f"-I{tmp_path}", *xp.cuda.DEBUG_OPTIONS) + assert _user_options(kernel) == (f"-I{tmp_path}", *_debug_options()) def test_debug_option_is_not_G(): @@ -1565,6 +1577,13 @@ def test_out_of_bounds_write_on_gpu(debug): """Under debug, the error is a RuntimeError naming the kernel and its shape; without debug, it is CuPy's own error (raised at the synchronization).""" _skip_without_cupy() + if xp.cuda.is_hip(): + pytest.skip( + "on HIP/ROCm an illegal device write is a fatal GPU memory-protection " + "fault that terminates the process (the driver prints it directly), " + "unlike NVIDIA's sticky CUDA context error that surfaces as a " + "catchable exception at the next synchronization", + ) output = "".join( _run_python( OUT_OF_BOUNDS.replace("DEBUG", str(debug)), @@ -2214,7 +2233,7 @@ def test_reduce_header_is_shipped(): @pytest.mark.parametrize("block_size", [32, 128, 1024]) def test_reductions_on_gpu(block_size): - _skip_without_cupy() + _skip_without_warp_shuffle() import cupy as cp n = 5000 # not a multiple of the block size: the last block is partial From 51dcfa966f1e8a6aeb875c9c69d8325ffbfc8648 Mon Sep 17 00:00:00 2001 From: Max Date: Sat, 10 Oct 2026 22:10:52 +0200 Subject: [PATCH 4/5] Fix tests --- docs/source/kernels/cuda-headers.md | 21 +++++- src/cunumpy/_cuda_kernel.py | 58 +++++++++++++++- src/cunumpy/_device.py | 16 ++++- src/cunumpy/kernel_testing.py | 102 +++++++++++++++++++++++++++- tests/unit/test_cuda_kernel.py | 6 ++ tests/unit/test_cunumpy.py | 15 ++-- tests/unit/test_emulation.py | 6 ++ 7 files changed, 211 insertions(+), 13 deletions(-) diff --git a/docs/source/kernels/cuda-headers.md b/docs/source/kernels/cuda-headers.md index 66e39b5..4eff27f 100644 --- a/docs/source/kernels/cuda-headers.md +++ b/docs/source/kernels/cuda-headers.md @@ -23,8 +23,25 @@ in [Writing CUDA kernels](cuda-kernel.md)). and the device/memory/stream helpers of `xp.cuda` work unchanged; check `xp.cuda.is_hip()` to tell the two apart where it matters. `CudaKernel` compiles plain CUDA C through `cupy.RawKernel`, which HIPRTC accepts for -most kernels; the one shipped header that is not yet portable is -`cunumpy/reduce.cuh` (and `cunumpy/scan.cuh`, which includes it), see below. +most kernels. Three known gaps, found on real HIP/ROCm hardware (an AMD +Instinct MI210) and all raising `NotImplementedError` at +`CudaKernel.compile()` on HIP rather than silently misbehaving, none yet +root-caused: + +* `cunumpy/reduce.cuh` (and `cunumpy/scan.cuh`, which includes it), see below. +* A kernel taking an `Array5D`/`CArray5D` view or higher (`Array1D` to + `Array4D` are unaffected): launching reliably corrupts the device, even for + an in-bounds access (a HIP code-generation or by-value struct-argument + issue in the generic, variadic-template `ArrayView` is suspected). +* A kernel that includes `cunumpy/array_view.cuh` and compiles with + `CUNUMPY_BOUNDS_CHECK` (directly, or via debug mode, which adds it): merely + compiling the resulting `printf()` and trap reliably corrupts the device on + any launch, even one that never takes an out-of-bounds index (a HIP + device-`printf` or trap-instruction issue is suspected). + +Use `cunumpy.kernel_testing`'s `requires_warp_shuffle`, `requires_high_dim_cuda_views` +and `requires_bounds_check_views` markers (or their `*_available()` functions) +to skip tests that hit these on HIP, the way cunumpy's own test suite does. ## `cunumpy/index.cuh` diff --git a/src/cunumpy/_cuda_kernel.py b/src/cunumpy/_cuda_kernel.py index 9394585..0451fb5 100644 --- a/src/cunumpy/_cuda_kernel.py +++ b/src/cunumpy/_cuda_kernel.py @@ -857,6 +857,17 @@ def check(value: Any) -> Any: return check +def _max_view_ndim(params: Iterable[CudaParameter]) -> int: + """Return the largest `view_ndim` among `params` and their struct fields (0 if none).""" + best = 0 + for p in params: + if p.view_ndim is not None: + best = max(best, p.view_ndim) + if p.struct is not None: + best = max(best, _max_view_ndim(p.struct.fields)) + return best + + def _checker(param: CudaParameter, index: int) -> Callable[[Any], Any]: if param.struct is not None: return _struct_checker(param, index) @@ -2300,7 +2311,21 @@ def compile(self, *, log_stream: Any = None) -> Any: their warp-level shuffles assume a 32-lane warp and CUDA's ``_sync`` intrinsics, neither of which hold on AMD wavefronts (commonly 64 lanes on CDNA GPUs), so compiling them as-is would - silently reduce or scan the wrong set of threads. + silently reduce or scan the wrong set of threads. Also raised, on + a HIP/ROCm build, for a kernel using an ``Array5D``/``CArray5D`` + view or higher: launching one reliably corrupts the device there + (observed as ``hipErrorIllegalState`` on a later launch), even for + an in-bounds access; ``Array1D`` to ``Array4D`` are unaffected. + Not yet root-caused (looks like a HIP code-generation or + by-value-struct-argument issue in the generic, variadic-template + ``ArrayView`` that implements dimensions 5 and up, see + ``cunumpy/array_view.cuh``). Also raised, on a HIP/ROCm build, + for a kernel that includes ``cunumpy/array_view.cuh`` and + compiles with ``-DCUNUMPY_BOUNDS_CHECK`` (directly or via debug + mode): merely compiling the resulting ``printf()`` and trap + reliably corrupts the device on any launch of that kernel, even + one that never takes an out-of-bounds index; also not yet + root-caused. """ from cunumpy.xp import cupy_available @@ -2329,7 +2354,38 @@ def compile(self, *, log_stream: Any = None) -> Any: "GPUs). Compiling them as-is would silently reduce/scan the " "wrong set of threads rather than fail loudly.", ) + if self._signature is not None: + max_view = _max_view_ndim(self._signature) + if max_view >= 5: + raise NotImplementedError( + f"CUDA kernel {self.expression!r} uses an " + f"Array{max_view}D/CArray{max_view}D view: on this " + "HIP/ROCm build, launching a kernel with a 5+ " + "dimensional array view reliably corrupts the device " + "(hipErrorIllegalState on a later launch), even for an " + "in-bounds access. Array1D to Array4D are unaffected; " + "this is not yet root-caused (likely a HIP code " + "generation or by-value struct argument issue in the " + "generic ArrayView template, not something " + "cunumpy can safely paper over).", + ) options = self.compile_options() + if ( + is_hip() + and "-DCUNUMPY_BOUNDS_CHECK" in options + and "array_view.cuh" in {h.name for h in self.included_headers} + ): + raise NotImplementedError( + f"CUDA kernel {self.expression!r} includes cunumpy/array_view.cuh " + "and compiles with -DCUNUMPY_BOUNDS_CHECK (directly, or via debug " + "mode) on this HIP/ROCm build: merely compiling the resulting " + "printf() + trap in its bounds check reliably corrupts the " + "device on any launch of that kernel, even one that never takes " + "an out-of-bounds index. Not yet root-caused (a HIP " + "device-printf or trap-instruction issue is suspected); compile " + "without CUNUMPY_BOUNDS_CHECK (and without debug mode, which adds " + "it) for a kernel using array views on this build.", + ) if self._template_args is None: raw = cp.RawKernel( self._source, diff --git a/src/cunumpy/_device.py b/src/cunumpy/_device.py index 915f546..3a48ddf 100644 --- a/src/cunumpy/_device.py +++ b/src/cunumpy/_device.py @@ -232,9 +232,19 @@ def max_shared_memory_per_block( return DEFAULT_SHARED_MEMORY_PER_BLOCK import cupy as cp - dev = cp.cuda.Device() if device is None else cp.cuda.Device(device) - key = "MaxSharedMemoryPerBlockOptin" if opt_in else "MaxSharedMemoryPerBlock" - return int(dev.attributes.get(key, DEFAULT_SHARED_MEMORY_PER_BLOCK)) + # `cp.cuda.Device().attributes` is unreliable here on a HIP/ROCm build (it + # can report 0 for these keys); `getDeviceProperties()`'s `sharedMemPerBlock` + # and `sharedMemPerBlockOptin` fields are the ones `CudaKernel` itself + # already relies on for the real launch-time check, on both backends. + dev_id = cp.cuda.Device().id if device is None else device + properties = cp.cuda.runtime.getDeviceProperties(dev_id) + base = int(properties.get("sharedMemPerBlock", 0)) or DEFAULT_SHARED_MEMORY_PER_BLOCK + if not opt_in: + return base + # The "opt in" larger dynamic-shared-memory limit is an NVIDIA concept + # (cudaFuncAttributeMaxDynamicSharedMemorySize); on a HIP build there is no + # larger limit to opt into, so the base limit is also the opt-in one. + return max(base, int(properties.get("sharedMemPerBlockOptin", 0))) def free_memory() -> None: diff --git a/src/cunumpy/kernel_testing.py b/src/cunumpy/kernel_testing.py index 24b0a96..b0c5192 100644 --- a/src/cunumpy/kernel_testing.py +++ b/src/cunumpy/kernel_testing.py @@ -22,11 +22,20 @@ ``requires_warp_shuffle`` A ``skipif`` marker for tests that compile ``cunumpy/reduce.cuh`` or ``cunumpy/scan.cuh``, see :func:`warp_shuffle_available`. + ``requires_high_dim_cuda_views`` + A ``skipif`` marker for tests that launch a kernel using an + ``Array5D``/``CArray5D`` view or higher, see + :func:`high_dim_cuda_views_available`. + ``requires_bounds_check_views`` + A ``skipif`` marker for tests that compile a kernel including + ``cunumpy/array_view.cuh`` with ``CUNUMPY_BOUNDS_CHECK`` (directly or + via debug mode), see :func:`bounds_check_views_available`. With ``CUNUMPY_REQUIRE_CUDA=1`` (:func:`cuda_required`) ``requires_cupy`` and ``requires_device_backend`` fail instead of skipping; - ``requires_warp_shuffle`` always skips on HIP, a known gap rather than a - broken GPU/driver. + ``requires_warp_shuffle``, ``requires_high_dim_cuda_views`` and + ``requires_bounds_check_views`` always skip on HIP, known gaps rather than + a broken GPU/driver. * Test ``__device__`` helpers: :func:`device_function_kernel`. * Run CUDA kernels on the CPU, without a GPU (from ``cunumpy._emulation``): :func:`emulate_cuda_kernel`, :func:`emulated_launches`, @@ -89,6 +98,7 @@ def test_parity(kernel): "BACKENDS", # noqa: F822 "assert_kernels_agree", "backend", # noqa: F822 + "bounds_check_views_available", "check_parity", "compile_for_emulation", "cuda_required", @@ -100,11 +110,14 @@ def test_parity(kernel): "emulation_compiler", "fake_cupy_active", "fake_cupy_session", + "high_dim_cuda_views_available", "host_buffer", "install_fake_cupy", "parity_cases", + "requires_bounds_check_views", # noqa: F822 "requires_cupy", # noqa: F822 "requires_device_backend", # noqa: F822 + "requires_high_dim_cuda_views", # noqa: F822 "requires_warp_shuffle", # noqa: F822 "run_in_fake_cupy_subprocess", "warp_shuffle_available", @@ -118,6 +131,16 @@ def test_parity(kernel): WARP_HIP_SKIP_REASON = ( "cunumpy/reduce.cuh and cunumpy/scan.cuh are not yet ported to HIP/ROCm" ) +HIGH_DIM_VIEW_HIP_SKIP_REASON = ( + "launching a kernel with an Array5D/CArray5D view or higher reliably " + "corrupts the device on this HIP/ROCm build (not yet root-caused)" +) +BOUNDS_CHECK_VIEW_HIP_SKIP_REASON = ( + "compiling a kernel that includes cunumpy/array_view.cuh with " + "CUNUMPY_BOUNDS_CHECK (directly or via debug mode) reliably corrupts the " + "device on this HIP/ROCm build, even without an out-of-bounds access " + "(not yet root-caused)" +) def fake_cupy_active() -> bool: @@ -243,6 +266,69 @@ def warp_shuffle_available() -> bool: return not is_hip() +def high_dim_cuda_views_available() -> bool: + """Whether a CUDA kernel using an ``Array5D``/``CArray5D`` view or higher can run here. + + False on CuPy's HIP/ROCm build, even with a working GPU: launching such a + kernel there reliably corrupts the device (observed as + ``hipErrorIllegalState`` on a later launch), even for an in-bounds + access, and is not yet root-caused (see + :meth:`~cunumpy.kernels.CudaKernel.compile`). ``Array1D`` to ``Array4D`` + are unaffected. The marker ``requires_high_dim_cuda_views`` skips tests + where this is false; like ``requires_warp_shuffle``, it always skips on + HIP (even under ``CUNUMPY_REQUIRE_CUDA``), since this is a known gap, not + a broken GPU/driver. + + Returns + ------- + bool + True if CUDA kernels can be launched and the backend is not HIP. + + Examples + -------- + >>> from cunumpy.kernel_testing import high_dim_cuda_views_available + >>> high_dim_cuda_views_available() # doctest: +SKIP + True + """ + if not _can_launch(): + return False + from cunumpy._device import is_hip + + return not is_hip() + + +def bounds_check_views_available() -> bool: + """Whether a kernel using ``cunumpy/array_view.cuh`` can compile with ``CUNUMPY_BOUNDS_CHECK`` here. + + False on CuPy's HIP/ROCm build, even with a working GPU: merely + compiling such a kernel with bounds checking on (directly, or via debug + mode, which adds it) reliably corrupts the device on any launch, even + one that never takes an out-of-bounds index, and is not yet root-caused + (see :meth:`~cunumpy.kernels.CudaKernel.compile`). A kernel that does not + include ``cunumpy/array_view.cuh`` is unaffected. The marker + ``requires_bounds_check_views`` skips tests where this is false; like + ``requires_warp_shuffle``, it always skips on HIP (even under + ``CUNUMPY_REQUIRE_CUDA``), since this is a known gap, not a broken + GPU/driver. + + Returns + ------- + bool + True if CUDA kernels can be launched and the backend is not HIP. + + Examples + -------- + >>> from cunumpy.kernel_testing import bounds_check_views_available + >>> bounds_check_views_available() # doctest: +SKIP + True + """ + if not _can_launch(): + return False + from cunumpy._device import is_hip + + return not is_hip() + + def device_backend_available() -> bool: """Tell whether a CuPy-backend program can run here: a GPU, or the fake CuPy. @@ -482,6 +568,14 @@ def _build_lazy() -> None: not warp_shuffle_available(), reason=SKIP_REASON if not _can_launch() else WARP_HIP_SKIP_REASON, ) + requires_high_dim_cuda_views = pytest.mark.skipif( + not high_dim_cuda_views_available(), + reason=SKIP_REASON if not _can_launch() else HIGH_DIM_VIEW_HIP_SKIP_REASON, + ) + requires_bounds_check_views = pytest.mark.skipif( + not bounds_check_views_available(), + reason=SKIP_REASON if not _can_launch() else BOUNDS_CHECK_VIEW_HIP_SKIP_REASON, + ) backends = ["numpy", pytest.param("cupy", marks=requires_cupy)] @pytest.fixture(params=backends) @@ -508,6 +602,8 @@ def backend(request): requires_cupy=requires_cupy, requires_device_backend=requires_device_backend, requires_warp_shuffle=requires_warp_shuffle, + requires_high_dim_cuda_views=requires_high_dim_cuda_views, + requires_bounds_check_views=requires_bounds_check_views, BACKENDS=backends, backend=backend, ) @@ -518,6 +614,8 @@ def __getattr__(name: str) -> Any: "requires_cupy", "requires_device_backend", "requires_warp_shuffle", + "requires_high_dim_cuda_views", + "requires_bounds_check_views", "BACKENDS", "backend", ): diff --git a/tests/unit/test_cuda_kernel.py b/tests/unit/test_cuda_kernel.py index 39d74a5..53bf69d 100644 --- a/tests/unit/test_cuda_kernel.py +++ b/tests/unit/test_cuda_kernel.py @@ -30,6 +30,10 @@ parse_cuda_signature, resolve_includes, ) +from cunumpy.kernel_testing import ( + requires_bounds_check_views, + requires_high_dim_cuda_views, +) from cunumpy.kernels import CudaKernel, CudaKernelVariants AXPY = r""" @@ -987,6 +991,7 @@ def test_struct_with_view_fields_on_gpu(): assert np.array_equal(base.get(), expected) +@requires_bounds_check_views def test_bounds_check_traps_on_gpu(): _skip_without_cupy() import cupy as cp @@ -2486,6 +2491,7 @@ def init(a): ... @pytest.mark.parametrize("ndim", [5, 6, 10, 16]) @pytest.mark.parametrize("contiguous", [False, True]) +@requires_high_dim_cuda_views def test_high_dimensional_struct_layout_and_execution_on_gpu(ndim, contiguous): _skip_without_cupy() import cupy as cp diff --git a/tests/unit/test_cunumpy.py b/tests/unit/test_cunumpy.py index 77e97ce..ce5120e 100644 --- a/tests/unit/test_cunumpy.py +++ b/tests/unit/test_cunumpy.py @@ -411,13 +411,18 @@ def test_max_shared_memory_per_block_reads_the_device(monkeypatch): class Device: def __init__(self, device_id=0): - self.attributes = { - "MaxSharedMemoryPerBlock": 49152 + device_id, - "MaxSharedMemoryPerBlockOptin": 232448, - } + self.id = device_id + + properties = { + 0: {"sharedMemPerBlock": 49152, "sharedMemPerBlockOptin": 232448}, + 1: {"sharedMemPerBlock": 49153, "sharedMemPerBlockOptin": 232449}, + } cupy = types.ModuleType("cupy") - cupy.cuda = types.SimpleNamespace(Device=Device) + cupy.cuda = types.SimpleNamespace( + Device=Device, + runtime=types.SimpleNamespace(getDeviceProperties=properties.__getitem__), + ) monkeypatch.setitem(sys.modules, "cupy", cupy) monkeypatch.setattr(device_module, "cupy_available", lambda: True) assert xp.cuda.max_shared_memory_per_block() == 49152 diff --git a/tests/unit/test_emulation.py b/tests/unit/test_emulation.py index 87fbac3..fe6c10a 100644 --- a/tests/unit/test_emulation.py +++ b/tests/unit/test_emulation.py @@ -435,6 +435,12 @@ def test_high_dimensional_view_indexing_and_conversion(ndim, contiguous, backend if not xp.cupy_available(): pytest.skip("CuPy not installed or not functional") + if xp.cuda.is_hip(): + pytest.skip( + "launching a kernel with an Array5D/CArray5D view or higher " + "reliably corrupts the device on this HIP/ROCm build (not yet " + "root-caused)", + ) import cupy as cp device_base = cp.asarray(base) From dbdebe8753765ddeb23b80d054d7bf268d4d8e30 Mon Sep 17 00:00:00 2001 From: Max Date: Sat, 10 Oct 2026 22:40:13 +0200 Subject: [PATCH 5/5] formatting --- src/cunumpy/_device.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/cunumpy/_device.py b/src/cunumpy/_device.py index 3a48ddf..dd4e489 100644 --- a/src/cunumpy/_device.py +++ b/src/cunumpy/_device.py @@ -238,7 +238,9 @@ def max_shared_memory_per_block( # already relies on for the real launch-time check, on both backends. dev_id = cp.cuda.Device().id if device is None else device properties = cp.cuda.runtime.getDeviceProperties(dev_id) - base = int(properties.get("sharedMemPerBlock", 0)) or DEFAULT_SHARED_MEMORY_PER_BLOCK + base = ( + int(properties.get("sharedMemPerBlock", 0)) or DEFAULT_SHARED_MEMORY_PER_BLOCK + ) if not opt_in: return base # The "opt in" larger dynamic-shared-memory limit is an NVIDIA concept