From 8dc5012c89b244ee8d380680679c1b5f2bdd610b Mon Sep 17 00:00:00 2001 From: Max Date: Wed, 7 Oct 2026 23:42:31 +0200 Subject: [PATCH 1/2] Add emulated launches, sync counting, compact_by_mask and output names --- CHANGELOG.md | 24 +++ docs/source/api.md | 51 +++++-- docs/source/guides/data-movement.md | 24 +++ docs/source/guides/execution-helpers.md | 15 +- docs/source/kernels/pyccel-kernel.md | 14 +- docs/source/kernels/testing.md | 54 ++++++- src/cunumpy/LLM_GUIDE.md | 4 + src/cunumpy/_algorithms.py | 46 ++++++ src/cunumpy/_cuda_kernel.py | 5 + src/cunumpy/_dispatch.py | 38 +++++ src/cunumpy/_emulation.py | 192 ++++++++++++++++++++++-- src/cunumpy/_fake_cupy.py | 28 +++- src/cunumpy/_fake_cupy_impl.py | 15 ++ src/cunumpy/_kernel.py | 152 ++++++++++++++++--- src/cunumpy/_mpi.py | 6 +- src/cunumpy/_transfers.py | 46 +++++- src/cunumpy/algorithms.py | 6 +- src/cunumpy/kernel_testing.py | 152 +++++++++++++++---- src/cunumpy/kernels.py | 2 + src/cunumpy/xp.py | 3 + tests/unit/test_compact_by_mask.py | 43 ++++++ tests/unit/test_emulated_launches.py | 159 ++++++++++++++++++++ tests/unit/test_kernel_outputs.py | 156 +++++++++++++++++++ tests/unit/test_kernel_testing.py | 4 +- tests/unit/test_porting_helpers.py | 5 +- tests/unit/test_pyccel_kernel.py | 2 +- tests/unit/test_require_cuda.py | 66 ++++++++ tests/unit/test_sync_counting.py | 76 ++++++++++ tests/unit/test_transfers.py | 4 +- 29 files changed, 1295 insertions(+), 97 deletions(-) create mode 100644 tests/unit/test_compact_by_mask.py create mode 100644 tests/unit/test_emulated_launches.py create mode 100644 tests/unit/test_kernel_outputs.py create mode 100644 tests/unit/test_require_cuda.py create mode 100644 tests/unit/test_sync_counting.py diff --git a/CHANGELOG.md b/CHANGELOG.md index 778b957..4b30ca5 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -8,6 +8,30 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [Unreleased] ### Added +- `kernel_testing.emulated_launches()` runs every `CudaKernel` launch in a block + on the CPU, on the host buffers of the fake CuPy arrays, so code that launches + kernels can be tested without a GPU. `emulate_cuda_kernel` now accepts struct + parameters (a mapping of field values, a `CudaStructValue`, or an object with + an attribute per field). `kernel_testing.host_buffer(array)` returns the NumPy + array behind a fake CuPy array. +- `CUNUMPY_REQUIRE_CUDA=1` makes the GPU markers of `kernel_testing` fail instead + of skipping: `requires_cupy`, the `cupy` run of the `backend` fixture (which + activates CuPy strictly) and `assert_kernels_agree`. New `kernel_testing.cuda_required()`. +- `count_transfers()` records `sync` events (the host waiting for the device): + `xp.synchronize()`, the waits of the MPI helpers and of the CUDA debug mode, and, + on the fake CuPy, scalar reads of device arrays. They are in `counter.syncs` + and the report but not in `total`; `assert_no_transfers(syncs=True)` rejects them. + The real CuPy's own `float(a)` cannot be observed from Python and is not counted. +- `xp.algorithms.compact_by_mask(mask, *arrays)` moves the masked rows of arrays to + the front, in place and in order, and returns their number. +- Kernel outputs: a name in `PyccelKernel(outputs=...)` also finds a positional + argument and an index a keyword argument, using the parameter names of the + function (or the new `parameters=`); a `Kernel` supplies those of its host + function. `xp.kernels.outputs_from_annotations()` reads the outputs from the + annotations (not `Final`, `const` or a scalar), and `Kernel.from_folder()` / + `KernelCatalog.from_package()` take `outputs=` (names, indices or + `"annotations"`). `assert_kernels_agree(outputs=...)` takes parameter names and + `"name.field"` to compare only some fields of a struct argument. - `xp.kernels.MetalKernel` runs a Metal Shading Language kernel on the GPU of an Apple silicon Mac through MLX (`pip install 'cunumpy[metal]'`). It takes and fills NumPy arrays, is float32 only (`float64="cast"` computes float64 data in diff --git a/docs/source/api.md b/docs/source/api.md index fba1154..195a2a1 100644 --- a/docs/source/api.md +++ b/docs/source/api.md @@ -44,7 +44,7 @@ The submodules are named so that they do not hide a NumPy name (`rng`, not | `cunumpy.arguments` | CUDA only | `CudaArguments`, `CudaStruct`, `CudaStructArguments`, `CudaStructValue`, `write_cuda_header` | | `cunumpy.cuda` | CUDA only | device selection and memory, `stream`, streams/events, `pin_memory`, debug mode, CUDA headers and source tools (`cuda_include_dir`, `parse_cuda_signature`) | | `cunumpy.rng` | both | `random_streams`, `get_rng`, `philox_*` | -| `cunumpy.algorithms` | both | `morton_*`, `sort_by_key`, `segment_sum` | +| `cunumpy.algorithms` | both | `morton_*`, `sort_by_key`, `segment_sum`, `compact_by_mask` | | `cunumpy.mpi` | both | `mpi_buffer`, CUDA-aware MPI detection, `local_rank`, `synchronize_for_mpi` | | `cunumpy.profiling` | both | `timed_region`, `nvtx_range`, `count_transfers`, `assert_no_transfers` | | `cunumpy.memory` | both | `HostStaging`, `DeviceMirror` | @@ -267,6 +267,21 @@ keys, order, positions, charges = xp.algorithms.sort_by_key(keys, positions, cha Returns `(keys[order], order, *(a[order] for a in arrays))`, `order` as `int64`. Equal keys keep their order, so the result is reproducible. +### `algorithms.compact_by_mask(mask, *arrays)` + +Moves the rows where the boolean `mask` is True to the front of every array, in +place and in their original order, and returns how many there are. Typical use: +keep the live particles at the front of the marker arrays. + +```python +n = xp.algorithms.compact_by_mask(alive, markers, weights) +markers, weights = markers[:n], weights[:n] +``` + +The rows after the first `n` are unspecified. The count is needed on the host, +so on CuPy each call synchronizes once. The mask and the arrays must be on the +same backend. + ## Count transfers A transfer inside a time loop is the classic performance bug of a GPU port: @@ -286,7 +301,7 @@ with xp.profiling.count_transfers() as counter: assert counter.total == 0, counter.report() ``` -Five kinds of events are recorded: +Six kinds of events are recorded: * `to_host`: an actual device-to-host copy through conversion, mirror, staging, serial MPI, or host kernel helpers; @@ -298,7 +313,12 @@ Five kinds of events are recorded: * `fallback`: a `Kernel` without CUDA kernel calling its host kernel on the CuPy backend (`missing_cuda="fallback"`), one event per call, naming the kernel. Physical copies and a `kernel_conversion` marker are recorded separately; -* `device_copy`: a device-only dtype/layout conversion through CuNumpy helpers. +* `device_copy`: a device-only dtype/layout conversion through CuNumpy helpers; +* `sync`: the host waited for the device: `xp.synchronize()`, the waits of the MPI + helpers (`synchronize_for_mpi()`, `mpi_buffer()` staging) and of the CUDA debug + mode, and, on the fake CuPy, a scalar read of a device array (`float(a)`, + `int(a)`, `bool(a)`, `a.item()`, `a.tolist()`). They are in `counter.syncs` and in + the report, but not in `total`. Only real transfers count: `to_numpy()` of a NumPy array or `to_cupy()` of a CuPy array records nothing. The counter has the attributes `to_host`, @@ -327,16 +347,17 @@ not thread-safe. **Limitation:** only transfers made through CuNumpy are seen. Raw `cupy.ndarray.get()`, `cupy.asarray(numpy_array)`, `numpy.asarray(cupy_array)`, -`float(device_array)`, forwarded backend calls such as `xp.asarray()`, and -implicit conversions inside other libraries are -not counted. Use `nsys` (or CuPy's profiling hooks) to find those. +forwarded backend calls such as `xp.asarray()`, and implicit conversions inside +other libraries are not counted. Neither is `float(device_array)` (an implicit +sync) with the real CuPy, which cannot be observed from Python; the fake CuPy +counts it. Use `nsys` (or CuPy's profiling hooks) to find those. -### `profiling.assert_no_transfers()` +### `profiling.assert_no_transfers(*, syncs=False)` Context manager that raises `AssertionError` with the counter's `report()` if the block makes a host/device transfer or host fallback through CuNumpy. -Device-only dtype/layout conversions are allowed. It yields the `TransferCounter` -too. An exception raised inside the block propagates as it is: +Device-only dtype/layout conversions are allowed, and so are syncs unless +`syncs=True`. It yields the `TransferCounter` too. An exception raised inside the block propagates as it is: ```python def test_time_step_stays_on_the_device(): @@ -808,8 +829,18 @@ CuPy arrays. * `is_array`: predicate for host array values to convert back to CuPy. The default is `isinstance(value, numpy.ndarray)`. * `outputs`: sequence of arguments the kernel may write to. Entries are - positional indices or keyword names. If omitted, every converted array is + indices or parameter names; a name also finds a positional argument, and an + index a keyword argument, when the parameter names are known (from the + Python signature, or `parameters`). If omitted, every converted array is copied back. +* `parameters`: the names of the positional parameters, or a function that + returns them, for compiled kernels without a Python signature. A `Kernel` + supplies the names of its host function. + +`kernels.outputs_from_annotations(function)` returns the parameters a kernel may +write to, read from its annotations: everything that is not `Final`, `const` or a +scalar. Use it as `outputs=outputs_from_annotations(push)`, or pass +`outputs="annotations"` to `Kernel.from_folder()` / `KernelCatalog.from_package()`. ### Example and output declarations diff --git a/docs/source/guides/data-movement.md b/docs/source/guides/data-movement.md index 07f3f78..72d5980 100644 --- a/docs/source/guides/data-movement.md +++ b/docs/source/guides/data-movement.md @@ -165,3 +165,27 @@ with xp.cuda.stream(): Pinned memory is a limited system resource; use it for large, repeatedly transferred buffers after a profile shows transfers matter. `pin_memory()` requires CuPy. + +## Host-only code: `host_call` + +Some code can only run on the host: a SciPy spline, a file reader, an external +equilibrium code. `xp.host_call(fun, *args, **kwargs)` calls it with arguments of +either backend. Device arrays are copied to the host, `fun` runs on the NumPy +backend, and array results are copied back, once per call. `@xp.evaluate_on_host` +does the same for a method, and `@xp.setup_on_host` runs an `__init__` on the NumPy +backend, so the object holds only host data. The copies are counted by +`count_transfers()`. Use it for setup and diagnostics, not in a time loop. + +```python +values = xp.host_call(spline, x) # x on the device -> values on the device + + +class Equilibrium: + @xp.setup_on_host + def __init__(self, path): + self.spline = read_spline(path) # NumPy and SciPy + + @xp.evaluate_on_host + def pressure(self, x): + return self.spline(x) +``` diff --git a/docs/source/guides/execution-helpers.md b/docs/source/guides/execution-helpers.md index fc67400..8c6083b 100644 --- a/docs/source/guides/execution-helpers.md +++ b/docs/source/guides/execution-helpers.md @@ -142,10 +142,23 @@ per device; catalog setup catches CUDA compiler failures before timesteps. checks actual block/grid, kernel thread, and static-plus-dynamic shared-memory limits. These execution checks are independent of scope-profiler. -GPU CI requires a real CUDA device (`CUNUMPY_REQUIRE_CUDA=1`) and runs focused +GPU CI requires a real CUDA device (`CUNUMPY_REQUIRE_CUDA=1`; the GPU markers of +`cunumpy.kernel_testing` then fail instead of skipping) and runs focused `memcheck`, `racecheck`, and `synccheck` jobs with nonzero sanitizer error exits. Numerical parity tests are separate from performance comparisons. +## Keep the live particles: `compact_by_mask` + +```python +n = xp.algorithms.compact_by_mask(alive, markers, weights) +markers, weights = markers[:n], weights[:n] +``` + +Moves the rows where the boolean mask is True to the front of every array, in +place and in order, and returns their number. The rows after the first `n` are +unspecified. The count is needed on the host, so on CuPy each call +synchronizes once. + ## Prepare cell ranges and segment reductions ```python diff --git a/docs/source/kernels/pyccel-kernel.md b/docs/source/kernels/pyccel-kernel.md index 22db221..5be0abd 100644 --- a/docs/source/kernels/pyccel-kernel.md +++ b/docs/source/kernels/pyccel-kernel.md @@ -50,10 +50,16 @@ xp.kernels.PyccelKernel(solve, outputs=("out",)) # solve(a, b, out=out) xp.kernels.PyccelKernel(norm, outputs=()) # writes nothing ``` -* Positional arguments are declared by index (negative indices count from the - end), keyword arguments by name. The two forms are not interchangeable, - because compiled functions usually do not expose a Python signature that - would map names to positions. +* Arguments are declared by index (negative indices count from the end) or by + name. A name finds the argument also when it is passed positionally, and an + index also when it is passed as a keyword, as long as the parameter names are + known: from the Python signature, from `parameters=[...]`, or, for a kernel in a + `Kernel`, from its host function. For a compiled function without any of those + the two forms are not interchangeable. +* `xp.kernels.outputs_from_annotations(function)` reads the outputs from the + annotations: every parameter that is not `Final`, `const` or a scalar. A + `Kernel.from_folder(..., outputs="annotations")` (and `KernelCatalog.from_package`) + applies it to the host function of each kernel folder. * A container or object declared as output has all its arrays copied back. * **A missing declaration is a silent bug**: if the function writes an argument that is not declared, the device array keeps its old values. When in diff --git a/docs/source/kernels/testing.md b/docs/source/kernels/testing.md index 18b135c..ddc6d1c 100644 --- a/docs/source/kernels/testing.md +++ b/docs/source/kernels/testing.md @@ -95,8 +95,11 @@ Things to know: * **Build random data on the host.** NumPy and CuPy generators produce different sequences from the same seed, so use `numpy.random.default_rng` and convert with `to_cunumpy()`, as above. -* **Which arguments are compared**: `outputs=(2,)` selects them by index; - otherwise the host kernel's declared `outputs` are used, and if there are +* **Which arguments are compared**: `outputs=(2,)` selects them by index and + `outputs=("markers",)` by parameter name (the names of the host function). + `"markers.positions"` compares only that field of a struct or argument + object, leaving out fields the two kernels fill differently (scratch buffers, + for instance); otherwise the host kernel's declared `outputs` are used, and if there are none, every array argument. Arrays held by argument objects (one level deep, e.g. a `CudaArguments` object or a list) are compared too. A `CudaStructArguments` object is read through its struct fields, so its @@ -216,6 +219,40 @@ intrinsics; kernels using the latter are refused with `NotImplementedError`, so those still need a GPU run. The compiler may fuse multiply-adds as NVRTC does, so compare with a tolerance of a few ulp. +### Struct parameters and `emulated_launches()` + +A kernel with struct parameters takes, for each struct, a dictionary of field +names to values, a `CudaStructValue`, or any object with an attribute per field +(a host argument class, a `CudaStructArguments`); the arrays in it are updated in +place: + +```python +emulate_cuda_kernel( + push, + {"markers": markers, "alive": alive, "n": 4}, # the struct Particles + 0.5, + total, + n_threads=4, +) +``` + +To test code that *launches* kernels (a `Kernel` on the CuPy backend, a +propagator), wrap it in `emulated_launches()`. With the fake CuPy (below) every +`CudaKernel` launch in the block then runs through the emulation on the host +buffers of the fake arrays, so the device arrays hold the results: + +```python +from cunumpy.kernel_testing import emulated_launches, host_buffer + +with emulated_launches(): + propagator(dt) # CuPy backend = the fake CuPy +np.testing.assert_allclose(host_buffer(markers), expected) +``` + +`host_buffer(array)` is the NumPy array behind a fake CuPy array (not a copy). +Launches are serial and compile the kernel as C++ every time, so use small +problems. + ## Test `__device__` helpers: `device_function_kernel` Helpers such as B-spline evaluation or coordinate maps are `__device__` @@ -314,6 +351,13 @@ This catches `to_numpy()` calls, `PyccelKernel` conversions and `Kernel` fallbacks that crept into the step. It does not see copies made outside CuNumpy (see [Data movement](../guides/data-movement.md)). +Syncs (the host waiting for the device) are recorded as well, in +`counter.syncs`, and are accepted unless you ask for none: +`assert_no_transfers(syncs=True)`. They include `xp.synchronize()` and the waits +of the MPI helpers; on the fake CuPy also `float(a)`, `int(a)`, `bool(a)`, +`a.item()` and `a.tolist()`, which stall the real CuPy too but cannot be +observed there from Python. + ## Test generated headers When struct headers are generated with `write_cuda_header()` and committed, @@ -324,6 +368,12 @@ offsets. ## CI setup +* With `CUNUMPY_REQUIRE_CUDA=1` the GPU markers fail instead of skipping: + `requires_cupy` (as an error when the test is set up), the `cupy` run of the + `backend` fixture (which also activates CuPy strictly, never falling back to + NumPy) and `assert_kernels_agree`. Set it on the GPU CI job, so a broken CuPy + or driver cannot pass as a set of skipped tests. + * Run the suite on a normal CPU runner: everything on NumPy runs, GPU cases are reported as skipped, and `emulate_cuda_kernel` tests check the CUDA kernels' arithmetic (the runner needs a C++ compiler, which Linux images diff --git a/src/cunumpy/LLM_GUIDE.md b/src/cunumpy/LLM_GUIDE.md index 54ad69d..b0db197 100644 --- a/src/cunumpy/LLM_GUIDE.md +++ b/src/cunumpy/LLM_GUIDE.md @@ -73,6 +73,10 @@ https://max-models.github.io/cunumpy/ and in `docs/source/` of the repository. | call an existing NumPy-only kernel with GPU arrays (slow, correct) | `xp.kernels.PyccelKernel(fn, outputs=(...))` | | launch a hand-written CUDA C kernel | `xp.kernels.CudaKernel(source, "name")` / `CudaKernel.from_file(path)` | | run a Metal (MSL) kernel on an Apple silicon GPU, NumPy float32 in and out (float64 raises; `float64="cast"` computes in float32); not part of `Kernel` dispatch | `xp.kernels.MetalKernel(body, inputs=[...], outputs=[...])(*args, out=arrays, n_threads=n)`; check `xp.kernels.metal_available()` | +| call host-only code (SciPy, file readers) with arguments of either backend | `xp.host_call(fun, *args)`; `@xp.evaluate_on_host` on a method; `@xp.setup_on_host` on `__init__` | +| keep the live rows of particle arrays at the front | `n = xp.algorithms.compact_by_mask(alive, markers, weights)` | +| run CUDA kernel launches on the CPU in a test (fake CuPy) | `with kernel_testing.emulated_launches(): ...`; `kernel_testing.host_buffer(a)` reads a fake array; struct arguments are read through their fields | +| find the arguments a host kernel writes (copy only those back) | `PyccelKernel(fn, outputs=("out",))` (names work positionally); `xp.kernels.outputs_from_annotations(fn)`; `Kernel.from_folder(..., outputs="annotations")` | | host kernel + CUDA port, chosen by backend | `xp.kernels.Kernel(host_fn, cuda_kernel_or_None)` | | many kernels in a package, ported incrementally | `xp.kernels.KernelCatalog.from_package(__name__, missing_cuda="fallback")` | | host kernels compiled at first call (your compile function), NumPy fallback | `from_package(..., host_suffix="_pyccel", compile_host=my_compile, host_fallback={...})` -> `xp.kernels.CompiledHostKernel` | diff --git a/src/cunumpy/_algorithms.py b/src/cunumpy/_algorithms.py index 353c83a..ab79486 100644 --- a/src/cunumpy/_algorithms.py +++ b/src/cunumpy/_algorithms.py @@ -251,3 +251,49 @@ def sort_by_key(keys: Any, *arrays: Any) -> tuple[Any, ...]: ) order = xpm.argsort(keys, kind="stable").astype(xpm.int64, copy=False) return (keys[order], order, *(array[order] for array in arrays)) + + +def compact_by_mask(mask: Any, *arrays: Any) -> int: + """Move the rows where `mask` is True to the front of every array, in place. + + The usual step after particles left the domain or were absorbed: keep the + live ones at the front of the marker array (and of the arrays that go with + it) and continue with ``markers[:n]``. The order of the kept rows is + preserved, so the result is reproducible:: + + n = xp.algorithms.compact_by_mask(alive, markers, weights) + markers, weights = markers[:n], weights[:n] + + Parameters + ---------- + mask : array of bool, shape (n,) + True for the rows to keep. + *arrays : arrays + Arrays with ``n`` rows (any further axes), on the backend of `mask`. + Rows ``[:count]`` hold the kept rows afterwards; the rows after them + are unspecified, so ignore them (or overwrite them). + + Returns + ------- + int + The number of kept rows. Its value is needed on the host, so on CuPy + the call synchronizes once per call (counted by + :func:`~cunumpy.profiling.count_transfers` where it can be seen). + """ + assert_same_backend(mask, *arrays) + xpm = get_array_module(mask) + mask = xpm.asarray(mask) + if mask.ndim != 1 or mask.dtype != np.bool_: + raise TypeError( + f"mask must be a 1D boolean array, got dtype {mask.dtype}, {mask.ndim}D", + ) + for i, array in enumerate(arrays): + if array.ndim < 1 or array.shape[0] != mask.shape[0]: + raise ValueError( + f"array {i} has shape {array.shape}, expected {mask.shape[0]} rows", + ) + rows = xpm.nonzero(mask)[0] + n_kept = int(rows.size) + for array in arrays: + array[:n_kept] = array[rows] # the right side is a copy: no overlap problem + return n_kept diff --git a/src/cunumpy/_cuda_kernel.py b/src/cunumpy/_cuda_kernel.py index e072149..7c3bd8b 100644 --- a/src/cunumpy/_cuda_kernel.py +++ b/src/cunumpy/_cuda_kernel.py @@ -66,6 +66,9 @@ class is the one definition of the arguments. import numpy as np +from cunumpy._transfers import _ACTIVE as _COUNTERS +from cunumpy._transfers import _record_sync + __all__ = [ "DEBUG_OPTIONS", "CudaArguments", @@ -2532,6 +2535,8 @@ def _synchronize_after_launch( stream = cp.cuda.get_current_stream() if _is_capturing(stream): return + if _COUNTERS: + _record_sync(f"debug synchronization after kernel {self.expression!r}") try: stream.synchronize() except Exception as error: diff --git a/src/cunumpy/_dispatch.py b/src/cunumpy/_dispatch.py index ed3e388..305082f 100644 --- a/src/cunumpy/_dispatch.py +++ b/src/cunumpy/_dispatch.py @@ -40,6 +40,7 @@ HostImplementations, PyccelKernel, get_device_kernel_implementation, + outputs_from_annotations, ) from cunumpy._transfers import _ACTIVE as _COUNTERS from cunumpy._transfers import _record @@ -219,6 +220,10 @@ def __init__( "host_options are for wrapping a plain callable; configure the " "given PyccelKernel directly", ) + if host_kernel._parameters is None: + # compiled kernels have no Python signature: name the parameters + # from the host function, so that outputs may be given by name + host_kernel._parameters = self.host_parameters self._host_kernel = host_kernel self._cuda_kernel = cuda_kernel self._name = name if name is not None else host_kernel.name @@ -239,6 +244,7 @@ def from_folder( check_name_length: bool = True, missing_cuda: str = "raise", host_options: Mapping[str, Any] | None = None, + outputs: Sequence[int | str] | str | None = None, include_dirs: Sequence[str | Path] | None = None, dispatch: str = "backend", compile_host: Callable[[Any], Any] | None = None, @@ -285,6 +291,14 @@ def from_folder( As for :meth:`KernelCatalog.from_package`. missing_cuda, host_options, dispatch Passed on to :class:`Kernel`. + outputs : Sequence[int | str] | "annotations" | None + The arguments the host kernel writes to (names or indices, see + :class:`~cunumpy.kernels.PyccelKernel`), so that only those arrays are + copied back to the device on the fallback path. ``"annotations"`` + reads them from the annotations of the host function: every + parameter that is not ``Final`` or a scalar is an output (see + :func:`~cunumpy.kernels.outputs_from_annotations`). A value in + `host_options` takes precedence. include_dirs : Sequence[str | Path] | None Include directories of the CUDA kernel, in addition to the folder itself; by default the source root of the top-level package (the @@ -334,6 +348,19 @@ def from_folder( test_args = f"{package}.{name}{test_args_suffix}" module = importlib.import_module(f"{package}.{name}{host_suffix}") python = getattr(module, name) + if outputs is not None and "outputs" not in (host_options or {}): + declared = ( + outputs_from_annotations(python) + if isinstance(outputs, str) and outputs == "annotations" + else outputs + ) + if isinstance(outputs, str) and outputs != "annotations": + raise ValueError( + f"outputs must be a sequence of names/indices or 'annotations', " + f"got {outputs!r}", + ) + if declared is not None: + host_options = {**(host_options or {}), "outputs": declared} loaders: dict[str, Callable[[], Callable[..., Any]]] = { "python": lambda: python, "pyccel": ( @@ -669,6 +696,12 @@ def from_package( host_options: ( Mapping[str, Any] | Callable[[str], Mapping[str, Any]] | None ) = None, + outputs: ( + Sequence[int | str] + | str + | Callable[[str], Sequence[int | str] | str | None] + | None + ) = None, include_dirs: Sequence[str | Path] | None = None, dispatch: str = "backend", compile_host: Callable[[Any], Any] | None = None, @@ -713,6 +746,10 @@ def from_package( Keyword arguments for the :class:`~cunumpy.kernels.PyccelKernel` wrapping each host kernel (see :class:`Kernel`): the same for all kernels, or a function of the kernel name, e.g. to declare per-kernel ``outputs``. + outputs : Sequence[int | str] | "annotations" | Callable | None + The arguments every host kernel writes to, or ``"annotations"`` to + read them from the annotations of each host function, or a function + of the kernel name returning either (see :meth:`Kernel.from_folder`). include_dirs : Sequence[str | Path] | None Include directories of the CUDA kernels, in addition to each kernel's own folder. By default the source root of the top-level @@ -757,6 +794,7 @@ def from_package( host_options=( host_options(name) if callable(host_options) else host_options ), + outputs=outputs(name) if callable(outputs) else outputs, include_dirs=include_dirs, dispatch=dispatch, compile_host=compile_host, diff --git a/src/cunumpy/_emulation.py b/src/cunumpy/_emulation.py index 56d47d0..dab6c43 100644 --- a/src/cunumpy/_emulation.py +++ b/src/cunumpy/_emulation.py @@ -28,11 +28,22 @@ GPU. Per-block deposits, shared-memory reductions and tiled kernels therefore work. +Struct parameters take, for each struct, a mapping from field name to value, a +:class:`~cunumpy.arguments.CudaStructValue`, or any object with an attribute +per field (e.g. a :class:`~cunumpy.arguments.CudaStructArguments` or a host +argument class); the arrays of the fields are updated in place like the others. + +:func:`emulated_launches` makes every :class:`~cunumpy.kernels.CudaKernel` launch +in a block run through the emulation, on the arrays of the fake CuPy +(:mod:`cunumpy._fake_cupy`), so that code that launches kernels (a +:class:`~cunumpy.kernels.Kernel` on the CuPy backend, a propagator) can be +tested without a GPU. + What it does not emulate: concurrency between the barriers (threads run one after another, so atomics are plain additions and races never show), warp intrinsics (``__shfl_*``, ``__syncwarp``, ``__ballot_sync``, ...; a kernel or an -included header using them is refused), structs and ``CudaArguments`` objects -(not supported), and ````. Use a GPU for those. +included header using them is refused), and ````. Use a GPU +for those. Floating point: like NVRTC (``--fmad=true`` by default) the C++ compiler may fuse ``a * b + c`` into one fused multiply-add, so results can differ from @@ -49,21 +60,25 @@ import shutil import subprocess import tempfile -from collections.abc import Sequence +from collections.abc import Generator, Mapping, Sequence +from contextlib import contextmanager from pathlib import Path from typing import Any import numpy as np +from cunumpy import _fake_cupy from cunumpy._cuda_kernel import ( CudaKernel, CudaParameter, + CudaStructArguments, + CudaStructValue, _scalar_checker, _strip_comments, cuda_include_dir, ) -__all__ = ["emulate_cuda_kernel", "emulation_compiler"] +__all__ = ["emulate_cuda_kernel", "emulated_launches", "emulation_compiler"] # CUDA constructs that serial emulation would get wrong _UNSUPPORTED = { @@ -258,6 +273,76 @@ def _check_supported(kernel: CudaKernel) -> None: ) +def _host(value: Any) -> Any: + """The NumPy buffer of a fake CuPy array; anything else as it is.""" + if _fake_cupy.is_active(): + try: + return _fake_cupy.host_buffer(value) + except TypeError: + pass + return value + + +def _field_value(struct_name: str, value: Any, field: CudaParameter) -> Any: + """The value of the struct field `field` in the struct argument `value`.""" + try: + if isinstance(value, (CudaStructValue, Mapping)): + return value[field.name] + return getattr(value, field.name) + except (KeyError, AttributeError): + raise TypeError( + f"struct argument {struct_name!r} has no value for the field " + f"{field.name!r} ({type(value).__name__})", + ) from None + + +def _declaration(param: CudaParameter, name: str) -> str: + if param.view_ndim is not None: + return f"{param.ctype} {name}" + return f"{param.ctype}{'*' if param.pointer else ''} {name}" + + +def _flatten_structs( + kernel: CudaKernel, + args: Sequence[Any], +) -> tuple[CudaKernel, list[Any]]: + """A kernel that takes every struct field as a parameter, and its arguments. + + The wrapper kernel (the source of `kernel` plus a ``__global__`` function with + one parameter per field) rebuilds the structs and calls `kernel`, so that the + emulation only has to deal with arrays and scalars. + """ + params, builds, call, flat = [], [], [], [] + for p, value in zip(kernel.signature, args): + if p.struct is None: + params.append(_declaration(p, p.name)) + call.append(p.name) + flat.append(value) + continue + names = [f"{p.name}__{f.name}" for f in p.struct.fields] + params += [_declaration(f, n) for f, n in zip(p.struct.fields, names)] + builds.append(f" {p.struct.name} {p.name}_struct{{{', '.join(names)}}};") + call.append(f"{p.name}_struct") + flat += [_field_value(p.name, value, f) for f in p.struct.fields] + name = f"emulated_{kernel.name}" + source = ( + kernel.source + + f'\nextern "C" __global__ void {name}({", ".join(params)}) {{\n' + + "\n".join(builds) + + f"\n {kernel.expression}({', '.join(call)});\n}}\n" + ) + wrapper = CudaKernel( + source, + name, + block_size=kernel.block_size, + options=[o for o in kernel.options if not o.startswith("-I")], + include_dirs=kernel.include_dirs, + source_dir=kernel.source_dir, + n_threads_from=kernel.n_threads_from, + ) + return wrapper, flat + + def emulate_cuda_kernel( kernel: CudaKernel, *args: Any, @@ -275,8 +360,11 @@ def emulate_cuda_kernel( kernel : CudaKernel The kernel; its signature must be parsed (the default). *args - The kernel arguments with NumPy arrays in place of CuPy arrays. Arrays - are updated in place with what the kernel wrote. + The kernel arguments with NumPy arrays in place of CuPy arrays (arrays of + the fake CuPy are used through their host buffer). Arrays are updated in + place with what the kernel wrote. A struct parameter takes a mapping of + field names to values, a :class:`~cunumpy.arguments.CudaStructValue`, or + an object with an attribute per field. n_threads, grid, block Launch shape, as for :meth:`CudaKernel.__call__`. compiler : str | None @@ -291,8 +379,8 @@ def emulate_cuda_kernel( Raises ------ NotImplementedError - If the kernel (or a header it includes) uses warp intrinsics, or the - kernel has struct parameters or complex scalars. + If the kernel (or a header it includes) uses warp intrinsics, or has + complex scalars. TypeError If an argument does not match its parameter (dtype, dimensions, a scalar that does not fit). @@ -309,6 +397,19 @@ def emulate_cuda_kernel( raise TypeError( f"kernel {kernel.name!r} takes {len(params)} arguments, got {len(args)}", ) + if any(p.struct is not None for p in params): + wrapper, flat = _flatten_structs(kernel, args) + return emulate_cuda_kernel( + wrapper, + *flat, + n_threads=n_threads, + grid=grid, + block=block, + compiler=compiler, + options=options, + shared_mem=shared_mem, + ) + args = tuple(_host(a) for a in args) compiler = compiler or emulation_compiler() if compiler is None: raise RuntimeError("emulation needs a C++ compiler (set CXX or install c++)") @@ -323,10 +424,6 @@ def emulate_cuda_kernel( globals_, inits, call_args, writes, arrays = [], [], [], [], [] for i, (param, value) in enumerate(zip(params, args)): name = f"cunumpy_arg{i}" - if param.struct is not None: - raise NotImplementedError( - "emulation does not support struct parameters", - ) if param.pointer or param.view_ndim is not None: if not isinstance(value, np.ndarray): raise TypeError( @@ -444,3 +541,74 @@ def emulate_cuda_kernel( for value, buffer, out in arrays: result = np.fromfile(out, dtype=buffer.dtype).reshape(buffer.shape) value[...] = result + + +@contextmanager +def emulated_launches( + *, + compiler: str | None = None, + options: Sequence[str] = (), +) -> Generator[None, None, None]: + """Run every :class:`~cunumpy.kernels.CudaKernel` launch in the block on the CPU. + + With the fake CuPy (:func:`cunumpy.kernel_testing.install_fake_cupy`) + CUDA kernels cannot run. Inside this block a launch is emulated instead + (:func:`emulate_cuda_kernel`), on the host buffers of the fake CuPy arrays + it is given, so code that launches kernels runs on a machine without a GPU + and its device arrays hold the results afterwards:: + + with emulated_launches(): + propagator(dt) # CuPy backend: the CUDA kernels run on the CPU + + The launch arguments are those of :meth:`CudaKernel.__call__` + (`stream` is ignored). Argument objects are flattened like in a launch; + struct values and :class:`~cunumpy.arguments.CudaStructArguments` objects + stay one argument and are read through their fields. Arrays must be arrays + of the fake CuPy or NumPy arrays. + + Launches are serial and slow (each one compiles the kernel as C++), so use it + on small problems. The limits of :func:`emulate_cuda_kernel` apply. + + Parameters + ---------- + compiler : str | None + C++ compiler; by default :func:`emulation_compiler`. + options : Sequence[str] + Additional compiler options for every launch, e.g. + ``("-ffp-contract=off",)``. + """ + original = CudaKernel.__call__ + + def launch( + self: CudaKernel, + *args: Any, + n_threads: int | Sequence[int] | None = None, + grid: int | Sequence[int] | None = None, + block: int | Sequence[int] | None = None, + shared_mem: int = 0, + stream: Any = None, + ) -> None: + flat: list[Any] = [] + for arg in args: + if isinstance(arg, (CudaStructValue, CudaStructArguments)): + flat.append(arg) + elif hasattr(arg, "__cuda_args__"): + flat.extend(arg.__cuda_args__()) + else: + flat.append(arg) + emulate_cuda_kernel( + self, + *flat, + n_threads=n_threads, + grid=grid, + block=block, + compiler=compiler, + options=options, + shared_mem=shared_mem, + ) + + CudaKernel.__call__ = launch # type: ignore[method-assign] + try: + yield + finally: + CudaKernel.__call__ = original # type: ignore[method-assign] diff --git a/src/cunumpy/_fake_cupy.py b/src/cunumpy/_fake_cupy.py index 6f84e4a..47d3e09 100644 --- a/src/cunumpy/_fake_cupy.py +++ b/src/cunumpy/_fake_cupy.py @@ -17,7 +17,9 @@ :class:`~cunumpy.kernels.CudaKernel` work; * CUDA kernels cannot run: ``RawKernel`` and friends raise ``NotImplementedError`` when called, and :func:`cunumpy.kernel_testing.requires_cupy` - skips tests while the fake is active. + skips tests while the fake is active. Inside + :func:`cunumpy.kernel_testing.emulated_launches` a ``CudaKernel`` launch runs + on the CPU instead, on the host buffers of the arrays (:func:`host_buffer`). Activate it before CuPy or cunumpy's backend is first used, either with the environment variable ``CUNUMPY_FAKE_CUPY=1`` (read when cunumpy is imported) @@ -35,7 +37,7 @@ import types from pathlib import Path -__all__ = ["install", "is_active", "uninstall"] +__all__ = ["host_buffer", "install", "is_active", "uninstall"] _IMPLEMENTATION = Path(__file__).with_name("_fake_cupy_impl.py") _SUBMODULES = ("cupy.cuda", "cupy.cuda.device", "cupy.cuda.runtime", "cupy.linalg") @@ -47,6 +49,28 @@ def is_active() -> bool: return bool(getattr(module, "__cunumpy_fake__", False)) +def host_buffer(array: object) -> object: + """The NumPy array that holds the data of a fake ``cupy.ndarray``. + + Not a copy: writing to the buffer changes the fake device array, so a + kernel emulation (:func:`cunumpy.kernel_testing.emulated_launches`) can + update the array in place. A test can also read results with it without a + ``.get()``, which is the point of a "device" array being visible in tests. + + Raises + ------ + TypeError + If `array` is not an array of the fake CuPy (including when the fake + is not active). + """ + module = sys.modules.get("cupy") + if is_active() and isinstance(array, module.ndarray): + return array._a + raise TypeError( + f"host_buffer needs an array of the fake CuPy, got {type(array).__name__}", + ) + + def install() -> types.ModuleType: """Install the fake ``cupy`` package into ``sys.modules`` and return it. diff --git a/src/cunumpy/_fake_cupy_impl.py b/src/cunumpy/_fake_cupy_impl.py index 0a4a13a..952cffa 100644 --- a/src/cunumpy/_fake_cupy_impl.py +++ b/src/cunumpy/_fake_cupy_impl.py @@ -15,6 +15,14 @@ _HOST = _np.ndarray +def _sync(what): + """Count a scalar read of a device array (an implicit sync of the real CuPy).""" + from cunumpy._transfers import _ACTIVE, _record_sync + + if _ACTIVE: + _record_sync(what) + + def _err(obj, where=""): return TypeError( f"Unsupported type {type(obj)}{where} (fake CuPy: host arrays/lists are " @@ -75,6 +83,8 @@ def __getattr__(self, name): # no __array_interface__ etc.: NumPy must not see the host buffer if name.startswith("__"): raise AttributeError(name) + if name in ("item", "tolist"): + _sync(f"ndarray.{name}()") attr = getattr(self._a, name) if callable(attr): return _wrap_callable(attr, strict=True, name=f"ndarray.{name}") @@ -101,18 +111,23 @@ def __format__(self, spec): return format(self._a.item() if self._a.ndim == 0 else self._a, spec) def __bool__(self): + _sync("bool(device array)") return bool(self._a) def __int__(self): + _sync("int(device array)") return int(self._a) def __float__(self): + _sync("float(device array)") return float(self._a) def __complex__(self): + _sync("complex(device array)") return complex(self._a) def __index__(self): + _sync("index of a device array") return operator.index(self._a.item() if self._a.ndim == 0 else self._a) __hash__ = None diff --git a/src/cunumpy/_kernel.py b/src/cunumpy/_kernel.py index a87c2de..735d600 100644 --- a/src/cunumpy/_kernel.py +++ b/src/cunumpy/_kernel.py @@ -17,6 +17,7 @@ import copy import importlib +import inspect import os import warnings from collections.abc import Callable, Generator, Iterator, Mapping, Sequence @@ -73,11 +74,21 @@ class PyccelKernel: interpolate = PyccelKernel(some_interpolation_kernel, outputs=(5,)) interpolate(x, y, z, basis, coeffs, out) # `out` is argument 5 - Pyccel-compiled kernels are builtins with no introspectable signature, - so an index and a name are *not* interchangeable: declare the form you - actually call with. An empty sequence declares that the kernel writes to - none of its arguments. By default (``None``) every converted array is - copied back, which is always correct but does more work. + A name also finds the argument when it is passed positionally, and an + index also finds a keyword argument, if the parameter names of the + kernel are known: from its Python signature, or from `parameters` + (Pyccel-compiled kernels are builtins with no introspectable + signature; a :class:`~cunumpy.kernels.Kernel` supplies the names of + its host function). Without parameter names an index and a name are + *not* interchangeable: declare the form you actually call with. An + empty sequence declares that the kernel writes to none of its + arguments. By default (``None``) every converted array is copied back, + which is always correct but does more work. + parameters : sequence of str or callable, optional + The names of the positional parameters of `kernel`, or a function that + returns them (or None if they are unknown), for resolving `outputs` + names and indices. By default they are read from the signature of + `kernel`, where it has one. Examples -------- @@ -93,8 +104,10 @@ def __init__( object_modules: Sequence[str] = (), is_array: Callable[[Any], bool] | None = None, outputs: Sequence[int | str] | None = None, + parameters: Sequence[str] | Callable[[], Sequence[str] | None] | None = None, ) -> None: self._kernel = kernel + self._parameters = parameters self._use_cupy = use_cupy self._object_modules = tuple(object_modules) self._is_array = is_array or (lambda value: isinstance(value, np.ndarray)) @@ -122,6 +135,23 @@ def __repr__(self) -> str: f"outputs={self._outputs!r})" ) + def _parameter_names(self) -> Sequence[str] | None: + """The names of the positional parameters, or None if they are unknown.""" + source = self._parameters + if source is not None: + return source() if callable(source) else tuple(source) + function = getattr(self._kernel, "python", self._kernel) + try: + signature = inspect.signature(function) + except (TypeError, ValueError): + return None + positional = ( + inspect.Parameter.POSITIONAL_ONLY, + inspect.Parameter.POSITIONAL_OR_KEYWORD, + ) + names = [p.name for p in signature.parameters.values() if p.kind in positional] + return names or None + def _convert_to_numpy( self, value: Any, @@ -261,31 +291,61 @@ def _output_host_arrays( IndexError, KeyError If a declared output does not correspond to an argument of this call -- typically because an argument declared by index was passed - as a keyword, or vice versa. + as a keyword, or vice versa, and the parameter names are unknown. """ found: set[int] = set() seen: set[int] = set() + names = self._parameter_names() if self._outputs else None for entry in self._outputs or (): if isinstance(entry, int): index = entry + len(args_np) if entry < 0 else entry - if not 0 <= index < len(args_np): - raise IndexError( - f"{self.name}() was declared with output argument " - f"{entry}, but was called with {len(args_np)} " - "positional argument(s). Note that an output passed as " - "a keyword must be declared by name, not by index.", - ) - self._collect_host_arrays(args_np[index], found, seen) - else: - if entry not in kwargs_np: - raise KeyError( - f"{self.name}() was declared with output argument " - f"{entry!r}, but no such keyword argument was passed. " - "Note that an output passed positionally must be " - "declared by index, not by name.", - ) + if 0 <= index < len(args_np): + self._collect_host_arrays(args_np[index], found, seen) + continue + if ( + names is not None + and 0 <= entry < len(names) + and names[entry] in kwargs_np + ): + self._collect_host_arrays(kwargs_np[names[entry]], found, seen) + continue + raise IndexError( + f"{self.name}() was declared with output argument " + f"{entry}, but was called with {len(args_np)} " + "positional argument(s)" + + ( + "" + if names is not None + else ". Note that an output passed as a keyword must " + "be declared by name, not by index (the parameter " + "names of the kernel are unknown)." + ), + ) + if entry in kwargs_np: self._collect_host_arrays(kwargs_np[entry], found, seen) + continue + if ( + names is not None + and entry in names + and names.index(entry) + < len( + args_np, + ) + ): + self._collect_host_arrays(args_np[names.index(entry)], found, seen) + continue + raise KeyError( + f"{self.name}() was declared with output argument " + f"{entry!r}, but " + + ( + "no argument of that name was passed" + if names is not None + else "no such keyword argument was passed. Note that an " + "output passed positionally must be declared by index, not " + "by name (the parameter names of the kernel are unknown)." + ), + ) return found @@ -783,3 +843,51 @@ def kernel_output(out: Any, like: Any, dtype: Any = None) -> Generator[Any]: out[...] = to_numpy(buffer) else: out[...] = buffer + + +_SCALAR_ANNOTATIONS = {"int", "float", "bool", "complex", "str"} + + +def outputs_from_annotations(function: Callable[..., Any]) -> tuple[str, ...] | None: + """The names of the parameters a kernel may write to, read from its annotations. + + Pyccel kernels mark what they only read with ``Final``: ``x: "Final[float[:]]"``. + A parameter is an output unless it is annotated ``Final`` (or ``const``) or + has the annotation of a scalar (``int``, ``float``, ``bool``, ``complex``, + ``str``); an annotated array or an argument object may be written to. Use the + result as the `outputs` of a :class:`PyccelKernel` so that only those arrays + are copied back to the device:: + + PyccelKernel(push, outputs=outputs_from_annotations(push)) + + Returns + ------- + tuple[str, ...] | None + The parameter names, in order, or None if `function` has no signature or + none of its parameters is annotated (so nothing can be said). + """ + try: + parameters = inspect.signature(function).parameters.values() + except (TypeError, ValueError): + return None + if all(p.annotation is inspect.Parameter.empty for p in parameters): + return None + outputs = [] + for p in parameters: + if isinstance(p.annotation, str): + text = p.annotation + elif isinstance(p.annotation, type): + text = p.annotation.__name__ + else: + text = str(p.annotation) + text = text.replace("typing.", "").strip().strip("'\"") + if p.annotation is inspect.Parameter.empty: + outputs.append(p.name) # nothing is known: assume it may be written + elif ( + text.startswith(("Final[", "const ", "Final ")) + or text in _SCALAR_ANNOTATIONS + ): + continue + else: + outputs.append(p.name) + return tuple(outputs) diff --git a/src/cunumpy/_mpi.py b/src/cunumpy/_mpi.py index 8dfaf98..fd46415 100644 --- a/src/cunumpy/_mpi.py +++ b/src/cunumpy/_mpi.py @@ -13,7 +13,7 @@ from maybempi import get_mpi from cunumpy._transfers import _ACTIVE as _COUNTERS -from cunumpy._transfers import _describe, _nbytes, _record +from cunumpy._transfers import _describe, _nbytes, _record, _record_sync from cunumpy.xp import array_backend, cupy_available, to_numpy _logger = logging.getLogger(__name__) @@ -45,6 +45,8 @@ def synchronize_for_mpi(*arrays: Any, stream: Any = None, event: Any = None) -> if isinstance(event, HostEvent) or isinstance(stream, HostStream): raise TypeError("device buffers require a CUDA producer stream or event") + if _COUNTERS: + _record_sync("synchronize_for_mpi()") if event is not None: event.synchronize() return @@ -262,6 +264,8 @@ def mpi_buffer( # Ensure MPI's host buffer can be reused immediately on context exit. if transfer_array is not array: array[...] = transfer_array + if _COUNTERS: + _record_sync("mpi_buffer() staging for recv") cp.cuda.get_current_stream().synchronize() diff --git a/src/cunumpy/_transfers.py b/src/cunumpy/_transfers.py index f2d642e..e7ad0c3 100644 --- a/src/cunumpy/_transfers.py +++ b/src/cunumpy/_transfers.py @@ -26,7 +26,13 @@ device arrays to the host (and back), one event per call; * ``fallback``: a :class:`~cunumpy.kernels.Kernel` without CUDA kernel calling its host kernel on the CuPy backend (``missing_cuda="fallback"``), one event per call; -* ``device_copy``: device-only dtype/layout conversions in CuNumpy helpers. +* ``device_copy``: device-only dtype/layout conversions in CuNumpy helpers; +* ``sync``: the host waited for the device: :func:`~cunumpy.synchronize`, the + waits of the MPI helpers and of the CUDA debug mode, and, on the fake CuPy, + a scalar read of a device array (``float(a)``, ``int(a)``, ``bool(a)``, + ``a.item()``, ``a.tolist()``). Syncs are listed in :attr:`TransferCounter.syncs` + and the report but not in :attr:`~TransferCounter.total`, and + :func:`assert_no_transfers` accepts them unless called with ``syncs=True``. Mirror refreshes, argument conversions, staging, serial MPI and kernel output copy-back are also counted. Each physical host/device copy is recorded with its @@ -37,7 +43,9 @@ Limitations ----------- Only transfers made *through cunumpy* are seen. Raw ``cupy.ndarray.get()``, -``cupy.asarray(numpy_array)``, ``numpy.asarray(cupy_array)``, ``float(device_array)``, +``cupy.asarray(numpy_array)``, ``numpy.asarray(cupy_array)``, ``float(device_array)`` +(an implicit sync of the real CuPy, which cannot be observed from Python; the fake +CuPy reports it), forwarded backend operations such as ``xp.asarray`` and implicit conversions inside other libraries are not counted; use ``nsys`` (or CuPy's own profiling hooks) to find those. @@ -65,7 +73,14 @@ ] #: Event kinds, in the order they are reported. -KINDS = ("to_host", "to_device", "kernel_conversion", "fallback", "device_copy") +KINDS = ( + "to_host", + "to_device", + "kernel_conversion", + "fallback", + "device_copy", + "sync", +) # The currently active counters, innermost last. Instrumented code checks # ``if _ACTIVE:`` before doing any work, so the overhead of an inactive counter @@ -151,10 +166,15 @@ def fallbacks(self) -> int: """Number of `Kernel` calls that fell back to the host kernel on CuPy.""" return self.count("fallback") + @property + def syncs(self) -> int: + """Number of times the host waited for the device (see the module documentation).""" + return self.count("sync") + @property def total(self) -> int: - """Number of observations, including conversion/fallback markers.""" - return len(self.events) + """Number of observations, including conversion/fallback markers (not syncs).""" + return sum(1 for event in self.events if event.kind != "sync") def bytes(self, kind: str) -> int: """Known bytes copied for `kind`; markers and unknown sizes add zero.""" @@ -212,6 +232,11 @@ def _record(kind: str, description: str, *, nbytes: int | None = None) -> None: counter._add(event) +def _record_sync(description: str) -> None: + """Record that the host waits for the device (call only ``if _ACTIVE:``).""" + _record("sync", description) + + def _describe(array: Any) -> str: """``shape=..., dtype=...`` of an array, or the type name otherwise.""" shape = getattr(array, "shape", None) @@ -277,9 +302,15 @@ def count_transfers() -> Generator[TransferCounter, None, None]: @contextmanager -def assert_no_transfers() -> Generator[TransferCounter, None, None]: +def assert_no_transfers( + *, syncs: bool = False +) -> Generator[TransferCounter, None, None]: """Raise ``AssertionError`` if the block makes a transfer through cunumpy. + Syncs (the host waiting for the device) are accepted unless `syncs` is + True, e.g. ``assert_no_transfers(syncs=True)`` for a time step that must + never stall on the device. + A :func:`count_transfers` block that, on exit, raises with the counter's :meth:`~TransferCounter.report` if a host/device copy or host conversion/ fallback was counted. Device-only conversions are allowed. Only checked if @@ -292,7 +323,8 @@ def assert_no_transfers() -> Generator[TransferCounter, None, None]: """ with count_transfers() as counter: yield counter - if any(event.kind != "device_copy" for event in counter.events): + ignored = ("device_copy",) if syncs else ("device_copy", "sync") + if any(event.kind not in ignored for event in counter.events): raise AssertionError( "host/device transfers inside a block that must not transfer:\n" + counter.report(), diff --git a/src/cunumpy/algorithms.py b/src/cunumpy/algorithms.py index c98421f..6ffc26c 100644 --- a/src/cunumpy/algorithms.py +++ b/src/cunumpy/algorithms.py @@ -1,18 +1,21 @@ """Array algorithms missing from NumPy/CuPy, on either backend. Morton (Z-order) keys (the same as ``cunumpy/morton.cuh`` computes in a -kernel), a stable sort of several arrays by one key, and sums per key:: +kernel), a stable sort of several arrays by one key, sums per key, and the compaction of +the live rows of particle arrays:: import cunumpy as xp keys = xp.algorithms.morton_keys(positions, lower, upper, levels) keys, order, positions = xp.algorithms.sort_by_key(keys, positions) charge = xp.algorithms.segment_sum(q, cell, n_cells) + n = xp.algorithms.compact_by_mask(alive, positions, charges) """ from cunumpy._algorithms import ( SegmentPlan, cell_offsets, + compact_by_mask, segment_boundaries, segment_sum, sort_by_key, @@ -29,6 +32,7 @@ "MAX_MORTON_LEVELS", "SegmentPlan", "cell_offsets", + "compact_by_mask", "morton_decode", "morton_encode", "morton_keys", diff --git a/src/cunumpy/kernel_testing.py b/src/cunumpy/kernel_testing.py index c26f3f5..1514d8c 100644 --- a/src/cunumpy/kernel_testing.py +++ b/src/cunumpy/kernel_testing.py @@ -49,6 +49,7 @@ def test_parity(name, kernel): from __future__ import annotations +import os import re from collections.abc import Callable, Sequence from typing import Any @@ -67,7 +68,12 @@ def test_parity(name, kernel): _strip_comments, ) from cunumpy._dispatch import Kernel -from cunumpy._emulation import emulate_cuda_kernel, emulation_compiler +from cunumpy._emulation import ( + emulate_cuda_kernel, + emulated_launches, + emulation_compiler, +) +from cunumpy._fake_cupy import host_buffer from cunumpy.xp import cupy_available, get_backend, to_numpy, use_backend # the pytest objects are created on first access, see __getattr__ @@ -76,10 +82,13 @@ def test_parity(name, kernel): "assert_kernels_agree", "backend", # noqa: F822 "check_parity", + "cuda_required", "device_function_kernel", "emulate_cuda_kernel", + "emulated_launches", "emulation_compiler", "fake_cupy_active", + "host_buffer", "install_fake_cupy", "parity_cases", "requires_cupy", # noqa: F822 @@ -109,6 +118,32 @@ def _can_launch() -> bool: return cupy_available() and not fake_cupy_active() +def cuda_required() -> bool: + """Whether ``CUNUMPY_REQUIRE_CUDA`` demands a real GPU (the CI guard). + + Then tests that need a GPU fail instead of being skipped where there is + none (:data:`requires_cupy`, :func:`assert_kernels_agree`, the ``cupy`` + parameter of :func:`backend`), so that a CI job on a GPU machine cannot + pass silently because CuPy or the driver is broken. + """ + return os.environ.get("CUNUMPY_REQUIRE_CUDA", "").lower() in {"1", "true", "yes"} + + +def _skip_or_fail(reason: str) -> None: + """``pytest.skip``, or ``pytest.fail`` if a GPU is required.""" + if cuda_required(): + _pytest().fail(f"CUNUMPY_REQUIRE_CUDA is set, but {reason}", pytrace=False) + _pytest().skip(reason) + + +def cuda_gate() -> bool: + """Condition of ``requires_cupy`` under ``CUNUMPY_REQUIRE_CUDA``: fail, or False.""" + if not _can_launch(): + reason = FAKE_SKIP_REASON if fake_cupy_active() else SKIP_REASON + _pytest().fail(f"CUNUMPY_REQUIRE_CUDA is set, but {reason}", pytrace=False) + return False + + # pytest objects, built on first use so that importing this module does not # import pytest (see __getattr__ below) _LAZY: dict[str, Any] = {} @@ -126,16 +161,28 @@ def _pytest() -> Any: def _build_lazy() -> None: pytest = _pytest() - requires_cupy = pytest.mark.skipif( - not _can_launch(), - reason=FAKE_SKIP_REASON if fake_cupy_active() else SKIP_REASON, - ) + if cuda_required(): + # a string condition is evaluated when the test is set up; it fails the + # test (instead of skipping it) if there is no real GPU + requires_cupy = pytest.mark.skipif( + "__import__('cunumpy.kernel_testing', fromlist=['_']).cuda_gate()", + reason=SKIP_REASON, + ) + else: + requires_cupy = pytest.mark.skipif( + not _can_launch(), + reason=FAKE_SKIP_REASON if fake_cupy_active() else SKIP_REASON, + ) backends = ["numpy", pytest.param("cupy", marks=requires_cupy)] @pytest.fixture(params=backends) def backend(request): - """Run the test once per backend, with that backend active.""" - with use_backend(request.param): + """Run the test once per backend, with that backend active. + + With ``CUNUMPY_REQUIRE_CUDA`` set, the ``cupy`` run fails if the CuPy + backend cannot be activated, instead of silently running on NumPy. + """ + with use_backend(request.param, strict=cuda_required()): yield request.param _LAZY.update(requires_cupy=requires_cupy, BACKENDS=backends, backend=backend) @@ -190,9 +237,43 @@ def _arrays_in(value: Any, name: str, found: dict[str, Any], depth: int) -> None _arrays_in(item, f"{name}.{attr}", found, depth - 1) +def _resolve_output( + entry: Any, + n_args: int, + parameters: Sequence[str] | None, +) -> tuple[int, str | None]: + """The argument index and the field filter (or None) of an `outputs` entry.""" + if isinstance(entry, bool) or not isinstance(entry, (int, str)): + raise TypeError( + "outputs entries must be argument indices (int) or names (str, " + f"optionally 'name.field'), got {entry!r}", + ) + field = None + if isinstance(entry, str): + head, _, field = entry.partition(".") + field = field or None + if head.lstrip("-").isdigit(): + entry = int(head) + elif parameters is not None and head in parameters: + entry = list(parameters).index(head) + else: + known = "unknown" if parameters is None else sorted(parameters) + raise KeyError( + f"output {head!r} is not a parameter of the kernel (parameters: " + f"{known}); give an index if the names are unknown", + ) + index = entry + n_args if entry < 0 else entry + if not 0 <= index < n_args: + raise IndexError( + f"output argument {entry} does not exist: there are {n_args} arguments", + ) + return index, field + + def _collect_arrays( args: Sequence[Any], - outputs: Sequence[int] | None = None, + outputs: Sequence[int | str] | None = None, + parameters: Sequence[str] | None = None, ) -> dict[str, Any]: """The arrays among `args` (or among the arguments `outputs`), by name. @@ -205,22 +286,29 @@ def _collect_arrays( through its struct fields, ``"argument ."``, so that its arrays get the names of the attributes of the host argument object it mirrors, also when the fields are properties. + + An entry of `outputs` is an argument index, or the name of a parameter (in + `parameters`), optionally followed by ``.`` to compare only that + field (attribute) of a struct or argument object, e.g. ``"markers.positions"``. """ - indices = range(len(args)) if outputs is None else outputs + entries = range(len(args)) if outputs is None else outputs found: dict[str, Any] = {} - for entry in indices: - if not isinstance(entry, int) or isinstance(entry, bool): - raise TypeError( - "outputs must be positional argument indices (kernels take " - f"positional arguments only), got {entry!r}", - ) - index = entry + len(args) if entry < 0 else entry - if not 0 <= index < len(args): - raise IndexError( - f"output argument {entry} does not exist: there are {len(args)} " - "arguments", - ) - _arrays_in(args[index], f"argument {index}", found, depth=2) + for entry in entries: + index, field = _resolve_output(entry, len(args), parameters) + local: dict[str, Any] = {} + _arrays_in(args[index], f"argument {index}", local, depth=2) + if field is not None: + prefix = f"argument {index}.{field}" + local = { + name: array + for name, array in local.items() + if name == prefix or name.startswith((prefix + ".", prefix + "[")) + } + if not local: + raise KeyError( + f"output {entry!r}: argument {index} has no array field {field!r}", + ) + found.update(local) return found @@ -264,7 +352,7 @@ def assert_kernels_agree( rtol: float = 1e-12, atol: float = 0.0, n_calls: int = 1, - outputs: Sequence[int] | None = None, + outputs: Sequence[int | str] | None = None, seed: int = 0, ) -> dict[str, np.ndarray]: """Check that the host and CUDA versions of `kernel` compute the same. @@ -301,9 +389,12 @@ def assert_kernels_agree( n_calls : int How many times the kernel is called on each backend (e.g. to test a kernel that accumulates). - outputs : Sequence[int] | None - Indices of the arguments to compare (negative indices count from the - end), like ``PyccelKernel(outputs=...)``. By default the ``outputs`` + outputs : Sequence[int | str] | None + The arguments to compare, like ``PyccelKernel(outputs=...)``: indices + (negative indices count from the end) or parameter names. A name with + ``.`` (``"markers.positions"``) compares only that field of a + struct or argument object, leaving the other fields (e.g. buffers the + two kernels fill differently) out. By default the ``outputs`` declared by the host kernel are used, and if it declares none, every argument. An argument that is an array is compared; for a tuple, list, dict or object argument (e.g. a ``CudaArguments`` object), the arrays @@ -328,7 +419,7 @@ def assert_kernels_agree( Notes ----- The test is skipped with ``pytest.skip`` if CuPy or a GPU is not available, - or if the fake CuPy is active. + or if the fake CuPy is active; with ``CUNUMPY_REQUIRE_CUDA=1`` it fails instead. """ if not isinstance(kernel, Kernel): raise TypeError(f"expected a Kernel, got {type(kernel).__name__}") @@ -341,10 +432,11 @@ def assert_kernels_agree( if outputs is None: outputs = kernel.host_kernel.outputs if fake_cupy_active(): - _pytest().skip(FAKE_SKIP_REASON) + _skip_or_fail(FAKE_SKIP_REASON) if not cupy_available(): - _pytest().skip(SKIP_REASON) + _skip_or_fail(SKIP_REASON) + parameters = kernel.host_parameters() results = {} for backend in ("numpy", "cupy"): with use_backend(backend): @@ -354,7 +446,7 @@ def assert_kernels_agree( launch = n_threads(args) if callable(n_threads) else n_threads for _ in range(n_calls): kernel(*args, n_threads=launch, grid=grid, block=block) - results[backend] = _collect_arrays(args, outputs) + results[backend] = _collect_arrays(args, outputs, parameters) host = {name: to_numpy(a) for name, a in results["numpy"].items()} _compare_results(host, results["cupy"], rtol, atol, kernel.name) diff --git a/src/cunumpy/kernels.py b/src/cunumpy/kernels.py index 46de0a0..30aef3f 100644 --- a/src/cunumpy/kernels.py +++ b/src/cunumpy/kernels.py @@ -40,6 +40,7 @@ get_device_kernel_implementation, get_host_kernel_implementation, kernel_output, + outputs_from_annotations, set_device_kernel_implementation, set_host_kernel_implementation, use_device_kernel_implementation, @@ -64,6 +65,7 @@ "get_host_kernel_implementation", "kernel_output", "metal_available", + "outputs_from_annotations", "set_device_kernel_implementation", "set_host_kernel_implementation", "use_device_kernel_implementation", diff --git a/src/cunumpy/xp.py b/src/cunumpy/xp.py index 93f0d48..d667280 100644 --- a/src/cunumpy/xp.py +++ b/src/cunumpy/xp.py @@ -20,6 +20,7 @@ _is_device_copy, _nbytes, _record, + _record_sync, ) if os.environ.get("CUNUMPY_FAKE_CUPY", "").strip().lower() in ("1", "true", "yes"): @@ -248,6 +249,8 @@ def default_float_dtype() -> Any: def synchronize() -> None: """Wait for all kernels in all streams on current device to complete.""" if array_backend.backend == "cupy": + if _COUNTERS: + _record_sync("synchronize()") try: import cupy as cp diff --git a/tests/unit/test_compact_by_mask.py b/tests/unit/test_compact_by_mask.py new file mode 100644 index 0000000..39014b9 --- /dev/null +++ b/tests/unit/test_compact_by_mask.py @@ -0,0 +1,43 @@ +"""compact_by_mask keeps the masked rows, in order, at the front of every array.""" + +import numpy as np +import pytest + +import cunumpy as xp + + +def test_rows_are_moved_to_the_front_in_order_for_every_array(): + markers = np.arange(12.0).reshape(4, 3) + weights = np.array([1.0, 2.0, 3.0, 4.0]) + ids = np.array([10, 11, 12, 13]) + alive = np.array([True, False, True, True]) + n = xp.algorithms.compact_by_mask(alive, markers, weights, ids) + assert n == 3 + np.testing.assert_array_equal(markers[:n], [[0, 1, 2], [6, 7, 8], [9, 10, 11]]) + np.testing.assert_array_equal(weights[:n], [1.0, 3.0, 4.0]) + np.testing.assert_array_equal(ids[:n], [10, 12, 13]) + + +def test_all_true_none_true_and_empty(): + a = np.arange(4) + assert xp.algorithms.compact_by_mask(np.ones(4, bool), a) == 4 + np.testing.assert_array_equal(a, np.arange(4)) + assert xp.algorithms.compact_by_mask(np.zeros(4, bool), a) == 0 + assert xp.algorithms.compact_by_mask(np.zeros(0, bool), np.zeros(0)) == 0 + + +def test_matches_boolean_indexing_on_random_masks(): + rng = np.random.default_rng(1) + for _ in range(20): + a = rng.random((50, 2)) + mask = rng.random(50) < 0.4 + expected = a[mask] + n = xp.algorithms.compact_by_mask(mask, a) + np.testing.assert_array_equal(a[:n], expected) + + +def test_argument_errors(): + with pytest.raises(TypeError, match="boolean"): + xp.algorithms.compact_by_mask(np.array([1, 0]), np.zeros(2)) + with pytest.raises(ValueError, match="2 rows"): + xp.algorithms.compact_by_mask(np.array([True, False]), np.zeros(3)) diff --git a/tests/unit/test_emulated_launches.py b/tests/unit/test_emulated_launches.py new file mode 100644 index 0000000..31aff8c --- /dev/null +++ b/tests/unit/test_emulated_launches.py @@ -0,0 +1,159 @@ +"""Struct parameters in the emulation, and `emulated_launches` on the fake CuPy.""" + +import os +import subprocess +import sys +from pathlib import Path + +import numpy as np +import pytest + +from cunumpy.arguments import CudaStruct +from cunumpy.kernel_testing import emulate_cuda_kernel, emulation_compiler +from cunumpy.kernels import CudaKernel + +pytestmark = pytest.mark.skipif( + emulation_compiler() is None, + reason="no C++ compiler for the emulation", +) + +PARTICLES = CudaStruct( + "Particles", + [("markers", "Array2D"), ("alive", "bool*"), ("n", "int")], +) +SOURCE = ( + '#include "cunumpy/array_view.cuh"\n' + + PARTICLES.declaration + + r""" +extern "C" __global__ +void push(Particles p, double dt, double* total) { + int i = blockDim.x * blockIdx.x + threadIdx.x; + if (i < p.n && p.alive[i]) { + p.markers(i, 0) += dt * p.markers(i, 1); + total[0] += p.markers(i, 0); + } +} +""" +) + + +def particles(): + markers = np.arange(12.0).reshape(4, 3) + return markers, np.array([True, True, False, True]) + + +def expected(markers, alive, dt): + out = markers.copy() + out[alive, 0] += dt * out[alive, 1] + return out + + +@pytest.mark.parametrize("kind", ["mapping", "object"]) +def test_struct_given_as_mapping_or_object(kind): + markers, alive = particles() + want = expected(markers, alive, 0.5) + total = np.zeros(1) + fields = {"markers": markers, "alive": alive, "n": 4} + value = fields if kind == "mapping" else type("Args", (), fields)() + push = CudaKernel(SOURCE, "push", structs=[PARTICLES]) + emulate_cuda_kernel(push, value, 0.5, total, n_threads=4) + np.testing.assert_allclose(markers, want) + assert total[0] == pytest.approx(want[alive, 0].sum()) + + +def test_struct_launch_shape_is_inferred_from_the_first_array(): + markers, alive = particles() + push = CudaKernel(SOURCE, "push", structs=[PARTICLES]) + emulate_cuda_kernel( + push, {"markers": markers, "alive": alive, "n": 4}, 1.0, np.zeros(1) + ) + np.testing.assert_allclose(markers, expected(*particles(), 1.0)) + + +def test_struct_with_a_missing_field_is_a_type_error(): + markers, alive = particles() + push = CudaKernel(SOURCE, "push", structs=[PARTICLES]) + with pytest.raises(TypeError, match="no value for the field 'n'"): + emulate_cuda_kernel( + push, {"markers": markers, "alive": alive}, 1.0, np.zeros(1), n_threads=4 + ) + + +SCRIPT = r""" +import numpy as np +import cunumpy as xp +from cunumpy.arguments import CudaStruct, CudaStructArguments +from cunumpy.kernel_testing import emulated_launches, host_buffer +from cunumpy.kernels import CudaKernel, Kernel + +Particles = CudaStruct("Particles", [("x", "double*"), ("n", "int")]) +source = Particles.declaration + ''' +extern "C" __global__ void scale(Particles p, double f) { + int i = blockDim.x * blockIdx.x + threadIdx.x; + if (i < p.n) p.x[i] *= f; +}''' +scale = CudaKernel(source, "scale", structs=[Particles]) + + +class Args(CudaStructArguments): + struct_name = "Particles" + fields = (("x", "double*"), ("n", "int")) + + def __init__(self, x): + self.x = x + self.n = x.shape[0] + self.pack() + + +x = xp.asarray(np.arange(5.0)) +assert xp.get_backend() == "cupy" +try: + scale(Particles(x=x, n=5), 2.0, n_threads=5) +except NotImplementedError: + pass +else: + raise AssertionError("the fake CuPy must refuse a launch outside the block") + +with emulated_launches(): + scale(Particles(x=x, n=5), 2.0, n_threads=5) + scale(Args(x), 3.0, n_threads=5) + host_view = host_buffer(x) +np.testing.assert_allclose(host_buffer(x), np.arange(5.0) * 6.0) +assert host_view is host_buffer(x) # not a copy + +# a Kernel dispatches to the CUDA kernel on the CuPy backend +def host_scale(args, f): + raise AssertionError("the CUDA kernel must run") + +kernel = Kernel(host_scale, scale) +with emulated_launches(): + kernel(Args(x), 0.5, n_threads=5) +np.testing.assert_allclose(host_buffer(x), np.arange(5.0) * 3.0) + +try: + host_buffer(np.zeros(2)) +except TypeError: + pass +else: + raise AssertionError("host_buffer must refuse a NumPy array") +print("emulated launches OK") +""" + + +def test_emulated_launches_on_the_fake_cupy(): + root = Path(__file__).resolve().parents[2] + env = dict(os.environ, CUNUMPY_FAKE_CUPY="1", CUNUMPY_BACKEND="cupy") + env["PYTHONPATH"] = os.pathsep.join( + p for p in (str(root / "src"), env.get("PYTHONPATH", "")) if p + ) + env.pop("CUNUMPY_CUDA_DEBUG", None) + result = subprocess.run( + [sys.executable, "-c", SCRIPT], + env=env, + capture_output=True, + text=True, + check=False, + cwd=str(root), + ) + assert result.returncode == 0, result.stdout + result.stderr + assert "emulated launches OK" in result.stdout diff --git a/tests/unit/test_kernel_outputs.py b/tests/unit/test_kernel_outputs.py new file mode 100644 index 0000000..358f0ac --- /dev/null +++ b/tests/unit/test_kernel_outputs.py @@ -0,0 +1,156 @@ +"""Kernel outputs by name, from annotations and from a kernel folder; struct fields in tests.""" + +import importlib +import sys +import textwrap +from typing import Final + +import numpy as np +import pytest + +from cunumpy.kernel_testing import _collect_arrays +from cunumpy.kernels import Kernel, PyccelKernel, outputs_from_annotations + + +def solve(a, b, out): + out[:] = a + b + + +def host_arrays(kernel, args, kwargs=None): + """The ids of the arrays the declared outputs of `kernel` reach.""" + return kernel._output_host_arrays(list(args), dict(kwargs or {})) + + +def test_output_name_finds_a_positional_argument(): + a, b, out = np.zeros(2), np.zeros(2), np.zeros(2) + kernel = PyccelKernel(solve, outputs=("out",)) + assert host_arrays(kernel, (a, b, out)) == {id(out)} + assert host_arrays(kernel, (a, b), {"out": out}) == {id(out)} + + +def test_output_index_finds_a_keyword_argument(): + a, b, out = np.zeros(2), np.zeros(2), np.zeros(2) + kernel = PyccelKernel(solve, outputs=(2,)) + assert host_arrays(kernel, (a, b, out)) == {id(out)} + assert host_arrays(kernel, (a, b), {"out": out}) == {id(out)} + + +def test_unknown_parameter_names_keep_the_old_rules(): + def hidden(*args): # no named parameters, like a compiled kernel + pass + + a, out = np.zeros(2), np.zeros(2) + with pytest.raises(KeyError, match="declared by index"): + host_arrays(PyccelKernel(hidden, outputs=("out",)), (a, out)) + with pytest.raises(IndexError, match="declared by name"): + host_arrays(PyccelKernel(hidden, outputs=(1,)), (a,), {"out": out}) + named = PyccelKernel(hidden, outputs=("out",), parameters=["a", "out"]) + assert host_arrays(named, (a, out)) == {id(out)} + + +def test_unknown_name_is_reported(): + kernel = PyccelKernel(solve, outputs=("missing",)) + with pytest.raises(KeyError, match="no argument of that name"): + host_arrays(kernel, (np.zeros(1),) * 3) + + +def test_outputs_from_annotations(): + def push( + markers: "float[:, :]", + weights: "Final[float[:]]", + dt: "float", + n: int, + args: "MarkerArgs", # noqa: F821 + scratch: Final[np.ndarray], + ): + pass + + assert outputs_from_annotations(push) == ("markers", "args") + + def unannotated(a, b): + pass + + assert outputs_from_annotations(unannotated) is None + assert outputs_from_annotations(len) is None # no signature to read + + +@pytest.fixture +def folder_kernel(tmp_path, monkeypatch): + folder = tmp_path / "output_pkg" / "scale" + folder.mkdir(parents=True) + (tmp_path / "output_pkg" / "__init__.py").write_text("") + (folder / "__init__.py").write_text("") + (folder / "scale_kernels.py").write_text( + textwrap.dedent(""" + def scale(x: "float[:]", factor: "Final[float[:]]", a: "float", n: "int"): + for i in range(n): + x[i] *= factor[i] * a + """), + ) + monkeypatch.syspath_prepend(str(tmp_path)) + yield "output_pkg.scale" + for module in [m for m in sys.modules if m.startswith("output_pkg")]: + del sys.modules[module] + importlib.invalidate_caches() + + +def test_from_folder_reads_outputs_from_annotations(folder_kernel): + kernel = Kernel.from_folder(folder_kernel, outputs="annotations") + assert kernel.host_kernel.outputs == ("x",) + x, factor = np.ones(3), np.full(3, 2.0) + assert host_arrays(kernel.host_kernel, (x, factor, 3.0, 3)) == {id(x)} + + +def test_from_folder_outputs_by_name_use_the_host_parameters(folder_kernel): + kernel = Kernel.from_folder(folder_kernel, outputs=("x",)) + x, factor = np.ones(3), np.full(3, 2.0) + assert host_arrays(kernel.host_kernel, (x, factor, 3.0, 3)) == {id(x)} + assert host_arrays(kernel.host_kernel, (), {"x": x}) == {id(x)} + + +def test_host_options_take_precedence_over_outputs(folder_kernel): + kernel = Kernel.from_folder( + folder_kernel, outputs="annotations", host_options={"outputs": ("factor",)} + ) + assert kernel.host_kernel.outputs == ("factor",) + + +def test_from_folder_rejects_an_unknown_outputs_string(folder_kernel): + with pytest.raises(ValueError, match="annotations"): + Kernel.from_folder(folder_kernel, outputs="all") + + +# --- comparing only some fields of a struct argument ------------------------- + + +class Markers: + def __init__(self): + self.positions = np.zeros((3, 2)) + self.buffer = np.ones(5) + self.n = 3 + + +def test_collect_arrays_by_parameter_name_and_field(): + markers, weights = Markers(), np.zeros(3) + args = (markers, weights) + names = ["markers", "weights"] + + both = _collect_arrays(args, ("markers",), names) + assert set(both) == {"argument 0.positions", "argument 0.buffer"} + + only = _collect_arrays(args, ("markers.positions", "weights"), names) + assert set(only) == {"argument 0.positions", "argument 1"} + assert only["argument 0.positions"] is markers.positions + + assert set(_collect_arrays(args, ("0.buffer",), names)) == {"argument 0.buffer"} + assert set(_collect_arrays(args, (-1,), names)) == {"argument 1"} + + +def test_collect_arrays_field_and_name_errors(): + args = (Markers(),) + with pytest.raises(KeyError, match="no array field 'velocities'"): + _collect_arrays(args, ("markers.velocities",), ["markers"]) + with pytest.raises(KeyError, match="not a parameter"): + _collect_arrays(args, ("particles",), ["markers"]) + with pytest.raises(KeyError, match="names are unknown|unknown"): + _collect_arrays(args, ("markers",), None) diff --git a/tests/unit/test_kernel_testing.py b/tests/unit/test_kernel_testing.py index ed01e9c..dceb1f7 100644 --- a/tests/unit/test_kernel_testing.py +++ b/tests/unit/test_kernel_testing.py @@ -127,8 +127,10 @@ def __init__(self, x, values): with pytest.raises(IndexError, match="output argument 6 does not exist"): _collect_arrays(args, outputs=(6,)) - with pytest.raises(TypeError, match="positional argument indices"): + with pytest.raises(KeyError, match="not a parameter of the kernel"): _collect_arrays(args, outputs=("out",)) + with pytest.raises(TypeError, match="indices \\(int\\) or names"): + _collect_arrays(args, outputs=(1.5,)) def test_compare_results(): diff --git a/tests/unit/test_porting_helpers.py b/tests/unit/test_porting_helpers.py index c6ac552..5c7ccb0 100644 --- a/tests/unit/test_porting_helpers.py +++ b/tests/unit/test_porting_helpers.py @@ -532,7 +532,10 @@ def test_require_version(monkeypatch): assert isinstance(buf, np.ndarray) and buf.tolist() == [1.0, 2.0] buf[:] = [5.0, 6.0] assert xp.to_numpy(d).tolist() == [5.0, 6.0] -assert sorted(e.kind for e in counter.events) == ["to_device", "to_host"] +assert sorted(e.kind for e in counter.events if e.kind != "sync") == [ + "to_device", + "to_host", +] with xp.mpi.mpi_buffer(d, cuda_aware=True) as buf: assert buf is d producer = xp.cuda.create_stream() diff --git a/tests/unit/test_pyccel_kernel.py b/tests/unit/test_pyccel_kernel.py index 43113b8..36492ed 100644 --- a/tests/unit/test_pyccel_kernel.py +++ b/tests/unit/test_pyccel_kernel.py @@ -423,7 +423,7 @@ def test_misdeclared_output_index_raises(): def test_misdeclared_output_name_raises(): wrapped = PyccelKernel(lambda out: None, use_cupy=True, outputs=("nope",)) - with pytest.raises(KeyError, match="no such keyword argument"): + with pytest.raises(KeyError, match="no argument of that name"): wrapped(np.zeros(2)) diff --git a/tests/unit/test_require_cuda.py b/tests/unit/test_require_cuda.py new file mode 100644 index 0000000..8325267 --- /dev/null +++ b/tests/unit/test_require_cuda.py @@ -0,0 +1,66 @@ +"""With CUNUMPY_REQUIRE_CUDA=1 the GPU markers of kernel_testing fail (as errors) instead of skipping.""" + +import os +import subprocess +import sys +from pathlib import Path + +import pytest + +import cunumpy as xp + +pytestmark = pytest.mark.skipif( + xp.cupy_available(), reason="checks the behavior without a GPU" +) + +TESTS = """ +import pytest +from cunumpy.kernel_testing import backend, requires_cupy + + +@requires_cupy +def test_marked(): + pass + + +def test_backend(backend): + pass +""" + + +def run(tmp_path, require): + (tmp_path / "test_gpu.py").write_text(TESTS) + root = Path(__file__).resolve().parents[2] + env = dict(os.environ) + env["PYTHONPATH"] = os.pathsep.join( + p for p in (str(root / "src"), env.get("PYTHONPATH", "")) if p + ) + env.pop("CUNUMPY_REQUIRE_CUDA", None) + env.pop("CUNUMPY_FAKE_CUPY", None) + env.pop("CUNUMPY_BACKEND", None) + if require: + env["CUNUMPY_REQUIRE_CUDA"] = "1" + return subprocess.run( + [sys.executable, "-m", "pytest", "-p", "no:cacheprovider", "-rA", "-q"], + cwd=tmp_path, + env=env, + capture_output=True, + text=True, + check=False, + ) + + +def test_markers_skip_without_a_gpu(tmp_path): + result = run(tmp_path, require=False) + assert result.returncode == 0, result.stdout + assert "SKIPPED" in result.stdout and "ERROR" not in result.stdout + + +def test_markers_error_when_a_gpu_is_required(tmp_path): + result = run(tmp_path, require=True) + assert result.returncode != 0, result.stdout + assert "CUNUMPY_REQUIRE_CUDA is set" in result.stdout + # the condition of the marker raises while the test is set up: an error + assert "ERROR test_gpu.py::test_marked" in result.stdout + assert "ERROR test_gpu.py::test_backend[cupy]" in result.stdout + assert "PASSED test_gpu.py::test_backend[numpy]" in result.stdout diff --git a/tests/unit/test_sync_counting.py b/tests/unit/test_sync_counting.py new file mode 100644 index 0000000..d0db679 --- /dev/null +++ b/tests/unit/test_sync_counting.py @@ -0,0 +1,76 @@ +"""count_transfers counts syncs: synchronize(), and scalar reads on the fake CuPy.""" + +import os +import subprocess +import sys +from pathlib import Path + +import pytest + +import cunumpy as xp +from cunumpy._transfers import TransferCounter, _record_sync + + +def test_sync_events_are_listed_but_not_in_the_total(): + with xp.profiling.count_transfers() as counter: + _record_sync("something") + assert counter.syncs == 1 and counter.total == 0 + assert "sync (1)" in counter.report() + + +def test_assert_no_transfers_accepts_syncs_unless_asked(): + with xp.profiling.assert_no_transfers(): + _record_sync("something") + with ( + pytest.raises(AssertionError, match="sync"), + xp.profiling.assert_no_transfers(syncs=True), + ): + _record_sync("something") + + +def test_numpy_backend_synchronize_is_not_a_sync(): + with xp.profiling.count_transfers() as counter: + xp.synchronize() + assert counter.syncs == 0 + assert isinstance(counter, TransferCounter) + + +SCRIPT = r""" +import cunumpy as xp + +a = xp.asarray([1.0, 2.0, 3.0]) +with xp.profiling.count_transfers() as counter: + float(a[0]); int(a[1]); bool(a[2]); a.sum().item(); a.tolist() + xp.synchronize() + b = a * 2 # device work: no sync +assert counter.syncs == 6, counter.report() +assert counter.total == 0 +with xp.profiling.assert_no_transfers(): + float(a[0]) +try: + with xp.profiling.assert_no_transfers(syncs=True): + float(a[0]) +except AssertionError: + pass +else: + raise AssertionError("syncs=True must reject a scalar read") +print("sync counting OK") +""" + + +def test_scalar_reads_and_synchronize_are_counted_on_the_fake_cupy(): + root = Path(__file__).resolve().parents[2] + env = dict(os.environ, CUNUMPY_FAKE_CUPY="1", CUNUMPY_BACKEND="cupy") + env["PYTHONPATH"] = os.pathsep.join( + p for p in (str(root / "src"), env.get("PYTHONPATH", "")) if p + ) + result = subprocess.run( + [sys.executable, "-c", SCRIPT], + env=env, + capture_output=True, + text=True, + check=False, + cwd=str(root), + ) + assert result.returncode == 0, result.stdout + result.stderr + assert "sync counting OK" in result.stdout diff --git a/tests/unit/test_transfers.py b/tests/unit/test_transfers.py index 773cc11..90201bf 100644 --- a/tests/unit/test_transfers.py +++ b/tests/unit/test_transfers.py @@ -84,7 +84,7 @@ def test_empty_counter(): assert counter.events == [] and counter.kernel_conversion_calls == [] assert counter.report().startswith("0 transfer(s) through cunumpy") assert repr(counter) == ( - "TransferCounter(to_host=0, to_device=0, kernel_conversion=0, fallback=0, device_copy=0)" + "TransferCounter(to_host=0, to_device=0, kernel_conversion=0, fallback=0, device_copy=0, sync=0)" ) @@ -189,7 +189,7 @@ def test_report_groups_events_by_kind_and_call_site(fake_device): lines = report.splitlines() assert lines[0] == ( "4 transfer(s) through cunumpy " - "(3 to_host, 1 to_device, 0 kernel_conversion, 0 fallback, 0 device_copy)" + "(3 to_host, 1 to_device, 0 kernel_conversion, 0 fallback, 0 device_copy, 0 sync)" ) assert " to_host (3):" in lines assert " to_device (1):" in lines From c6ce9d02bc991649a73c023b93342a7444a69301 Mon Sep 17 00:00:00 2001 From: Max Date: Thu, 8 Oct 2026 00:08:50 +0200 Subject: [PATCH 2/2] Fix gpu test --- src/cunumpy/kernel_testing.py | 5 +++-- tests/unit/test_porting_helpers.py | 1 + 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/src/cunumpy/kernel_testing.py b/src/cunumpy/kernel_testing.py index 1514d8c..0b269ca 100644 --- a/src/cunumpy/kernel_testing.py +++ b/src/cunumpy/kernel_testing.py @@ -161,9 +161,10 @@ def _pytest() -> Any: def _build_lazy() -> None: pytest = _pytest() - if cuda_required(): + if cuda_required() and not _can_launch(): # a string condition is evaluated when the test is set up; it fails the - # test (instead of skipping it) if there is no real GPU + # test (instead of skipping it) because there is no real GPU. With a + # working GPU the marker is the ordinary one below. requires_cupy = pytest.mark.skipif( "__import__('cunumpy.kernel_testing', fromlist=['_']).cuda_gate()", reason=SKIP_REASON, diff --git a/tests/unit/test_porting_helpers.py b/tests/unit/test_porting_helpers.py index 5c7ccb0..baaef36 100644 --- a/tests/unit/test_porting_helpers.py +++ b/tests/unit/test_porting_helpers.py @@ -553,6 +553,7 @@ def test_require_version(monkeypatch): def test_fake_cupy_in_subprocess(): root = Path(__file__).resolve().parents[2] env = dict(os.environ, CUNUMPY_FAKE_CUPY="1", CUNUMPY_BACKEND="cupy") + env.pop("CUNUMPY_REQUIRE_CUDA", None) # the fake CuPy skips by design env["PYTHONPATH"] = os.pathsep.join( p for p in (str(root / "src"), env.get("PYTHONPATH", "")) if p )