Skip to content

Submodule namespaces #60

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
max-models merged 2 commits into from
Oct 3, 2026
Merged

Submodule namespaces #60

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
1b5645a
Refactor: Moved helper methods to modules
max-models Oct 3, 2026
a1f2bcd
formatting
max-models Oct 3, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,11 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0

## [Unreleased]

### Changed (breaking, with deprecation)
- The helpers moved from the top level of `cunumpy` to submodules, so that the top level is the NumPy/CuPy namespace plus backend selection and array conversion, and no helper hides a NumPy or CuPy name (`xp.fuse` hid `cupy.fuse`): `cunumpy.cuda` (CUDA only: `CudaKernel`, `CudaKernelVariants`, `CudaStruct*`, `CudaArguments`, `CudaParameter`, header tools, debug mode, device selection and memory, `stream`, `pin_memory`), `cunumpy.kernels` (`Kernel`, `KernelCatalog`, `PyccelKernel`, `KernelArguments`, `PyccelStructArguments`, host implementations, `as_kernel_array`, `kernel_output`, `fuse`), `cunumpy.rng` (`random_streams`, `RandomStreams`, `get_rng`, `philox_*`), `cunumpy.algorithms` (`morton_*`, `sort_by_key`, `segment_sum`), `cunumpy.mpi` (`mpi_buffer`, CUDA-aware MPI, `local_rank`, `synchronize_for_mpi`), `cunumpy.profiling` (`timed_region`, `Timing`, `nvtx_range`, transfer counting), `cunumpy.memory` (`HostStaging`, `StagedCopy`, `DeviceMirror`) and `cunumpy.petsc` (`petsc_vec`). All are imported by `import cunumpy`. The old top-level names still work and raise a `DeprecationWarning` naming the new place; they will be removed in 0.6.
- `cunumpy.testing` is now `cunumpy.kernel_testing`. Once imported, `cunumpy.testing` replaced NumPy's `xp.testing`, so `xp.testing.assert_allclose` failed in every test that ran after an `import cunumpy.testing`. `cunumpy.testing` still works, with a `DeprecationWarning`, until 0.6.
- Importing `cunumpy.cuda` makes `xp.cuda` the cunumpy submodule instead of CuPy's `cupy.cuda`; use `import cupy` for the latter.

### Fixed
- `CudaKernel`'s header hash (`-DCUNUMPY_INCLUDE_HASH`) now covers the headers shipped with cunumpy (`cunumpy/atomic.cuh`, `reduce.cuh`, ...), also when included in angle brackets. Before, an upgrade of cunumpy that changed one of them left CuPy's kernel cache serving the kernel compiled with the old header. `resolve_includes(..., angle_dirs=...)` tracks angle-bracket includes found in the given directories.
- `xp.testing.assert_kernels_agree` reads `CudaStructArguments` objects and struct values through their struct fields, so their arrays get the same names as the attributes of the host argument object (before, arrays behind properties were named after the private attribute holding the owner, and the comparison failed with "do not have the same array arguments").
Expand Down
64 changes: 41 additions & 23 deletions README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,24 @@ move existing arrays just because the selected backend changes. This guide
covers backend selection, array movement, mixed CPU/GPU workflows, and the
helper APIs CuNumpy provides around NumPy and CuPy.

The top level of `cunumpy` is the NumPy (or CuPy) namespace plus backend
selection and array conversion. The helpers are in submodules, so that they
never hide a NumPy name:

| Submodule | Contents |
|---|---|
| `xp.cuda` | CUDA only: `CudaKernel`, `CudaStruct`, CUDA headers, devices, streams |
| `xp.kernels` | `Kernel`, `KernelCatalog`, `PyccelKernel`, host implementations, `fuse` |
| `xp.rng` | `random_streams`, `get_rng`, `philox_*` |
| `xp.algorithms` | `morton_*`, `sort_by_key`, `segment_sum` |
| `xp.mpi` | `mpi_buffer`, CUDA-aware MPI |
| `xp.profiling` | `timed_region`, `nvtx_range`, `count_transfers` |
| `xp.memory` | `HostStaging`, `DeviceMirror` |
| `xp.petsc` | `petsc_vec` |
| `cunumpy.kernel_testing` | pytest helpers for host/CUDA kernel pairs |

Everything except `xp.cuda` works on both backends.

## Install

```bash
Expand Down Expand Up @@ -132,7 +150,7 @@ anything was counted. Only transfers made through CuNumpy are seen; raw
`cupy.ndarray.get()` or `cupy.asarray()` calls need a profiler such as `nsys`.

```python
with xp.count_transfers() as counter:
with xp.profiling.count_transfers() as counter:
propagator(dt)

assert counter.total == 0, counter.report()
Expand All @@ -145,7 +163,7 @@ CuPy have similar generator APIs, though exact bit-for-bit sequences are not
guaranteed to match between libraries:

```python
rng = xp.get_rng(seed=42)
rng = xp.rng.get_rng(seed=42)
samples = rng.normal(size=1000)
```

Expand All @@ -162,9 +180,9 @@ These helpers are useful for multi-GPU programs and for understanding CuPy's
memory behavior:

```python
print("visible GPUs:", xp.device_count())
xp.set_device(0) # selects CUDA device 0 when CuPy is active
print("memory (free, total):", xp.memory_info())
print("visible GPUs:", xp.cuda.device_count())
xp.cuda.set_device(0) # selects CUDA device 0 when CuPy is active
print("memory (free, total):", xp.cuda.memory_info())
```

`set_device()` is a no-op on NumPy. `device_count()` checks visible CUDA
Expand Down Expand Up @@ -193,12 +211,12 @@ For MPI programs with one rank per GPU, the startup sequence is:

```python
xp.set_backend("cupy")
xp.bind_local_device() # before MPI_Init
xp.cuda.bind_local_device() # before MPI_Init
from mpi4py import MPI # MPI_Init

xp.require_cuda_aware_mpi() # once, on all ranks
xp.mpi.require_cuda_aware_mpi() # once, on all ranks

xp.synchronize_for_mpi(send, recv)
xp.mpi.synchronize_for_mpi(send, recv)
MPI.COMM_WORLD.Sendrecv(send, dest, recvbuf=recv, source=source)
```

Expand All @@ -214,7 +232,7 @@ on NumPy. GPU work is asynchronous, so synchronize before reading results on
the host:

```python
with xp.stream():
with xp.cuda.stream():
device = xp.to_cupy(host)
transformed = xp.fft.fft(device)

Expand All @@ -229,12 +247,12 @@ device before reading the clock (on NumPy it is a plain timer), and
no-ops or plain timers on NumPy, and `nvtx_range` also works as a decorator:

```python
with xp.timed_region("fft") as timing:
with xp.profiling.timed_region("fft") as timing:
transformed = xp.fft.fft(device)
print(timing.elapsed, timing.synced)


@xp.nvtx_range("step")
@xp.profiling.nvtx_range("step")
def step(dt):
...
```
Expand All @@ -256,7 +274,7 @@ def scale_in_place(values, factor):
return values


scale = xp.PyccelKernel(scale_in_place, outputs=(0,))
scale = xp.kernels.PyccelKernel(scale_in_place, outputs=(0,))

with xp.use_backend("cupy"):
values = xp.arange(5, dtype=xp.float64)
Expand Down Expand Up @@ -304,7 +322,7 @@ def axpy(a, x, y, n): # host version, e.g. compiled with Pyccel
y[i] += a * x[i]


kernel = xp.Kernel(axpy, xp.CudaKernel(AXPY, "axpy"))
kernel = xp.kernels.Kernel(axpy, xp.cuda.CudaKernel(AXPY, "axpy"))

with xp.use_backend("cupy"):
x = xp.arange(1000, dtype=xp.float64)
Expand All @@ -327,8 +345,8 @@ the matching memory layout) and packs values into it, which the kernel takes
as one parameter:

```python
Vec = xp.CudaStruct("Vec", [("data", "double*"), ("n", "int")])
scale = xp.CudaKernel(
Vec = xp.cuda.CudaStruct("Vec", [("data", "double*"), ("n", "int")])
scale = xp.cuda.CudaKernel(
Vec.declaration
+ r"""
extern "C" __global__ void scale(Vec v, double a) {
Expand Down Expand Up @@ -357,7 +375,7 @@ on the CUDA path, so the call site is the same on both backends and each form
can be built lazily on first access (a CPU run never builds device arguments):

```python
class ParticleArguments(xp.KernelArguments):
class ParticleArguments(xp.kernels.KernelArguments):
def __init__(self, markers):
self.markers = markers
self._host = None
Expand Down Expand Up @@ -387,9 +405,9 @@ class MarkerArguments:
def __init__(self, markers: "float[:, :]", n_markers: int, valid: "bool[:]"):
...

MarkerArgs = xp.CudaStruct.from_signature(MarkerArguments.__init__, "MarkerArgs")
MarkerArgs = xp.cuda.CudaStruct.from_signature(MarkerArguments.__init__, "MarkerArgs")
MarkerArgs.to_header("marker_args.cuh") # Array2D<double> markers; long long n_markers; ...
push = xp.CudaKernel(
push = xp.cuda.CudaKernel(
r"""
#include "marker_args.cuh"
#include <cunumpy/index.cuh>
Expand All @@ -414,8 +432,8 @@ details.

Kernels run asynchronously, so a CUDA error (an illegal memory access, say)
normally surfaces at a later `.get()` or MPI call, far from the kernel that
caused it. In debug mode, enabled with `xp.set_cuda_debug(True)`, the
context manager `xp.cuda_debug()`, `CudaKernel(..., debug=True)` or the
caused it. In debug mode, enabled with `xp.cuda.set_cuda_debug(True)`, the
context manager `xp.cuda.cuda_debug()`, `CudaKernel(..., debug=True)` or the
environment variable `CUNUMPY_CUDA_DEBUG=1`, kernels are compiled with
`-lineinfo` and `-DCUNUMPY_BOUNDS_CHECK` and every launch is synchronized, so
the error is raised as a `RuntimeError` naming the kernel and its launch shape.
Expand All @@ -425,7 +443,7 @@ next step is NVIDIA's memory checker:

## Test kernel pairs

`cunumpy.testing` helps to test the ports with pytest. `assert_kernels_agree`
`cunumpy.kernel_testing` helps to test the ports with pytest. `assert_kernels_agree`
builds the arguments on both backends, runs the host and the CUDA kernel and
compares the arrays they wrote; with `catalog.parity_cases()`, one
parametrised test covers every ported kernel of a catalog. `BACKENDS` and
Expand All @@ -436,7 +454,7 @@ kernel:

```python
import pytest
from cunumpy.testing import assert_kernels_agree
from cunumpy.kernel_testing import assert_kernels_agree


def make_args(backend, seed):
Expand All @@ -460,7 +478,7 @@ the one transfer per accumulation is explicit. The shipped header
writes:

```python
mirror = xp.DeviceMirror(vector._data)
mirror = xp.memory.DeviceMirror(vector._data)
mirror.zero()
accumulate(markers, mirror.device, n_threads=n_markers)
mirror.to_host() # vector._data holds the result on both backends
Expand Down
Loading
Loading