From 37c3d6816382ad017c922f672088efa0566bd6b6 Mon Sep 17 00:00:00 2001
From: Max <max.lindqvist@ipp.mpg.de>
Date: Tue, 27 Oct 2026 16:33:56 +0100
Subject: [PATCH] Run ssort, ruff and numpydoc as pre-commit hooks

---
 .pre-commit-config.yaml        |  19 +
 src/cunumpy/_algorithms.py     |  30 +-
 src/cunumpy/_cuda_kernel.py    | 634 ++++++++++++++++-----------------
 src/cunumpy/_dispatch.py       |  98 ++---
 src/cunumpy/_fake_cupy.py      |  20 +-
 src/cunumpy/_fake_cupy_impl.py | 146 ++++----
 src/cunumpy/_fake_cupyx.py     |  72 ++--
 src/cunumpy/_kernel.py         | 116 +++---
 src/cunumpy/_metal_kernel.py   |  24 +-
 src/cunumpy/_mirror.py         |  48 +--
 src/cunumpy/_pyccel_cache.py   |   6 +-
 src/cunumpy/_random_streams.py |  16 +-
 src/cunumpy/_scipy_backend.py  |  50 +--
 src/cunumpy/_staging.py        |  12 +-
 src/cunumpy/_streams.py        |  12 +-
 src/cunumpy/_transfers.py      |  33 +-
 src/cunumpy/kernel_testing.py  |  20 +-
 src/cunumpy/xp.py              | 112 +++---
 testing/smoke_wheel.py         |   1 +
 tests/unit/test_cuda_kernel.py |  88 ++---
 tests/unit/test_mpi_serial.py  |   6 +-
 tests/unit/test_staging.py     |  12 +-
 tests/unit/test_transfers.py   |  14 +-
 23 files changed, 806 insertions(+), 783 deletions(-)

diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index af67085..6125aeb 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -13,3 +13,22 @@ repos:
     rev: 0.8.1
     hooks:
       - id: nbstripout # remove jupyter notebook cell output
+
+  - repo: https://github.com/bwhmather/ssort
+    rev: v0.12.0
+    hooks:
+      - id: ssort # sort statements so definitions come before their use
+
+  - repo: https://github.com/astral-sh/ruff-pre-commit
+    rev: v0.17.3
+    hooks:
+      - id: ruff-check # lint, including import sorting
+        args: ["--fix"]
+      - id: ruff-format
+
+  - repo: https://github.com/numpy/numpydoc
+    rev: v1.11.0
+    hooks:
+      - id: numpydoc-validation # docstring sections, configured in pyproject.toml
+        files: ^src/cunumpy/
+        exclude: _fake_cupy
diff --git a/src/cunumpy/_algorithms.py b/src/cunumpy/_algorithms.py
index 76355df..15876e7 100644
--- a/src/cunumpy/_algorithms.py
+++ b/src/cunumpy/_algorithms.py
@@ -374,6 +374,21 @@ def _radix_argsort(keys: np.ndarray) -> np.ndarray:
     return order.astype(np.int64, copy=False)
 
 
+def _take_along_axis(xpm: Any, array: Any, order: Any, axis: int) -> Any:
+    """Return `array` taken at `order` along `axis`, as a new C-contiguous array.
+
+    A component-major ``(ncomp, n)`` NumPy array is gathered one contiguous row at a
+    time, which is about twice as fast as ``take`` along its last axis; everything
+    else, and CuPy, goes through ``take``.
+    """
+    if xpm is np and array.ndim == 2 and axis == 1:
+        gathered = np.empty((array.shape[0], order.shape[0]), dtype=array.dtype)
+        for row in range(array.shape[0]):
+            gathered[row] = array[row][order]
+        return gathered
+    return xpm.take(array, order, axis=axis)
+
+
 @sp.profile("cunumpy.sort_by_key", tags=_PROFILE_TAGS)
 def sort_by_key(keys: Any, *arrays: Any, axis: int = 0) -> tuple[Any, ...]:
     """Sort `keys` and reorder every array the same way, in one stable argsort.
@@ -449,21 +464,6 @@ def sort_by_key(keys: Any, *arrays: Any, axis: int = 0) -> tuple[Any, ...]:
     )
 
 
-def _take_along_axis(xpm: Any, array: Any, order: Any, axis: int) -> Any:
-    """Return `array` taken at `order` along `axis`, as a new C-contiguous array.
-
-    A component-major ``(ncomp, n)`` NumPy array is gathered one contiguous row at a
-    time, which is about twice as fast as ``take`` along its last axis; everything
-    else, and CuPy, goes through ``take``.
-    """
-    if xpm is np and array.ndim == 2 and axis == 1:
-        gathered = np.empty((array.shape[0], order.shape[0]), dtype=array.dtype)
-        for row in range(array.shape[0]):
-            gathered[row] = array[row][order]
-        return gathered
-    return xpm.take(array, order, axis=axis)
-
-
 @sp.profile("cunumpy.compact_by_mask", tags=_PROFILE_TAGS)
 def compact_by_mask(mask: Any, *arrays: Any, axis: int = 0) -> int:
     """Move the entries where `mask` is True to the front of every array, in place.
diff --git a/src/cunumpy/_cuda_kernel.py b/src/cunumpy/_cuda_kernel.py
index 3e72f02..e6ae384 100644
--- a/src/cunumpy/_cuda_kernel.py
+++ b/src/cunumpy/_cuda_kernel.py
@@ -584,98 +584,6 @@ def _template_arg(value: Any) -> str:
     return ctype_of(value)
 
 
-def parse_cuda_signature(
-    source: str,
-    name: str,
-    *,
-    structs: Iterable[CudaStruct] = (),
-    template_args: Sequence[Any] | None = None,
-) -> tuple[CudaParameter, ...]:
-    """Parse the parameters of a ``__global__`` function.
-
-    C types map to dtypes as on Linux (LP64): ``int`` is int32, ``long`` and
-    ``long long`` are int64.
-
-    Parameters
-    ----------
-    source : str
-        CUDA C source code.
-    name : str
-        Name of the ``__global__`` function.
-    structs : iterable of CudaStruct, optional
-        Struct types that may appear as parameters (by value). A definition
-        of the struct in `source` must match.
-    template_args : sequence, optional
-        Template arguments if `name` is a function template: C types or
-        dtypes (see :func:`~cunumpy.cuda.ctype_of`) for type parameters,
-        ints or bools for non-type parameters.
-
-    Returns
-    -------
-    tuple of CudaParameter
-        The parameters, in order.
-
-    Raises
-    ------
-    ValueError
-        If there is no such function, `template_args` do not fit, a struct
-        definition does not match its :class:`~cunumpy.arguments.CudaStruct`,
-        or a parameter type cannot be checked (e.g. a macro or ``double**``).
-
-    Examples
-    --------
-    >>> source = '''
-    ... template <typename T>
-    ... __global__ void scale(T* x, T factor, int n) {}'''
-    >>> [(p.name, p.ctype, p.pointer) for p in xp.cuda.parse_cuda_signature(
-    ...     source, "scale", template_args=[np.float32])]
-    [('x', 'float', True), ('factor', 'float', False), ('n', 'int', False)]
-    """
-    code = _strip_comments(source)
-    structs = {s.name: s for s in structs}
-    for struct in structs.values():
-        struct.check_source(code)
-
-    pattern = r"(?:template\s*<(?P<template>[^{};]*?)>\s*)?__global__\s+void\s+"
-    match = re.search(pattern + re.escape(name) + r"\s*\(", code)
-    if match is None:
-        raise ValueError(f"no __global__ function {name!r} found in the CUDA source")
-
-    depth, start = 1, match.end()
-    for pos in range(start, len(code)):
-        if code[pos] == "(":
-            depth += 1
-        elif code[pos] == ")":
-            depth -= 1
-            if depth == 0:
-                break
-    else:
-        raise ValueError(f"unbalanced parentheses in the signature of {name!r}")
-    params = code[start:pos].strip()
-
-    template = match.group("template")
-    if template is not None:
-        template_params = [p.split() for p in _split_top_level(template) if p.strip()]
-        if template_args is None or len(template_args) != len(template_params):
-            raise ValueError(
-                f"{name!r} is a template with {len(template_params)} parameters; "
-                f"pass them as template_args",
-            )
-        for words, value in zip(template_params, template_args):
-            if words[0] in ("typename", "class"):
-                params = re.sub(
-                    r"\b" + re.escape(words[-1]) + r"\b",
-                    _template_arg(value),
-                    params,
-                )
-    elif template_args:
-        raise ValueError(f"{name!r} is not a template, but template_args were given")
-
-    if params in ("", "void"):
-        return ()
-    return tuple(_parse_parameter(p, structs) for p in _split_top_level(params))
-
-
 def _describe(param: CudaParameter, index: int) -> str:
     ctype = param.ctype + ("*" if param.pointer else "")
     return f"argument {index} ({ctype} {param.name})"
@@ -860,17 +768,6 @@ 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)
@@ -996,60 +893,6 @@ def _header_source(
     return "\n".join(lines)
 
 
-def write_cuda_header(
-    path: str | Path,
-    structs: Iterable[CudaStruct],
-    guard: str | None = None,
-    *,
-    includes: Iterable[str] = (),
-) -> str:
-    """Write the definitions of several structs to one header file.
-
-    The header has an include guard, ``#include "cunumpy/array_view.cuh"`` if
-    a struct has array view fields, any `includes`, then the structs in
-    order. Commit it next to the kernels and let a test regenerate and
-    compare it (see :meth:`CudaStruct.to_header`).
-
-    Parameters
-    ----------
-    path : str or Path
-        File to write.
-    structs : iterable of CudaStruct
-        The structs, in declaration order.
-    guard : str, optional
-        Include guard macro; by default from the file name
-        (``pusher_args.cuh`` -> ``PUSHER_ARGS_CUH``).
-    includes : iterable of str, optional
-        Additional headers, as file names or ``#include`` lines.
-
-    Returns
-    -------
-    str
-        The header source that was written.
-
-    Examples
-    --------
-    >>> import os, tempfile
-    >>> Vec = xp.arguments.CudaStruct("Vec", [("data", "double*"), ("n", "int")])
-    >>> path = os.path.join(tempfile.mkdtemp(), "vec.cuh")
-    >>> print(xp.arguments.write_cuda_header(path, [Vec]), end="")
-    // Generated by cunumpy.arguments.CudaStruct from the Python definition; do not edit.
-    #ifndef VEC_CUH
-    #define VEC_CUH
-    <BLANKLINE>
-    struct Vec {
-        double* data;
-        int n;
-    };
-    <BLANKLINE>
-    #endif  // VEC_CUH
-    """
-    path = Path(path)
-    source = _header_source(tuple(structs), guard or _header_guard(path.name), includes)
-    path.write_text(source)
-    return source
-
-
 def _read_python_source(source: str | Path) -> tuple[str, str]:
     """Return ``(text, description)`` of a Python source given as a path or as code."""
     if isinstance(source, Path) or (
@@ -1325,9 +1168,6 @@ def from_pyccel_class(
         fields = _apply_contiguous(fields, contiguous, f"{class_name} in {where}")
         return cls(class_name if name is None else name, fields)
 
-    def __repr__(self) -> str:
-        return f"CudaStruct({self._name!r}, {len(self._fields)} fields)"
-
     @property
     def name(self) -> str:
         """Name of the struct type in C."""
@@ -1592,6 +1432,166 @@ def __call__(self, **values: Any) -> CudaStructValue:
             packed[field.name] = value.data.ptr if field.pointer else value
         return CudaStructValue(self, packed[()], values)
 
+    def __repr__(self) -> str:
+        return f"CudaStruct({self._name!r}, {len(self._fields)} fields)"
+
+
+def parse_cuda_signature(
+    source: str,
+    name: str,
+    *,
+    structs: Iterable[CudaStruct] = (),
+    template_args: Sequence[Any] | None = None,
+) -> tuple[CudaParameter, ...]:
+    """Parse the parameters of a ``__global__`` function.
+
+    C types map to dtypes as on Linux (LP64): ``int`` is int32, ``long`` and
+    ``long long`` are int64.
+
+    Parameters
+    ----------
+    source : str
+        CUDA C source code.
+    name : str
+        Name of the ``__global__`` function.
+    structs : iterable of CudaStruct, optional
+        Struct types that may appear as parameters (by value). A definition
+        of the struct in `source` must match.
+    template_args : sequence, optional
+        Template arguments if `name` is a function template: C types or
+        dtypes (see :func:`~cunumpy.cuda.ctype_of`) for type parameters,
+        ints or bools for non-type parameters.
+
+    Returns
+    -------
+    tuple of CudaParameter
+        The parameters, in order.
+
+    Raises
+    ------
+    ValueError
+        If there is no such function, `template_args` do not fit, a struct
+        definition does not match its :class:`~cunumpy.arguments.CudaStruct`,
+        or a parameter type cannot be checked (e.g. a macro or ``double**``).
+
+    Examples
+    --------
+    >>> source = '''
+    ... template <typename T>
+    ... __global__ void scale(T* x, T factor, int n) {}'''
+    >>> [(p.name, p.ctype, p.pointer) for p in xp.cuda.parse_cuda_signature(
+    ...     source, "scale", template_args=[np.float32])]
+    [('x', 'float', True), ('factor', 'float', False), ('n', 'int', False)]
+    """
+    code = _strip_comments(source)
+    structs = {s.name: s for s in structs}
+    for struct in structs.values():
+        struct.check_source(code)
+
+    pattern = r"(?:template\s*<(?P<template>[^{};]*?)>\s*)?__global__\s+void\s+"
+    match = re.search(pattern + re.escape(name) + r"\s*\(", code)
+    if match is None:
+        raise ValueError(f"no __global__ function {name!r} found in the CUDA source")
+
+    depth, start = 1, match.end()
+    for pos in range(start, len(code)):
+        if code[pos] == "(":
+            depth += 1
+        elif code[pos] == ")":
+            depth -= 1
+            if depth == 0:
+                break
+    else:
+        raise ValueError(f"unbalanced parentheses in the signature of {name!r}")
+    params = code[start:pos].strip()
+
+    template = match.group("template")
+    if template is not None:
+        template_params = [p.split() for p in _split_top_level(template) if p.strip()]
+        if template_args is None or len(template_args) != len(template_params):
+            raise ValueError(
+                f"{name!r} is a template with {len(template_params)} parameters; "
+                f"pass them as template_args",
+            )
+        for words, value in zip(template_params, template_args):
+            if words[0] in ("typename", "class"):
+                params = re.sub(
+                    r"\b" + re.escape(words[-1]) + r"\b",
+                    _template_arg(value),
+                    params,
+                )
+    elif template_args:
+        raise ValueError(f"{name!r} is not a template, but template_args were given")
+
+    if params in ("", "void"):
+        return ()
+    return tuple(_parse_parameter(p, structs) for p in _split_top_level(params))
+
+
+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 write_cuda_header(
+    path: str | Path,
+    structs: Iterable[CudaStruct],
+    guard: str | None = None,
+    *,
+    includes: Iterable[str] = (),
+) -> str:
+    """Write the definitions of several structs to one header file.
+
+    The header has an include guard, ``#include "cunumpy/array_view.cuh"`` if
+    a struct has array view fields, any `includes`, then the structs in
+    order. Commit it next to the kernels and let a test regenerate and
+    compare it (see :meth:`CudaStruct.to_header`).
+
+    Parameters
+    ----------
+    path : str or Path
+        File to write.
+    structs : iterable of CudaStruct
+        The structs, in declaration order.
+    guard : str, optional
+        Include guard macro; by default from the file name
+        (``pusher_args.cuh`` -> ``PUSHER_ARGS_CUH``).
+    includes : iterable of str, optional
+        Additional headers, as file names or ``#include`` lines.
+
+    Returns
+    -------
+    str
+        The header source that was written.
+
+    Examples
+    --------
+    >>> import os, tempfile
+    >>> Vec = xp.arguments.CudaStruct("Vec", [("data", "double*"), ("n", "int")])
+    >>> path = os.path.join(tempfile.mkdtemp(), "vec.cuh")
+    >>> print(xp.arguments.write_cuda_header(path, [Vec]), end="")
+    // Generated by cunumpy.arguments.CudaStruct from the Python definition; do not edit.
+    #ifndef VEC_CUH
+    #define VEC_CUH
+    <BLANKLINE>
+    struct Vec {
+        double* data;
+        int n;
+    };
+    <BLANKLINE>
+    #endif  // VEC_CUH
+    """
+    path = Path(path)
+    source = _header_source(tuple(structs), guard or _header_guard(path.name), includes)
+    path.write_text(source)
+    return source
+
 
 class CudaStructValue:
     """A value of a :class:`CudaStruct`, passed to kernels as one argument.
@@ -1622,9 +1622,6 @@ def __init__(self, struct: CudaStruct, packed: np.void, values: dict) -> None:
         self._packed = packed
         self._values = dict(values)
 
-    def __repr__(self) -> str:
-        return f"<{self._struct.name} value>"
-
     @property
     def struct(self) -> CudaStruct:
         """The struct type."""
@@ -1635,13 +1632,36 @@ def packed(self) -> np.void:
         """The packed struct, with the memory layout of the C struct."""
         return self._packed
 
-    def __getitem__(self, field: str) -> Any:
-        return self._values[field]
-
     def __cuda_args__(self) -> tuple[np.void]:
         """Return the packed struct, the one kernel argument this value stands for."""
         return (self._packed,)
 
+    def __getitem__(self, field: str) -> Any:
+        return self._values[field]
+
+    def __repr__(self) -> str:
+        return f"<{self._struct.name} value>"
+
+
+def _field_state(struct: CudaStruct, values: Mapping[str, Any]) -> tuple[Any, ...]:
+    """Return what the packed struct depends on, to detect changed fields."""
+    state = []
+    for field in struct.fields:
+        value = values[field.name]
+        if field.pointer or field.view_ndim is not None:
+            ptr = getattr(getattr(value, "data", None), "ptr", None)
+            if ptr is None:
+                state.append(("id", id(value)))
+            elif field.pointer:
+                state.append(ptr)
+            else:
+                state.append((ptr, tuple(value.shape), tuple(value.strides)))
+        elif isinstance(value, (bool, int, float, complex, np.generic)):
+            state.append((type(value), value))
+        else:
+            state.append(("id", id(value)))
+    return tuple(state)
+
 
 class CudaStructArguments(CudaArguments):
     """Base class for argument objects passed to CUDA kernels as one C struct.
@@ -1722,6 +1742,19 @@ def __init_subclass__(cls, **kwargs: Any) -> None:
             )
         cls.struct = CudaStruct(cls.struct_name, cls.fields)
 
+    def _field_values(self, struct: CudaStruct) -> dict[str, Any]:
+        """Return the current value of every field attribute, by field name."""
+        values = {}
+        for field in struct.fields:
+            try:
+                values[field.name] = getattr(self, field.name)
+            except AttributeError:
+                raise AttributeError(
+                    f"{type(self).__qualname__} has no attribute {field.name!r} "
+                    f"for the field of struct {struct.name}",
+                ) from None
+        return values
+
     def pack(self) -> None:
         """Pack the field attributes into the struct.
 
@@ -1745,19 +1778,6 @@ def pack(self) -> None:
         self._struct_value = struct(**values)
         self._packed_state = _field_state(struct, values)
 
-    def _field_values(self, struct: CudaStruct) -> dict[str, Any]:
-        """Return the current value of every field attribute, by field name."""
-        values = {}
-        for field in struct.fields:
-            try:
-                values[field.name] = getattr(self, field.name)
-            except AttributeError:
-                raise AttributeError(
-                    f"{type(self).__qualname__} has no attribute {field.name!r} "
-                    f"for the field of struct {struct.name}",
-                ) from None
-        return values
-
     @property
     def packed(self) -> np.void:
         """The packed struct, repacked first if a field attribute changed."""
@@ -1789,26 +1809,6 @@ def __setstate__(self, state: dict[str, Any]) -> None:
         self.pack()
 
 
-def _field_state(struct: CudaStruct, values: Mapping[str, Any]) -> tuple[Any, ...]:
-    """Return what the packed struct depends on, to detect changed fields."""
-    state = []
-    for field in struct.fields:
-        value = values[field.name]
-        if field.pointer or field.view_ndim is not None:
-            ptr = getattr(getattr(value, "data", None), "ptr", None)
-            if ptr is None:
-                state.append(("id", id(value)))
-            elif field.pointer:
-                state.append(ptr)
-            else:
-                state.append((ptr, tuple(value.shape), tuple(value.strides)))
-        elif isinstance(value, (bool, int, float, complex, np.generic)):
-            state.append((type(value), value))
-        else:
-            state.append(("id", id(value)))
-    return tuple(state)
-
-
 def _is_device_array(value: Any) -> bool:
     return hasattr(type(value), "__cuda_array_interface__") or hasattr(
         value,
@@ -1895,6 +1895,17 @@ def _device_arrays_in(args: Sequence[Any]) -> Iterator[tuple[str, Any]]:
                     yield f"{label}[{j}]", value
 
 
+def _is_capturing(stream: Any) -> bool:
+    """Whether `stream` is being captured into a CUDA graph."""
+    is_capturing = getattr(stream, "is_capturing", None)
+    if is_capturing is None:
+        return False
+    try:
+        return bool(is_capturing())
+    except Exception:  # noqa: BLE001 - e.g. the legacy null stream cannot capture
+        return False
+
+
 class CudaKernel:
     """A CUDA C kernel, compiled with NVRTC through CuPy.
 
@@ -1976,6 +1987,17 @@ class CudaKernel:
     >>> axpy(2.0, x, y, x.size, n_threads=x.size)  # doctest: +SKIP
     """
 
+    @staticmethod
+    def _check_block(block: tuple[int, ...]) -> tuple[int, ...]:
+        if any(n <= 0 for n in block):
+            raise ValueError(f"block sizes must be positive, got {block}")
+        if math.prod(block) > _MAX_THREADS_PER_BLOCK:
+            raise ValueError(
+                f"a block has at most {_MAX_THREADS_PER_BLOCK} threads, got "
+                f"{block} = {math.prod(block)}",
+            )
+        return block
+
     def __init__(
         self,
         source: str,
@@ -2112,20 +2134,6 @@ def all_from_file(cls, path: str | Path, **kwargs: Any) -> dict[str, CudaKernel]
             for name in names
         }
 
-    @staticmethod
-    def _check_block(block: tuple[int, ...]) -> tuple[int, ...]:
-        if any(n <= 0 for n in block):
-            raise ValueError(f"block sizes must be positive, got {block}")
-        if math.prod(block) > _MAX_THREADS_PER_BLOCK:
-            raise ValueError(
-                f"a block has at most {_MAX_THREADS_PER_BLOCK} threads, got "
-                f"{block} = {math.prod(block)}",
-            )
-        return block
-
-    def __repr__(self) -> str:
-        return f"CudaKernel(name={self.expression!r}, block_size={self.block_size})"
-
     @property
     def name(self) -> str:
         """Name of the kernel function."""
@@ -2239,6 +2247,9 @@ def n_threads_from(self) -> Callable[[tuple[Any, ...]], Any] | None:
         """The launch size inference: a callable of the argument tuple, or None; settable as in the constructor."""
         return self._n_threads_from
 
+    def _default_n_threads(self, args: tuple[Any, ...]) -> int | tuple[int, ...]:
+        return _first_array_shape(args, len(self._block))
+
     @n_threads_from.setter
     def n_threads_from(
         self,
@@ -2259,9 +2270,6 @@ def n_threads_from(
         self._automatic_threads = automatic
         self._n_threads_from = value
 
-    def _default_n_threads(self, args: tuple[Any, ...]) -> int | tuple[int, ...]:
-        return _first_array_shape(args, len(self._block))
-
     @property
     def check_finite(self) -> bool:
         """Whether every launch synchronizes and raises on NaN or inf in its float arrays (debugging; settable)."""
@@ -2561,88 +2569,6 @@ def launch_shape(
         grid_shape = tuple((n + b - 1) // b for n, b in zip(threads, block_shape))
         return grid_shape, block_shape
 
-    def __call__(
-        self,
-        *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:
-        """Launch the kernel (asynchronously, unless in debug mode).
-
-        The grid is `n_threads` divided by the block, rounded up, or an
-        explicit `grid`; with neither, the size is inferred
-        (:attr:`n_threads_from`).
-
-        Parameters
-        ----------
-        *args
-            Kernel arguments: CuPy arrays, scalars and argument objects, see
-            :meth:`prepare_args`.
-        n_threads : int or sequence of int, optional
-            Number of threads in 1 to 3 dimensions, e.g. ``(nx, ny)``. With a
-            1D block size the block is ``(block_size, 1, ...)``.
-        grid : int or sequence of int, optional
-            Number of blocks in 1 to 3 dimensions, instead of `n_threads`.
-        block : int or sequence of int, optional
-            Block shape for this call, instead of :attr:`block_size`.
-        shared_mem : int, optional
-            Dynamic shared memory per block in bytes (``extern __shared__``).
-        stream : cupy.cuda.Stream, optional
-            Stream to launch on; the current stream by default.
-
-        Raises
-        ------
-        RuntimeError
-            In debug mode, a CUDA error found by synchronizing after the
-            launch (CuPy's error chained); otherwise such errors surface at a
-            later synchronization. Also a NaN or inf with :attr:`check_finite`.
-        ValueError
-            If the launch exceeds the device or kernel limits (dimensions,
-            threads, shared memory), or arrays or stream are on another device.
-        """
-        grid_shape, block_shape = self.launch_shape(
-            n_threads, grid=grid, block=block, args=args
-        )
-        shared_mem = operator.index(shared_mem)
-        if shared_mem < 0:
-            raise ValueError(f"shared_mem must be non-negative, got {shared_mem}")
-        values = self.prepare_args(*args)
-        # RawKernel accepts size-one NumPy arrays for structs passed by value,
-        # but not structured NumPy scalars (np.void). Keep their packed bytes
-        # and alignment intact, including array-view pointers and strides.
-        values = tuple(np.asarray(v) if isinstance(v, np.void) else v for v in values)
-        if 0 in grid_shape:
-            return
-
-        kernel = self.compile()
-        device = _current_device()
-        state = self._compiled[device]
-        self._validate_launch(state, device, grid_shape, block_shape, shared_mem)
-        if stream is not None and getattr(stream, "device_id", device) not in (
-            device,
-            -1,
-        ):
-            raise ValueError(
-                f"kernel {self.expression!r}: stream belongs to another device"
-            )
-        for label, array in _device_arrays_in(args):
-            array_device = getattr(getattr(array, "device", None), "id", None)
-            if array_device is not None and array_device != device:
-                raise ValueError(
-                    f"kernel {self.expression!r} on CUDA device {device}: "
-                    f"{label} belongs to CUDA device {array_device}",
-                )
-        debug = self.debug_active()
-        with stream if stream is not None else nullcontext():
-            kernel(grid_shape, block_shape, values, shared_mem=shared_mem)
-            if debug or self._check_finite:
-                self._synchronize_after_launch(stream, grid_shape, block_shape)
-            if self._check_finite:
-                self._check_finite_arrays(args)
-
     def _validate_launch(
         self,
         state: _CompiledKernel,
@@ -2732,16 +2658,90 @@ def _synchronize_after_launch(
                 f"grid {grid} and block {block}: {error}",
             ) from error
 
+    def __call__(
+        self,
+        *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:
+        """Launch the kernel (asynchronously, unless in debug mode).
 
-def _is_capturing(stream: Any) -> bool:
-    """Whether `stream` is being captured into a CUDA graph."""
-    is_capturing = getattr(stream, "is_capturing", None)
-    if is_capturing is None:
-        return False
-    try:
-        return bool(is_capturing())
-    except Exception:  # noqa: BLE001 - e.g. the legacy null stream cannot capture
-        return False
+        The grid is `n_threads` divided by the block, rounded up, or an
+        explicit `grid`; with neither, the size is inferred
+        (:attr:`n_threads_from`).
+
+        Parameters
+        ----------
+        *args
+            Kernel arguments: CuPy arrays, scalars and argument objects, see
+            :meth:`prepare_args`.
+        n_threads : int or sequence of int, optional
+            Number of threads in 1 to 3 dimensions, e.g. ``(nx, ny)``. With a
+            1D block size the block is ``(block_size, 1, ...)``.
+        grid : int or sequence of int, optional
+            Number of blocks in 1 to 3 dimensions, instead of `n_threads`.
+        block : int or sequence of int, optional
+            Block shape for this call, instead of :attr:`block_size`.
+        shared_mem : int, optional
+            Dynamic shared memory per block in bytes (``extern __shared__``).
+        stream : cupy.cuda.Stream, optional
+            Stream to launch on; the current stream by default.
+
+        Raises
+        ------
+        RuntimeError
+            In debug mode, a CUDA error found by synchronizing after the
+            launch (CuPy's error chained); otherwise such errors surface at a
+            later synchronization. Also a NaN or inf with :attr:`check_finite`.
+        ValueError
+            If the launch exceeds the device or kernel limits (dimensions,
+            threads, shared memory), or arrays or stream are on another device.
+        """
+        grid_shape, block_shape = self.launch_shape(
+            n_threads, grid=grid, block=block, args=args
+        )
+        shared_mem = operator.index(shared_mem)
+        if shared_mem < 0:
+            raise ValueError(f"shared_mem must be non-negative, got {shared_mem}")
+        values = self.prepare_args(*args)
+        # RawKernel accepts size-one NumPy arrays for structs passed by value,
+        # but not structured NumPy scalars (np.void). Keep their packed bytes
+        # and alignment intact, including array-view pointers and strides.
+        values = tuple(np.asarray(v) if isinstance(v, np.void) else v for v in values)
+        if 0 in grid_shape:
+            return
+
+        kernel = self.compile()
+        device = _current_device()
+        state = self._compiled[device]
+        self._validate_launch(state, device, grid_shape, block_shape, shared_mem)
+        if stream is not None and getattr(stream, "device_id", device) not in (
+            device,
+            -1,
+        ):
+            raise ValueError(
+                f"kernel {self.expression!r}: stream belongs to another device"
+            )
+        for label, array in _device_arrays_in(args):
+            array_device = getattr(getattr(array, "device", None), "id", None)
+            if array_device is not None and array_device != device:
+                raise ValueError(
+                    f"kernel {self.expression!r} on CUDA device {device}: "
+                    f"{label} belongs to CUDA device {array_device}",
+                )
+        debug = self.debug_active()
+        with stream if stream is not None else nullcontext():
+            kernel(grid_shape, block_shape, values, shared_mem=shared_mem)
+            if debug or self._check_finite:
+                self._synchronize_after_launch(stream, grid_shape, block_shape)
+            if self._check_finite:
+                self._check_finite_arrays(args)
+
+    def __repr__(self) -> str:
+        return f"CudaKernel(name={self.expression!r}, block_size={self.block_size})"
 
 
 class CudaKernelVariants:
@@ -2771,9 +2771,6 @@ def __init__(self, factory: Callable[..., CudaKernel]) -> None:
         self._factory = factory
         self._kernels: dict[tuple[Hashable, ...], CudaKernel] = {}
 
-    def __repr__(self) -> str:
-        return f"CudaKernelVariants({len(self._kernels)} variants)"
-
     def get(self, *key: Hashable) -> CudaKernel:
         """Return the kernel for a variant, created on first use.
 
@@ -2802,13 +2799,6 @@ def get(self, *key: Hashable) -> CudaKernel:
             self._kernels[key] = kernel
         return kernel
 
-    def __len__(self) -> int:
-        return len(self._kernels)
-
-    def __iter__(self) -> Iterator[tuple[Hashable, ...]]:
-        """Iterate over the keys of the variants created so far."""
-        return iter(list(self._kernels))
-
     def keys(self) -> list[tuple[Hashable, ...]]:
         """Return the keys of the variants created so far.
 
@@ -2843,3 +2833,13 @@ def compile_all(
             {key: kernel.compile for key, kernel in self._kernels.items()},
             jobs,
         )
+
+    def __iter__(self) -> Iterator[tuple[Hashable, ...]]:
+        """Iterate over the keys of the variants created so far."""
+        return iter(list(self._kernels))
+
+    def __len__(self) -> int:
+        return len(self._kernels)
+
+    def __repr__(self) -> str:
+        return f"CudaKernelVariants({len(self._kernels)} variants)"
diff --git a/src/cunumpy/_dispatch.py b/src/cunumpy/_dispatch.py
index b3bc8a2..9b78cc7 100644
--- a/src/cunumpy/_dispatch.py
+++ b/src/cunumpy/_dispatch.py
@@ -419,12 +419,6 @@ def from_folder(
             test_args=test_args,
         )
 
-    def __repr__(self) -> str:
-        return (
-            f"Kernel(name={self._name!r}, cuda={self._cuda_kernel is not None}, "
-            f"missing_cuda={self._missing_cuda!r})"
-        )
-
     @property
     def name(self) -> str:
         """Name of the kernel."""
@@ -549,6 +543,32 @@ def implementations(self) -> tuple[str, ...]:  # numpydoc ignore=RT01
         )
         return (*host, "cuda") if self._cuda_kernel is not None else host
 
+    def _device_kernel(self) -> PyccelKernel | CudaKernel:
+        """Resolve the device selection, honoring fallback only in automatic mode."""
+        if self._cuda_kernel is not None:
+            return self._cuda_kernel
+        if get_device_kernel_implementation() == "cuda":
+            raise LookupError(
+                f"kernel {self._name!r} has no 'cuda' implementation "
+                "(explicitly selected device implementation)",
+            )
+        if self._missing_cuda == "raise":
+            expected = (
+                "" if self._cuda_path is None else f" (expected {self._cuda_path})"
+            )
+            raise NotImplementedError(
+                f"No CUDA version of kernel {self._name!r}{expected}.",
+            )
+        if not self._warned:
+            warnings.warn(
+                f"No CUDA version of kernel {self._name!r}: calling the host kernel, "
+                "which copies its arrays to the host and back at every call.",
+                RuntimeWarning,
+                stacklevel=3,
+            )
+            self._warned = True
+        return self._host_kernel
+
     def selected(self, device: bool = False) -> str:
         """Return the implementation a call runs now, e.g. to rule out a slow path.
 
@@ -608,32 +628,6 @@ def get_kernel(self) -> PyccelKernel | CudaKernel:
             return self._host_kernel
         return self._device_kernel()
 
-    def _device_kernel(self) -> PyccelKernel | CudaKernel:
-        """Resolve the device selection, honoring fallback only in automatic mode."""
-        if self._cuda_kernel is not None:
-            return self._cuda_kernel
-        if get_device_kernel_implementation() == "cuda":
-            raise LookupError(
-                f"kernel {self._name!r} has no 'cuda' implementation "
-                "(explicitly selected device implementation)",
-            )
-        if self._missing_cuda == "raise":
-            expected = (
-                "" if self._cuda_path is None else f" (expected {self._cuda_path})"
-            )
-            raise NotImplementedError(
-                f"No CUDA version of kernel {self._name!r}{expected}.",
-            )
-        if not self._warned:
-            warnings.warn(
-                f"No CUDA version of kernel {self._name!r}: calling the host kernel, "
-                "which copies its arrays to the host and back at every call.",
-                RuntimeWarning,
-                stacklevel=3,
-            )
-            self._warned = True
-        return self._host_kernel
-
     def compile(self) -> bool:
         """Compile the CUDA kernel now, if there is one.
 
@@ -715,6 +709,12 @@ def __call__(
                 stream=stream,
             )
 
+    def __repr__(self) -> str:
+        return (
+            f"Kernel(name={self._name!r}, cuda={self._cuda_kernel is not None}, "
+            f"missing_cuda={self._missing_cuda!r})"
+        )
+
 
 class KernelCatalog(Mapping):
     """A read-only mapping from names to :class:`~cunumpy.kernels.Kernel` objects.
@@ -889,23 +889,6 @@ def register(self, kernel: Kernel, name: str | None = None) -> Kernel:
         self._kernels[name] = kernel
         return kernel
 
-    def __getitem__(self, name: str) -> Kernel:
-        return self._kernels[name]
-
-    def __iter__(self) -> Iterator[str]:
-        return iter(self._kernels)
-
-    def __len__(self) -> int:
-        return len(self._kernels)
-
-    def __repr__(self) -> str:
-        return (
-            f"KernelCatalog({len(self)} kernels, {len(self.without_cuda)} without CUDA)"
-        )
-
-    def __str__(self) -> str:
-        return self.summary()
-
     def summary(self, max_missing: int = 10) -> str:
         """Return one line on the porting status, e.g. for a ``--status`` command.
 
@@ -1015,3 +998,20 @@ def compile_all(self, jobs: int | None = 1) -> list[str]:
             },
             jobs,
         )
+
+    def __getitem__(self, name: str) -> Kernel:
+        return self._kernels[name]
+
+    def __iter__(self) -> Iterator[str]:
+        return iter(self._kernels)
+
+    def __len__(self) -> int:
+        return len(self._kernels)
+
+    def __repr__(self) -> str:
+        return (
+            f"KernelCatalog({len(self)} kernels, {len(self.without_cuda)} without CUDA)"
+        )
+
+    def __str__(self) -> str:
+        return self.summary()
diff --git a/src/cunumpy/_fake_cupy.py b/src/cunumpy/_fake_cupy.py
index a96b1be..73e22ed 100644
--- a/src/cunumpy/_fake_cupy.py
+++ b/src/cunumpy/_fake_cupy.py
@@ -75,6 +75,16 @@ def host_buffer(array: object) -> object:
     )
 
 
+def uninstall() -> None:
+    """Remove the fake ``cupy`` package from ``sys.modules`` (no-op otherwise)."""
+    if not is_active():
+        return
+    for name in [n for n in sys.modules if n.split(".")[0] in ("cupy", "cupyx")]:
+        del sys.modules[name]
+    for name in [n for n in sys.modules if n.startswith("array_api_compat.cupy")]:
+        del sys.modules[name]
+
+
 def install() -> types.ModuleType:
     """Install the fake ``cupy`` package into ``sys.modules`` and return it.
 
@@ -114,13 +124,3 @@ def install() -> types.ModuleType:
         uninstall()
         raise
     return module
-
-
-def uninstall() -> None:
-    """Remove the fake ``cupy`` package from ``sys.modules`` (no-op otherwise)."""
-    if not is_active():
-        return
-    for name in [n for n in sys.modules if n.split(".")[0] in ("cupy", "cupyx")]:
-        del sys.modules[name]
-    for name in [n for n in sys.modules if n.startswith("array_api_compat.cupy")]:
-        del sys.modules[name]
diff --git a/src/cunumpy/_fake_cupy_impl.py b/src/cunumpy/_fake_cupy_impl.py
index 7f6fc00..c25fb5b 100644
--- a/src/cunumpy/_fake_cupy_impl.py
+++ b/src/cunumpy/_fake_cupy_impl.py
@@ -30,6 +30,24 @@ def _err(obj, where=""):
     )
 
 
+def _strict(args, where, first_is_data=False):
+    """Raise on NumPy arrays anywhere (top level or inside tuples/lists).
+
+    0-d NumPy arrays pass, like NumPy scalars: CuPy treats them as scalars.
+    With `first_is_data`, a list or tuple as first argument raises too.
+    """
+    for i, a in enumerate(args):
+        if isinstance(a, _HOST) and a.ndim > 0:
+            raise _err(a, f" in {where}")
+        if isinstance(a, (list, tuple)):
+            if i == 0 and first_is_data:
+                raise _err(a, f" as array argument of {where}")
+            _strict(a, where)
+
+
+cuda = types.ModuleType("cupy.cuda")
+
+
 class ndarray:
     __slots__ = ("__weakref__", "_a")
     __array_priority__ = 100
@@ -38,6 +56,12 @@ def __init__(self, a):
         assert isinstance(a, _HOST)
         object.__setattr__(self, "_a", a)
 
+    def __copy__(self):
+        return ndarray(self._a.copy())
+
+    def __deepcopy__(self, memo):
+        return ndarray(self._a.copy())
+
     # --- conversions ---------------------------------------------------------
     def __array__(self, *args, **kwargs):
         raise TypeError(
@@ -56,9 +80,6 @@ def get(self, stream=None, order="C", out=None, blocking=True):
     def set(self, arr, *args, **kwargs):
         self._a[...] = arr
 
-    def __reduce__(self):
-        return (_from_host, (self._a.copy(),))
-
     @property
     def __cuda_array_interface__(self):
         a = self._a
@@ -78,6 +99,21 @@ def data(self):
     def device(self):
         return cuda.Device(0)
 
+    # --- numpy protocols -----------------------------------------------------
+    def __array_ufunc__(self, ufunc, method, *inputs, **kwargs):
+        _strict(inputs, ufunc.__name__)
+        out = kwargs.get("out")
+        res = getattr(ufunc, method)(*_unwrap(inputs), **_unwrap(kwargs))
+        if out is not None:
+            return out[0] if isinstance(out, tuple) and len(out) == 1 else out
+        return _wrap(res)
+
+    def __array_function__(self, func, types_, args, kwargs):
+        mine = globals().get(func.__name__)
+        if mine is None:
+            mine = __getattr__(func.__name__)
+        return mine(*args, **kwargs)
+
     # --- attributes ----------------------------------------------------------
     def __getattr__(self, name):
         # no __array_interface__ etc.: NumPy must not see the host buffer
@@ -93,27 +129,33 @@ def __getattr__(self, name):
     def __setattr__(self, name, value):
         setattr(self._a, name, _unwrap(value))
 
-    # --- python protocol -----------------------------------------------------
-    def __len__(self):
-        return len(self._a)
+    def __getitem__(self, key):
+        return _wrap(self._a[_unwrap(key)])
+
+    def __setitem__(self, key, value):
+        self._a[_unwrap(key)] = _unwrap(value)
 
     def __iter__(self):
         for x in self._a:
             yield _wrap(x)
 
-    def __repr__(self):
-        return "fakecupy." + repr(self._a)
+    def __contains__(self, x):
+        return _unwrap(x) in self._a
 
-    def __str__(self):
-        return str(self._a)
+    # --- python protocol -----------------------------------------------------
+    def __len__(self):
+        return len(self._a)
 
-    def __format__(self, spec):
-        return format(self._a.item() if self._a.ndim == 0 else self._a, spec)
+    __hash__ = None
 
     def __bool__(self):
         _sync("bool(device array)")
         return bool(self._a)
 
+    def __complex__(self):
+        _sync("complex(device array)")
+        return complex(self._a)
+
     def __int__(self):
         _sync("int(device array)")
         return int(self._a)
@@ -122,45 +164,21 @@ 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
-
-    def __contains__(self, x):
-        return _unwrap(x) in self._a
-
-    def __getitem__(self, key):
-        return _wrap(self._a[_unwrap(key)])
-
-    def __setitem__(self, key, value):
-        self._a[_unwrap(key)] = _unwrap(value)
-
-    def __copy__(self):
-        return ndarray(self._a.copy())
+    def __reduce__(self):
+        return (_from_host, (self._a.copy(),))
 
-    def __deepcopy__(self, memo):
-        return ndarray(self._a.copy())
+    def __repr__(self):
+        return "fakecupy." + repr(self._a)
 
-    # --- numpy protocols -----------------------------------------------------
-    def __array_ufunc__(self, ufunc, method, *inputs, **kwargs):
-        _strict(inputs, ufunc.__name__)
-        out = kwargs.get("out")
-        res = getattr(ufunc, method)(*_unwrap(inputs), **_unwrap(kwargs))
-        if out is not None:
-            return out[0] if isinstance(out, tuple) and len(out) == 1 else out
-        return _wrap(res)
+    def __str__(self):
+        return str(self._a)
 
-    def __array_function__(self, func, types_, args, kwargs):
-        mine = globals().get(func.__name__)
-        if mine is None:
-            mine = __getattr__(func.__name__)
-        return mine(*args, **kwargs)
+    def __format__(self, spec):
+        return format(self._a.item() if self._a.ndim == 0 else self._a, spec)
 
 
 def _from_host(a):
@@ -195,21 +213,6 @@ def _unwrap(x):
     return x
 
 
-def _strict(args, where, first_is_data=False):
-    """Raise on NumPy arrays anywhere (top level or inside tuples/lists).
-
-    0-d NumPy arrays pass, like NumPy scalars: CuPy treats them as scalars.
-    With `first_is_data`, a list or tuple as first argument raises too.
-    """
-    for i, a in enumerate(args):
-        if isinstance(a, _HOST) and a.ndim > 0:
-            raise _err(a, f" in {where}")
-        if isinstance(a, (list, tuple)):
-            if i == 0 and first_is_data:
-                raise _err(a, f" as array argument of {where}")
-            _strict(a, where)
-
-
 # binary / unary operators ------------------------------------------------------
 def _binop(op):
     def f(self, other):
@@ -413,15 +416,15 @@ class _NoGPU:
     def __init__(self, *args, **kwargs):
         pass
 
-    def __call__(self, *args, **kwargs):
-        raise NotImplementedError("fake CuPy cannot run CUDA kernels")
-
     def compile(self, log_stream=None):
         raise NotImplementedError("fake CuPy cannot compile CUDA kernels")
 
     def get_function(self, name):
         return _NoGPU()
 
+    def __call__(self, *args, **kwargs):
+        raise NotImplementedError("fake CuPy cannot run CUDA kernels")
+
 
 RawKernel = RawModule = ElementwiseKernel = ReductionKernel = _NoGPU
 
@@ -430,12 +433,6 @@ class _Device:
     def __init__(self, id=0):
         self.id = 0 if id is None else id
 
-    def __enter__(self):
-        return self
-
-    def __exit__(self, *exc):
-        return False
-
     def use(self):
         pass
 
@@ -446,6 +443,12 @@ def synchronize(self):
     def attributes(self):
         return {"MaxSharedMemoryPerBlock": 48 * 1024}
 
+    def __enter__(self):
+        return self
+
+    def __exit__(self, *exc):
+        return False
+
 
 class _Event:
     def __init__(self, **kwargs):
@@ -466,9 +469,6 @@ class _Stream(_Device):
     def __init__(self, *args, **kwargs):
         super().__init__()
 
-    def __enter__(self):
-        return self
-
     def record(self, event=None):
         event = _Event() if event is None else event
         event.record(self)
@@ -481,6 +481,9 @@ def wait_event(self, event):
     def done(self):
         return True
 
+    def __enter__(self):
+        return self
+
 
 _Stream.null = _Stream()
 
@@ -489,7 +492,6 @@ def _alloc_pinned_memory(nbytes):
     raise NotImplementedError("fake CuPy has no pinned memory")
 
 
-cuda = types.ModuleType("cupy.cuda")
 cuda.Device = _Device
 cuda.Stream = _Stream
 cuda.Event = _Event
diff --git a/src/cunumpy/_fake_cupyx.py b/src/cunumpy/_fake_cupyx.py
index cd8954c..b861f84 100644
--- a/src/cunumpy/_fake_cupyx.py
+++ b/src/cunumpy/_fake_cupyx.py
@@ -248,24 +248,6 @@ def __init__(self, obj: Any, cupy: Any) -> None:
         object.__setattr__(self, "_obj", obj)
         object.__setattr__(self, "_cupy", cupy)
 
-    def __getattr__(self, name: str) -> Any:
-        attr = getattr(self._obj, name)
-        if callable(attr):
-            return _device_callable(
-                attr, self._cupy, f"{type(self._obj).__name__}.{name}"
-            )
-        return _to_device(attr, self._cupy)
-
-    def __setattr__(self, name: str, value: Any) -> None:
-        setattr(self._obj, name, _to_host(value, self._cupy))
-
-    def __call__(self, *args: Any, **kwargs: Any) -> Any:
-        where = f"{type(self._obj).__name__}.__call__"
-        return _device_callable(self._obj, self._cupy, where)(*args, **kwargs)
-
-    def __repr__(self) -> str:
-        return f"<fake cupyx {self._obj!r}>"
-
     # copies wrap a copy of the SciPy object and share the fake CuPy module (a
     # module cannot be copied), as real CuPy splines can be copied
     def __copy__(self) -> _Proxy:
@@ -274,31 +256,23 @@ def __copy__(self) -> _Proxy:
     def __deepcopy__(self, memo: dict) -> _Proxy:
         return _Proxy(copy.deepcopy(self._obj, memo), self._cupy)
 
-
-class _ProxyClass:
-    """A SciPy class seen as a CuPy class: instances are :class:`_Proxy`."""
-
-    def __init__(self, cls: type, cupy: Any, where: str) -> None:
-        self._cls = cls
-        self._cupy = cupy
-        self._where = where
-        self.__name__ = cls.__name__
-        self.__doc__ = cls.__doc__
-
     def __call__(self, *args: Any, **kwargs: Any) -> Any:
-        return _device_callable(self._cls, self._cupy, self._where)(*args, **kwargs)
+        where = f"{type(self._obj).__name__}.__call__"
+        return _device_callable(self._obj, self._cupy, where)(*args, **kwargs)
 
     def __getattr__(self, name: str) -> Any:
-        attr = getattr(self._cls, name)
+        attr = getattr(self._obj, name)
         if callable(attr):
-            return _device_callable(attr, self._cupy, f"{self._where}.{name}")
-        return attr
+            return _device_callable(
+                attr, self._cupy, f"{type(self._obj).__name__}.{name}"
+            )
+        return _to_device(attr, self._cupy)
 
-    def __instancecheck__(self, obj: Any) -> bool:
-        return isinstance(obj, _Proxy) and isinstance(obj._obj, self._cls)
+    def __setattr__(self, name: str, value: Any) -> None:
+        setattr(self._obj, name, _to_host(value, self._cupy))
 
     def __repr__(self) -> str:
-        return f"<fake {self._where}>"
+        return f"<fake cupyx {self._obj!r}>"
 
 
 def _is_scipy_object(value: Any) -> bool:
@@ -335,6 +309,32 @@ def call(*args: Any, **kwargs: Any) -> Any:
     return call
 
 
+class _ProxyClass:
+    """A SciPy class seen as a CuPy class: instances are :class:`_Proxy`."""
+
+    def __init__(self, cls: type, cupy: Any, where: str) -> None:
+        self._cls = cls
+        self._cupy = cupy
+        self._where = where
+        self.__name__ = cls.__name__
+        self.__doc__ = cls.__doc__
+
+    def __instancecheck__(self, obj: Any) -> bool:
+        return isinstance(obj, _Proxy) and isinstance(obj._obj, self._cls)
+
+    def __call__(self, *args: Any, **kwargs: Any) -> Any:
+        return _device_callable(self._cls, self._cupy, self._where)(*args, **kwargs)
+
+    def __getattr__(self, name: str) -> Any:
+        attr = getattr(self._cls, name)
+        if callable(attr):
+            return _device_callable(attr, self._cupy, f"{self._where}.{name}")
+        return attr
+
+    def __repr__(self) -> str:
+        return f"<fake {self._where}>"
+
+
 def _subpackage(path: str, cupy: Any) -> types.ModuleType:
     module = types.ModuleType(f"cupyx.scipy.{path}")
     allowed = NAMES[path]
diff --git a/src/cunumpy/_kernel.py b/src/cunumpy/_kernel.py
index 4ae7c3c..bc46aa1 100644
--- a/src/cunumpy/_kernel.py
+++ b/src/cunumpy/_kernel.py
@@ -137,12 +137,6 @@ def __init__(
                     )
             self._outputs = tuple(outputs)
 
-    def __repr__(self) -> str:
-        return (
-            f"PyccelKernel(kernel={self.name!r}, use_cupy={self.use_cupy!r}, "
-            f"outputs={self._outputs!r})"
-        )
-
     def _parameter_names(self) -> Sequence[str] | None:
         """Return the names of the positional parameters, or None if unknown."""
         source = self._parameters
@@ -221,16 +215,6 @@ def _convert_to_numpy(
 
         return value
 
-    def _convert_from_numpy(self, value: Any) -> Any:
-        """Move host arrays returned by the kernel back to the device."""
-        if self._is_array(value):
-            return self._copy_to_device(value)
-        if isinstance(value, tuple):
-            return tuple(self._convert_from_numpy(item) for item in value)
-        if isinstance(value, list):
-            return [self._convert_from_numpy(item) for item in value]
-        return value
-
     def _copy_to_device(self, value: Any) -> Any:
         result = _host_to_device(value)
         if _COUNTERS:
@@ -241,6 +225,16 @@ def _copy_to_device(self, value: Any) -> Any:
             )
         return result
 
+    def _convert_from_numpy(self, value: Any) -> Any:
+        """Move host arrays returned by the kernel back to the device."""
+        if self._is_array(value):
+            return self._copy_to_device(value)
+        if isinstance(value, tuple):
+            return tuple(self._convert_from_numpy(item) for item in value)
+        if isinstance(value, list):
+            return [self._convert_from_numpy(item) for item in value]
+        return value
+
     def _collect_host_arrays(self, value: Any, found: set[int], seen: set[int]) -> None:
         """Record the id of every host array reachable from `value`."""
         if self._is_array(value):
@@ -366,6 +360,38 @@ def _needs_conversion(self, args: tuple[Any, ...], kwargs: dict[str, Any]) -> bo
         # passed in explicitly.
         return any(self._contains_cupy(value) for value in (*args, *kwargs.values()))
 
+    @property
+    def name(self) -> str:
+        """The name of the wrapped kernel."""
+        return getattr(self._kernel, "__name__", type(self._kernel).__name__)
+
+    @property
+    def kernel(self) -> Callable[..., Any]:
+        """The wrapped kernel."""
+        return self._kernel
+
+    @property
+    def use_cupy(self) -> bool:
+        """Whether calls currently convert between device and host arrays."""
+        if self._use_cupy is not None:
+            return self._use_cupy
+        return _cupy_backend()
+
+    @property
+    def object_modules(self) -> tuple[str, ...]:
+        """The module prefixes whose instances are searched for arrays."""
+        return self._object_modules
+
+    @property
+    def is_array(self) -> Callable[[Any], bool]:
+        """The predicate for host values to move back to the device."""
+        return self._is_array
+
+    @property
+    def outputs(self) -> tuple[int | str, ...] | None:
+        """The declared output arguments, or None if every array is copied back."""
+        return self._outputs
+
     def __call__(self, *args: Any, **kwargs: Any) -> Any:
         if not self._needs_conversion(args, kwargs):
             return self._kernel(*args, **kwargs)
@@ -405,37 +431,11 @@ def __call__(self, *args: Any, **kwargs: Any) -> Any:
 
         return self._convert_from_numpy(result)
 
-    @property
-    def name(self) -> str:
-        """The name of the wrapped kernel."""
-        return getattr(self._kernel, "__name__", type(self._kernel).__name__)
-
-    @property
-    def kernel(self) -> Callable[..., Any]:
-        """The wrapped kernel."""
-        return self._kernel
-
-    @property
-    def use_cupy(self) -> bool:
-        """Whether calls currently convert between device and host arrays."""
-        if self._use_cupy is not None:
-            return self._use_cupy
-        return _cupy_backend()
-
-    @property
-    def object_modules(self) -> tuple[str, ...]:
-        """The module prefixes whose instances are searched for arrays."""
-        return self._object_modules
-
-    @property
-    def is_array(self) -> Callable[[Any], bool]:
-        """The predicate for host values to move back to the device."""
-        return self._is_array
-
-    @property
-    def outputs(self) -> tuple[int | str, ...] | None:
-        """The declared output arguments, or None if every array is copied back."""
-        return self._outputs
+    def __repr__(self) -> str:
+        return (
+            f"PyccelKernel(kernel={self.name!r}, use_cupy={self.use_cupy!r}, "
+            f"outputs={self._outputs!r})"
+        )
 
 
 HOST_IMPLEMENTATIONS = ("pyccel", "numba", "numpy", "python")
@@ -693,9 +693,6 @@ def __init__(
         self.errors: dict[str, BaseException] = {}
         self._default: str | None = None
 
-    def __repr__(self) -> str:
-        return f"HostImplementations({self.__name__!r}, {list(self.names)})"
-
     @property
     def names(self) -> tuple[str, ...]:
         """The implementations this kernel has, loaded or not."""
@@ -813,6 +810,9 @@ def build(self) -> str:
     def __call__(self, *args: Any, **kwargs: Any) -> Any:
         return self.get(self.selected())(*args, **kwargs)
 
+    def __repr__(self) -> str:
+        return f"HostImplementations({self.__name__!r}, {list(self.names)})"
+
 
 class CompiledHostKernel:
     """Call a host kernel that is compiled on its first call, with a fallback.
@@ -867,14 +867,6 @@ def __init__(
         self.error: BaseException | None = None
         self._warned = False
 
-    def __repr__(self) -> str:
-        state = (
-            "not built"
-            if not self._built
-            else ("compiled" if self._compiled else "failed")
-        )
-        return f"CompiledHostKernel({self.__name__!r}, {state})"
-
     @property
     def module(self) -> ModuleType:
         """The (uncompiled) module that defines the kernel."""
@@ -930,6 +922,14 @@ def __call__(self, *args: Any, **kwargs: Any) -> Any:
                 kernel = self.python
         return kernel(*args, **kwargs)
 
+    def __repr__(self) -> str:
+        state = (
+            "not built"
+            if not self._built
+            else ("compiled" if self._compiled else "failed")
+        )
+        return f"CompiledHostKernel({self.__name__!r}, {state})"
+
 
 def _c_ordered_with_gaps(array: np.ndarray) -> bool:
     """Return whether `array` is in C order, possibly with gaps (e.g. ``a[:, :n]``)."""
diff --git a/src/cunumpy/_metal_kernel.py b/src/cunumpy/_metal_kernel.py
index b011d3a..d195e74 100644
--- a/src/cunumpy/_metal_kernel.py
+++ b/src/cunumpy/_metal_kernel.py
@@ -105,6 +105,13 @@ class MetalKernel:
     array([0., 2., 4., 6.], dtype=float32)
     """
 
+    @staticmethod
+    def _as_triple(value: int | Sequence[int], what: str, default: int) -> tuple:
+        values = (int(value),) if isinstance(value, (int, np.integer)) else tuple(value)
+        if not 1 <= len(values) <= 3 or any(v < 1 for v in values):
+            raise ValueError(f"{what} must be 1 to 3 positive integers, got {value!r}")
+        return values + (default,) * (3 - len(values))
+
     def __init__(
         self,
         source: str,
@@ -136,18 +143,6 @@ def __init__(
         self.init_value = init_value
         self._kernel: Any = None
 
-    def __repr__(self) -> str:
-        return (
-            f"MetalKernel({self.name!r}, inputs={self.inputs}, outputs={self.outputs})"
-        )
-
-    @staticmethod
-    def _as_triple(value: int | Sequence[int], what: str, default: int) -> tuple:
-        values = (int(value),) if isinstance(value, (int, np.integer)) else tuple(value)
-        if not 1 <= len(values) <= 3 or any(v < 1 for v in values):
-            raise ValueError(f"{what} must be 1 to 3 positive integers, got {value!r}")
-        return values + (default,) * (3 - len(values))
-
     @staticmethod
     def _as_input(value: Any) -> np.ndarray:
         """Return `value` as a NumPy array; Python scalars become 1-element arrays."""
@@ -277,3 +272,8 @@ def __call__(
                     nbytes=_nbytes(host),
                 )
         return outs[0] if len(outs) == 1 else outs
+
+    def __repr__(self) -> str:
+        return (
+            f"MetalKernel({self.name!r}, inputs={self.inputs}, outputs={self.outputs})"
+        )
diff --git a/src/cunumpy/_mirror.py b/src/cunumpy/_mirror.py
index 746ded7..87ea95a 100644
--- a/src/cunumpy/_mirror.py
+++ b/src/cunumpy/_mirror.py
@@ -51,12 +51,6 @@ class DeviceMirror:
     True
     """
 
-    def __init__(self, host: np.ndarray) -> None:
-        self._host = self._check_host(host)
-        self._shape = host.shape
-        self._dtype = host.dtype
-        self._device: Any = None
-
     @staticmethod
     def _check_host(host: Any) -> np.ndarray:
         if not isinstance(host, np.ndarray):
@@ -66,11 +60,11 @@ def _check_host(host: Any) -> np.ndarray:
             )
         return host
 
-    def __repr__(self) -> str:
-        where = "host" if self._device is None else "device"
-        return (
-            f"DeviceMirror(shape={self._shape}, dtype={self._dtype}, buffer={where!r})"
-        )
+    def __init__(self, host: np.ndarray) -> None:
+        self._host = self._check_host(host)
+        self._shape = host.shape
+        self._dtype = host.dtype
+        self._device: Any = None
 
     @property
     def host(self) -> np.ndarray:
@@ -87,19 +81,6 @@ def dtype(self) -> np.dtype:
         """Dtype of the mirrored array."""
         return self._dtype
 
-    @property
-    def device(self) -> Any:
-        """The array kernels write into: CuPy (copied from the host on first use), or the host on NumPy."""
-        if not _cupy_backend():
-            return self._host
-        self._check_bound()
-        self._check_device()
-        if self._device is None:
-            # The only host -> device allocation; it goes through to_cupy so
-            # that it is visible wherever transfers are accounted for.
-            self._device = to_cupy(self._host)
-        return self._device
-
     def _check_device(self) -> None:
         if self._device is not None:
             import cupy as cp
@@ -120,6 +101,19 @@ def _check_bound(self) -> None:
                 "call rebind(host) with the new array",
             )
 
+    @property
+    def device(self) -> Any:
+        """The array kernels write into: CuPy (copied from the host on first use), or the host on NumPy."""
+        if not _cupy_backend():
+            return self._host
+        self._check_bound()
+        self._check_device()
+        if self._device is None:
+            # The only host -> device allocation; it goes through to_cupy so
+            # that it is visible wherever transfers are accounted for.
+            self._device = to_cupy(self._host)
+        return self._device
+
     def rebind(self, host: np.ndarray) -> DeviceMirror:
         """Bind to a new host array, e.g. after its owner reallocated it.
 
@@ -260,3 +254,9 @@ def zero(self) -> DeviceMirror:
         else:
             self._device.fill(0)
         return self
+
+    def __repr__(self) -> str:
+        where = "host" if self._device is None else "device"
+        return (
+            f"DeviceMirror(shape={self._shape}, dtype={self._dtype}, buffer={where!r})"
+        )
diff --git a/src/cunumpy/_pyccel_cache.py b/src/cunumpy/_pyccel_cache.py
index 563c2b4..5d3bc32 100644
--- a/src/cunumpy/_pyccel_cache.py
+++ b/src/cunumpy/_pyccel_cache.py
@@ -210,9 +210,6 @@ def __init__(
         )
         self.language = language
 
-    def __repr__(self) -> str:
-        return f"PyccelCache({self.name!r}, env_prefix={self.env_prefix!r})"
-
     @property
     def cache_env(self) -> str:
         """Name of the environment variable that moves the cache directory."""
@@ -371,3 +368,6 @@ def compile(self, module: ModuleType, *, language: str | None = None) -> ModuleT
     def __call__(self, module: ModuleType) -> ModuleType:
         """:meth:`compile`, so the instance can be a ``compile_host``."""
         return self.compile(module)
+
+    def __repr__(self) -> str:
+        return f"PyccelCache({self.name!r}, env_prefix={self.env_prefix!r})"
diff --git a/src/cunumpy/_random_streams.py b/src/cunumpy/_random_streams.py
index 07398d2..e6e7c7e 100644
--- a/src/cunumpy/_random_streams.py
+++ b/src/cunumpy/_random_streams.py
@@ -51,14 +51,6 @@ def __init__(self) -> None:
         self._bit_generator = "PCG64"
         self._generators: dict[str, Any] = {}
 
-    def __repr__(self) -> str:
-        if self._sequence is None:
-            return "RandomStreams(not seeded)"
-        return (
-            f"RandomStreams(entropy={self._sequence.entropy}, "
-            f"spawn_key={self._sequence.spawn_key}, bit_generator={self._bit_generator!r})"
-        )
-
     def seed(
         self,
         value: int | None,
@@ -277,6 +269,14 @@ def uniform(
             return rng.uniform(low=low, high=high, size=size)
         return low + (high - low) * rng.random(size=size)
 
+    def __repr__(self) -> str:
+        if self._sequence is None:
+            return "RandomStreams(not seeded)"
+        return (
+            f"RandomStreams(entropy={self._sequence.entropy}, "
+            f"spawn_key={self._sequence.spawn_key}, bit_generator={self._bit_generator!r})"
+        )
+
 
 random_streams = RandomStreams()
 """The random streams of this process."""
diff --git a/src/cunumpy/_scipy_backend.py b/src/cunumpy/_scipy_backend.py
index 027fe78..33514f6 100644
--- a/src/cunumpy/_scipy_backend.py
+++ b/src/cunumpy/_scipy_backend.py
@@ -94,9 +94,6 @@ def __init__(self, path: str = "") -> None:
         self._path = path
         self._children: dict[str, ScipyNamespace] = {}
 
-    def __repr__(self) -> str:
-        return f"<cunumpy scipy namespace {self._name('<backend>')}>"
-
     def _name(self, root: str) -> str:
         return f"{root}.{self._path}" if self._path else root
 
@@ -122,28 +119,6 @@ def resolve(self) -> ModuleType:
                 f"xp.scipy on the {backend} backend needs {name}: {_INSTALL[backend]}",
             ) from error
 
-    def __getattr__(self, name: str) -> Any:
-        if name.startswith("__"):
-            raise AttributeError(name)
-        child = f"{self._path}.{name}" if self._path else name
-        if child in SUBMODULES:
-            if name not in self._children:
-                self._children[name] = ScipyNamespace(child)
-            return self._children[name]
-        module = self.resolve()
-        try:
-            return getattr(module, name)
-        except AttributeError:
-            fill = self._fill(name)
-            if fill is not None:
-                return fill
-            other = "cupy" if get_backend() == "numpy" else "numpy"
-            raise AttributeError(
-                f"{module.__name__} has no attribute {name!r}: it is not available "
-                f"on the {get_backend()} backend (it may exist in "
-                f"{self._name(_ROOTS[other])})",
-            ) from None
-
     def _fill(self, name: str) -> Any:
         if get_backend() != "cupy":
             return None
@@ -179,6 +154,28 @@ def available(self, name: str) -> bool:
         except ImportError:
             return False
 
+    def __getattr__(self, name: str) -> Any:
+        if name.startswith("__"):
+            raise AttributeError(name)
+        child = f"{self._path}.{name}" if self._path else name
+        if child in SUBMODULES:
+            if name not in self._children:
+                self._children[name] = ScipyNamespace(child)
+            return self._children[name]
+        module = self.resolve()
+        try:
+            return getattr(module, name)
+        except AttributeError:
+            fill = self._fill(name)
+            if fill is not None:
+                return fill
+            other = "cupy" if get_backend() == "numpy" else "numpy"
+            raise AttributeError(
+                f"{module.__name__} has no attribute {name!r}: it is not available "
+                f"on the {get_backend()} backend (it may exist in "
+                f"{self._name(_ROOTS[other])})",
+            ) from None
+
     def __dir__(self) -> list[str]:
         prefix = f"{self._path}." if self._path else ""
         children = [
@@ -194,6 +191,9 @@ def __dir__(self) -> list[str]:
             names = [*names, *FILLS.get(self._path, {})]
         return sorted(set(names) | set(children) | {"available", "resolve"})
 
+    def __repr__(self) -> str:
+        return f"<cunumpy scipy namespace {self._name('<backend>')}>"
+
 
 scipy = ScipyNamespace()
 """SciPy of the active backend, available as ``xp.scipy``."""
diff --git a/src/cunumpy/_staging.py b/src/cunumpy/_staging.py
index 6fdfa36..b886abb 100644
--- a/src/cunumpy/_staging.py
+++ b/src/cunumpy/_staging.py
@@ -182,12 +182,6 @@ def __init__(
         self._stream: Any = None
         self._device_id: int | None = None
 
-    def __repr__(self) -> str:
-        return (
-            f"HostStaging(shape={self.shape}, dtype={self.dtype}, "
-            f"buffers={self._n_buffers})"
-        )
-
     @property
     def shape(self) -> tuple[int, ...]:
         """The fixed shape of each snapshot."""
@@ -321,6 +315,12 @@ def synchronize(self) -> None:
                 with slot.context():
                     slot.event.synchronize()
 
+    def __repr__(self) -> str:
+        return (
+            f"HostStaging(shape={self.shape}, dtype={self.dtype}, "
+            f"buffers={self._n_buffers})"
+        )
+
 
 class HostCopy:
     """A device-to-host copy started by :func:`~cunumpy.to_host_async`.
diff --git a/src/cunumpy/_streams.py b/src/cunumpy/_streams.py
index ab6e245..74af35b 100644
--- a/src/cunumpy/_streams.py
+++ b/src/cunumpy/_streams.py
@@ -59,12 +59,6 @@ class HostStream:
     True
     """
 
-    def __enter__(self) -> Any:
-        return self
-
-    def __exit__(self, *exc: object) -> None:
-        pass
-
     @property
     def done(self) -> bool:
         """Whether all work on the stream has completed; always True."""
@@ -106,6 +100,12 @@ def wait_event(self, event: Any) -> None:
         if not isinstance(event, HostEvent):
             raise TypeError("a host stream requires a host event")
 
+    def __enter__(self) -> Any:
+        return self
+
+    def __exit__(self, *exc: object) -> None:
+        pass
+
 
 def create_stream(*, non_blocking: bool = True) -> Any:
     """Allocate a reusable CUDA stream, or a :class:`HostStream` on NumPy.
diff --git a/src/cunumpy/_transfers.py b/src/cunumpy/_transfers.py
index 25a5ee8..1165369 100644
--- a/src/cunumpy/_transfers.py
+++ b/src/cunumpy/_transfers.py
@@ -42,14 +42,6 @@
 )
 """Event kinds, in the order they are reported."""
 
-# The currently active counters, innermost last. Instrumented code checks
-# ``if _ACTIVE:`` before doing any work, so the overhead of an inactive counter
-# is a single truth test. The list object is shared with the instrumented
-# modules and mutated in place; never rebind it.
-_ACTIVE: list[TransferCounter] = []
-
-_PACKAGE_DIR = os.path.dirname(os.path.abspath(__file__))
-
 
 @dataclass(frozen=True)
 class TransferEvent:
@@ -125,10 +117,6 @@ class TransferCounter:
     def __init__(self) -> None:
         self.events: list[TransferEvent] = []
 
-    def __repr__(self) -> str:
-        counts = ", ".join(f"{kind}={self.count(kind)}" for kind in KINDS)
-        return f"TransferCounter({counts})"
-
     def _add(self, event: TransferEvent) -> None:
         self.events.append(event)
 
@@ -237,6 +225,19 @@ def report(self) -> str:
                 lines.append(f"    {where}: {description}{times}")
         return "\n".join(lines)
 
+    def __repr__(self) -> str:
+        counts = ", ".join(f"{kind}={self.count(kind)}" for kind in KINDS)
+        return f"TransferCounter({counts})"
+
+
+# The currently active counters, innermost last. Instrumented code checks
+# ``if _ACTIVE:`` before doing any work, so the overhead of an inactive counter
+# is a single truth test. The list object is shared with the instrumented
+# modules and mutated in place; never rebind it.
+_ACTIVE: list[TransferCounter] = []
+
+_PACKAGE_DIR = os.path.dirname(os.path.abspath(__file__))
+
 
 def _caller() -> str:
     """Return the innermost ``file:line`` on the stack outside cunumpy."""
@@ -486,10 +487,6 @@ def __init__(self, *, started: bool = True) -> None:
         self._stack: list[str] = []
         self._router = _PhaseRouter(self)
 
-    def __getitem__(self, phase: str) -> TransferCounter:
-        """Return the counter of `phase` (empty if it has not run)."""
-        return self.phases.setdefault(phase, TransferCounter())
-
     def start(self) -> None:
         """Count the phases from now on."""
         self.started = True
@@ -665,3 +662,7 @@ def check(self) -> None:
         """
         if self.violations():
             raise AssertionError("transfer budget exceeded:\n" + self.report())
+
+    def __getitem__(self, phase: str) -> TransferCounter:
+        """Return the counter of `phase` (empty if it has not run)."""
+        return self.phases.setdefault(phase, TransferCounter())
diff --git a/src/cunumpy/kernel_testing.py b/src/cunumpy/kernel_testing.py
index b3d2f61..bae6923 100644
--- a/src/cunumpy/kernel_testing.py
+++ b/src/cunumpy/kernel_testing.py
@@ -215,6 +215,16 @@ def cuda_required() -> bool:
     return os.environ.get("CUNUMPY_REQUIRE_CUDA", "").lower() in {"1", "true", "yes"}
 
 
+def _pytest() -> Any:
+    try:
+        import pytest
+    except ImportError:  # pragma: no cover - pytest is installed in the test suite
+        raise ImportError(
+            "cunumpy.kernel_testing needs pytest for this feature: pip install pytest",
+        ) from None
+    return pytest
+
+
 def _skip_or_fail(reason: str) -> None:
     """``pytest.skip``, or ``pytest.fail`` if a GPU is required."""
     if cuda_required():
@@ -529,16 +539,6 @@ def run_in_fake_cupy_subprocess(
 _LAZY: dict[str, Any] = {}
 
 
-def _pytest() -> Any:
-    try:
-        import pytest
-    except ImportError:  # pragma: no cover - pytest is installed in the test suite
-        raise ImportError(
-            "cunumpy.kernel_testing needs pytest for this feature: pip install pytest",
-        ) from None
-    return pytest
-
-
 def _build_lazy() -> None:
     pytest = _pytest()
     if cuda_required() and not _can_launch():
diff --git a/src/cunumpy/xp.py b/src/cunumpy/xp.py
index d348f2c..739c7f8 100644
--- a/src/cunumpy/xp.py
+++ b/src/cunumpy/xp.py
@@ -159,21 +159,6 @@ class ArrayBackend:
         If `backend` is neither ``"numpy"`` nor ``"cupy"``.
     """
 
-    def __init__(
-        self,
-        backend: BackendType = "numpy",
-        verbose: bool = False,
-    ) -> None:
-        if backend.lower() not in ("numpy", "cupy"):
-            raise ValueError("Array backend must be either 'numpy' or 'cupy'.")
-
-        self._backend: BackendType = "cupy" if backend.lower() == "cupy" else "numpy"
-        self._xp: ModuleType = np  # Placeholder
-        self._listeners: list[Callable[[ModuleType], None]] = []
-
-        # Import numpy/cupy
-        self._xp = self._load_backend(self._backend, verbose)
-
     def _load_backend(self, backend: BackendType, verbose: bool = False) -> ModuleType:
         if backend == "cupy":
             if cupy_available():
@@ -191,8 +176,20 @@ def _load_backend(self, backend: BackendType, verbose: bool = False) -> ModuleTy
         self._backend = "numpy"
         return np
 
-    def __repr__(self) -> str:
-        return f"ArrayBackend(backend={self._backend!r}, module={self._xp.__name__!r})"
+    def __init__(
+        self,
+        backend: BackendType = "numpy",
+        verbose: bool = False,
+    ) -> None:
+        if backend.lower() not in ("numpy", "cupy"):
+            raise ValueError("Array backend must be either 'numpy' or 'cupy'.")
+
+        self._backend: BackendType = "cupy" if backend.lower() == "cupy" else "numpy"
+        self._xp: ModuleType = np  # Placeholder
+        self._listeners: list[Callable[[ModuleType], None]] = []
+
+        # Import numpy/cupy
+        self._xp = self._load_backend(self._backend, verbose)
 
     @property
     def backend(self) -> BackendType:
@@ -279,6 +276,9 @@ def use_backend(
         finally:
             self._set(old_backend, old_xp)
 
+    def __repr__(self) -> str:
+        return f"ArrayBackend(backend={self._backend!r}, module={self._xp.__name__!r})"
+
 
 array_backend = ArrayBackend(
     backend=(
@@ -366,6 +366,22 @@ def set_backend(backend: BackendType, *, strict: bool = False) -> None:
     array_backend.set(backend, strict=strict)
 
 
+def get_backend() -> BackendType:
+    """Return the name of the active process-wide backend.
+
+    Returns
+    -------
+    {"numpy", "cupy"}
+        The effective backend, after any fallback of :func:`set_backend`.
+
+    Examples
+    --------
+    >>> xp.get_backend()
+    'numpy'
+    """
+    return array_backend.backend
+
+
 def backend_info() -> dict[str, Any]:
     """Return JSON-compatible backend, dependency and CUDA diagnostics.
 
@@ -424,22 +440,6 @@ def backend_info() -> dict[str, Any]:
     return info
 
 
-def get_backend() -> BackendType:
-    """Return the name of the active process-wide backend.
-
-    Returns
-    -------
-    {"numpy", "cupy"}
-        The effective backend, after any fallback of :func:`set_backend`.
-
-    Examples
-    --------
-    >>> xp.get_backend()
-    'numpy'
-    """
-    return array_backend.backend
-
-
 def _cupy_backend() -> bool:
     """Check if the active global backend is CuPy."""
     return array_backend.backend == "cupy"
@@ -498,6 +498,29 @@ def synchronize() -> None:
             )
 
 
+def get_array_backend(array: Any) -> BackendType:
+    """Return the backend of `array`, independent of the active backend.
+
+    Parameters
+    ----------
+    array : object
+        Any object; only CuPy arrays count as ``"cupy"``.
+
+    Returns
+    -------
+    {"numpy", "cupy"}
+        ``"cupy"`` for a CuPy array, ``"numpy"`` for anything else.
+
+    Examples
+    --------
+    >>> xp.get_array_backend(np.zeros(3))
+    'numpy'
+    """
+    if type(array) is _NDARRAY:
+        return "numpy"
+    return "cupy" if array_api_compat.is_cupy_array(array) else "numpy"
+
+
 def _to_numpy(array: Any) -> np.ndarray:
     """`to_numpy` without transfer counting, for internal use."""
     if type(array) is _NDARRAY:
@@ -712,29 +735,6 @@ def to_cunumpy(array: Any) -> Any:
     return to_numpy(array)
 
 
-def get_array_backend(array: Any) -> BackendType:
-    """Return the backend of `array`, independent of the active backend.
-
-    Parameters
-    ----------
-    array : object
-        Any object; only CuPy arrays count as ``"cupy"``.
-
-    Returns
-    -------
-    {"numpy", "cupy"}
-        ``"cupy"`` for a CuPy array, ``"numpy"`` for anything else.
-
-    Examples
-    --------
-    >>> xp.get_array_backend(np.zeros(3))
-    'numpy'
-    """
-    if type(array) is _NDARRAY:
-        return "numpy"
-    return "cupy" if array_api_compat.is_cupy_array(array) else "numpy"
-
-
 def get_array_module(array: Any) -> ModuleType:
     """Return the ``array-api-compat`` module of `array`'s own backend.
 
diff --git a/testing/smoke_wheel.py b/testing/smoke_wheel.py
index 8ba93c6..c2e2bd3 100644
--- a/testing/smoke_wheel.py
+++ b/testing/smoke_wheel.py
@@ -14,6 +14,7 @@
 
 
 def main():
+    """Check that the installed wheel ships its assets and runs on the fake CuPy."""
     # Catch accidental imports from the checkout or another environment.
     assert Path(xp.__file__).resolve().is_relative_to(Path(sys.prefix).resolve())
     package = files("cunumpy")
diff --git a/tests/unit/test_cuda_kernel.py b/tests/unit/test_cuda_kernel.py
index 53bf69d..481783e 100644
--- a/tests/unit/test_cuda_kernel.py
+++ b/tests/unit/test_cuda_kernel.py
@@ -991,6 +991,50 @@ def test_struct_with_view_fields_on_gpu():
     assert np.array_equal(base.get(), expected)
 
 
+BOUNDS_TRAP = f"""
+import cupy as cp
+from cunumpy.kernels import CudaKernel
+
+kernel = CudaKernel({SCALE_COLUMN!r}, "scale_column", options=("-DCUNUMPY_BOUNDS_CHECK",))
+a = cp.ones((4, 3))
+try:
+    kernel(a, 5, 2.0, n_threads=4)  # column out of bounds
+    cp.cuda.Device().synchronize()
+except Exception as error:
+    print("trapped:", type(error).__name__, flush=True)
+else:
+    print("no error", flush=True)
+"""
+
+
+# ---------------------------------------------------------------------------
+# debug mode
+# ---------------------------------------------------------------------------
+
+
+def _run_python(code, env=None):
+    """Run `code` in a fresh interpreter that imports this cunumpy.
+
+    Returns its stdout and stderr."""
+    import os
+    import subprocess
+    import sys
+    from pathlib import Path
+
+    environment = {**os.environ, **(env or {})}
+    environment["PYTHONPATH"] = os.pathsep.join(
+        [str(Path(xp.__file__).parents[1]), environment.get("PYTHONPATH", "")],
+    )
+    result = subprocess.run(
+        [sys.executable, "-c", code],
+        capture_output=True,
+        text=True,
+        env=environment,
+        check=False,
+    )
+    return result.stdout, result.stderr
+
+
 @requires_bounds_check_views
 def test_bounds_check_traps_on_gpu():
     _skip_without_cupy()
@@ -1013,22 +1057,6 @@ def test_bounds_check_traps_on_gpu():
     assert "trapped" in output, output
 
 
-BOUNDS_TRAP = f"""
-import cupy as cp
-from cunumpy.kernels import CudaKernel
-
-kernel = CudaKernel({SCALE_COLUMN!r}, "scale_column", options=("-DCUNUMPY_BOUNDS_CHECK",))
-a = cp.ones((4, 3))
-try:
-    kernel(a, 5, 2.0, n_threads=4)  # column out of bounds
-    cp.cuda.Device().synchronize()
-except Exception as error:
-    print("trapped:", type(error).__name__, flush=True)
-else:
-    print("no error", flush=True)
-"""
-
-
 # ---------------------------------------------------------------------------
 # structs from pyccel annotations and header generation
 # ---------------------------------------------------------------------------
@@ -1392,34 +1420,6 @@ def test_shared_memory_on_gpu():
     assert float(out.sum()) == n * (n - 1) / 2
 
 
-# ---------------------------------------------------------------------------
-# debug mode
-# ---------------------------------------------------------------------------
-
-
-def _run_python(code, env=None):
-    """Run `code` in a fresh interpreter that imports this cunumpy.
-
-    Returns its stdout and stderr."""
-    import os
-    import subprocess
-    import sys
-    from pathlib import Path
-
-    environment = {**os.environ, **(env or {})}
-    environment["PYTHONPATH"] = os.pathsep.join(
-        [str(Path(xp.__file__).parents[1]), environment.get("PYTHONPATH", "")],
-    )
-    result = subprocess.run(
-        [sys.executable, "-c", code],
-        capture_output=True,
-        text=True,
-        env=environment,
-        check=False,
-    )
-    return result.stdout, result.stderr
-
-
 @pytest.fixture
 def debug_off():
     """Global debug mode off during the test, restored afterwards."""
diff --git a/tests/unit/test_mpi_serial.py b/tests/unit/test_mpi_serial.py
index f616e16..d134b65 100644
--- a/tests/unit/test_mpi_serial.py
+++ b/tests/unit/test_mpi_serial.py
@@ -59,13 +59,13 @@ def get(self):
     def reshape(self, *shape):
         return _DeviceArray(self._data.reshape(*shape))
 
-    def __setitem__(self, key, value):
-        self._data[key] = value
-
     @property
     def size(self):
         return self._data.size
 
+    def __setitem__(self, key, value):
+        self._data[key] = value
+
 
 def test_serial_copies_between_host_and_device_are_counted():
     with xp.profiling.count_transfers() as counter:
diff --git a/tests/unit/test_staging.py b/tests/unit/test_staging.py
index cc0e613..b70f592 100644
--- a/tests/unit/test_staging.py
+++ b/tests/unit/test_staging.py
@@ -77,12 +77,6 @@ def __init__(self, name):
         self.count = 0
         self.device_id = 0
 
-    def __enter__(self):
-        return self
-
-    def __exit__(self, *exc):
-        pass
-
     def record(self):
         self.count += 1
         event = FakeEvent(f"{self.name}#{self.count}")
@@ -92,6 +86,12 @@ def record(self):
     def wait_event(self, event):
         DeviceArray.log.append(("wait", self.name, event.name))
 
+    def __enter__(self):
+        return self
+
+    def __exit__(self, *exc):
+        pass
+
 
 @pytest.fixture
 def fake_device(monkeypatch):
diff --git a/tests/unit/test_transfers.py b/tests/unit/test_transfers.py
index f72493e..c5b764f 100644
--- a/tests/unit/test_transfers.py
+++ b/tests/unit/test_transfers.py
@@ -154,6 +154,13 @@ def test_to_cunumpy_counts_the_direction_it_delegates_to(fake_device, monkeypatc
     assert counter.to_device == 1 and counter.to_host == 0
 
 
+def _current_line() -> int:
+    """Line number of the caller."""
+    import sys
+
+    return sys._getframe(1).f_lineno
+
+
 # ---------------------------------------------------------------------------
 # Call sites and report
 # ---------------------------------------------------------------------------
@@ -415,10 +422,3 @@ def test_assert_no_transfers_on_device_only_work():
         values = xp.arange(10, dtype=xp.float64)
         values = values * 2 + 1
         xp.synchronize()
-
-
-def _current_line() -> int:
-    """Line number of the caller."""
-    import sys
-
-    return sys._getframe(1).f_lineno