Version 0.3.0

.github/workflows/gpu_ci_trigger.yml

@@ -64,6 +64,7 @@ jobs:
 
           # 4. Force push (This automatically starts the GitLab Pipeline)
           git push -f gitlab HEAD:refs/heads/$TARGET_BRANCH
+          echo "PUSHED_SHA=$(git rev-parse HEAD)" >> $GITHUB_ENV
 
           # 5. Provide the direct link
           PIPELINE_URL="https://gitlab.mpcdf.mpg.de/maxlin/cunumpy/-/pipelines?ref=$TARGET_BRANCH"
@@ -72,10 +73,41 @@ jobs:
           echo "::notice::View Pipeline: $PIPELINE_URL"
 
       - name: Wait for GitLab Pipeline
-        uses: docker://gitlab/glab:latest
+        timeout-minutes: 90
         env:
           GITLAB_TOKEN: ${{ secrets.GITLAB_TOKEN }}
-          GITLAB_HOST: gitlab.mpcdf.mpg.de
-        with:
-          entrypoint: glab
-          args: ci status --live --branch ${{ env.TARGET_BRANCH }} --repo maxlin/cunumpy
+        run: |
+          API="https://gitlab.mpcdf.mpg.de/api/v4/projects/maxlin%2Fcunumpy"
+          AUTH=()
+          if [ -n "$GITLAB_TOKEN" ]; then AUTH=(--header "PRIVATE-TOKEN: $GITLAB_TOKEN"); fi
+
+          # 1. GitLab creates the pipeline asynchronously after the push: wait until
+          #    the pipeline for the pushed commit exists (asking right away finds none).
+          PIPELINE=""
+          for i in $(seq 1 60); do
+            PIPELINE=$(curl -sf "${AUTH[@]}" "$API/pipelines?ref=$TARGET_BRANCH&sha=$PUSHED_SHA&per_page=1" | jq -r '.[0].id // empty' || true)
+            if [ -n "$PIPELINE" ]; then break; fi
+            sleep 5
+          done
+          if [ -z "$PIPELINE" ]; then
+            echo "::error::No GitLab pipeline for $PUSHED_SHA on $TARGET_BRANCH after 5 minutes"
+            exit 1
+          fi
+          URL="https://gitlab.mpcdf.mpg.de/maxlin/cunumpy/-/pipelines/$PIPELINE"
+          echo "::notice::GitLab pipeline: $URL"
+
+          # 2. Follow the pipeline until it has finished.
+          while true; do
+            STATUS=$(curl -sf "${AUTH[@]}" "$API/pipelines/$PIPELINE" | jq -r '.status // empty' || true)
+            case "$STATUS" in
+              success)
+                echo "GitLab pipeline $PIPELINE succeeded: $URL"
+                exit 0 ;;
+              failed|canceled|skipped|manual|scheduled)
+                echo "::error::GitLab pipeline $PIPELINE is $STATUS: $URL"
+                exit 1 ;;
+              *)
+                echo "GitLab pipeline $PIPELINE: ${STATUS:-status not available yet}"
+                sleep 20 ;;
+            esac
+          done

.github/workflows/testing.yml

@@ -17,7 +17,7 @@ jobs:
     strategy:
       fail-fast: false
       matrix:
-        python-version: ["3.8", "3.10", "3.13"]
+        python-version: ["3.10", "3.11", "3.12", "3.13", "3.14"]
 
     steps:
       # Checkout the repository

.gitlab-ci.yml

@@ -29,18 +29,19 @@ gpu_tests:
     - git --version
 
     - echo "--- Pytest Execution ---"
-    # The MPCDF image likely has a specific python environment. 
+    # The MPCDF image likely has a specific python environment.
     # We install our dependencies into the user directory or a virtualenv.
-    - python3 -m pip install --user cupy-cuda12x
-    - python3 -m pip install --user nvidia-cublas-cu12 nvidia-cufft-cu12 nvidia-curand-cu12 nvidia-cusolver-cu12 nvidia-cusparse-cu12
+    # One CUDA version only: nvhpcsdk/26 provides CUDA 13.2 (its headers are used
+    # when CuPy compiles kernels with NVRTC), so install CuPy for CUDA 13 with the
+    # CUDA 13.2 libraries and NVRTC. Mixing CUDA 12 (cupy-cuda12x) with the CUDA 13.2
+    # headers fails to compile CuPy's own kernels (e.g. CUB reductions).
+    - python3 -m pip install --user "cupy-cuda13x[ctk]" "cuda-toolkit==13.2.*"
     - python3 -m pip install --user -e .
-    
+
     # Add the user bin to PATH for pytest
     - export PATH="$HOME/.local/bin:$PATH"
-
-    # Try to find libcublas and other libraries in the HPC environment
-    - export LD_LIBRARY_PATH=$(find /mpcdf/soft /opt/nvidia -name libcublas.so.12 -exec dirname {} \; 2>/dev/null | head -n 1):$LD_LIBRARY_PATH
-
+
     - export ARRAY_BACKEND=cupy
+    - python3 -c "import cupy; cupy.show_config()"
 
     - pytest -xvs .

CHANGELOG.md

@@ -7,6 +7,30 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ## [Unreleased]
 
+### Removed
+- Support for Python 3.8 and 3.9 (both end-of-life); `cunumpy` now requires Python 3.10 or newer.
+
+### Changed
+- Python 3.14 is supported.
+- CI now tests every supported Python version (3.10, 3.11, 3.12, 3.13 and 3.14) instead of 3.8/3.10/3.13.
+
+### Added
+- `xp.CudaKernel`: Wraps a CUDA C kernel (`cupy.RawKernel`, compiled lazily with NVRTC) so it can be called with the same arguments as the host kernel it mirrors, plus `n_threads`. The `extern "C" __global__` signature is parsed once and every call is checked against it: argument count, array dtypes (host arrays raise, they are never copied), and scalars (Python scalars are cast to the declared C types with range checks; lossy or mismatching scalars raise instead of reaching the kernel as silently wrong values). Supports `block_size`, NVRTC `options`, `include_dirs`, `shared_mem`, `stream`, `CudaKernel.from_file()` (`<name>_cuda.cu`), `compile()` and `prepare_args()`; `check_signature=False` skips the checks.
+- `xp.CudaArguments`: Base class for argument objects that are flattened into several CUDA kernel arguments; any object with a `__cuda_args__()` method is flattened.
+- `xp.parse_cuda_signature(source, name)` and `xp.CudaParameter`: Parse the parameters of a `__global__` function.
+- `xp.Kernel`: A host kernel (`PyccelKernel`) and its CUDA counterpart, calling the one matching the active backend. Without a CUDA kernel on the CuPy backend it raises `NotImplementedError` (`missing_cuda="raise"`, default) or falls back to the host kernel with host copies (`missing_cuda="fallback"`).
+- `xp.KernelCatalog`: Read-only mapping of `Kernel`s; `KernelCatalog.from_package()` collects them from a package with one folder per kernel (`name/name_kernels.py`, `name/name_cuda.cu`); `without_cuda` lists the kernels still to port.
+- `xp.CudaStruct` and `xp.CudaStructValue`: C structs passed to CUDA kernels by value. A `CudaStruct` is defined once from `(field, C type)` pairs; it provides the C `declaration`, the NumPy `dtype` with the C memory layout, and packs values (device arrays as addresses, scalars checked and cast) into a `CudaStructValue` that is passed as one kernel argument. `CudaKernel(..., structs=[...])` checks struct parameters and that a struct definition in the source matches.
+- `CudaKernel(..., template_args=...)`: Instantiate C++ function templates (e.g. `template_args=(np.float64, 3)` for `name<double, 3>`); the template parameters are substituted into the checked signature.
+- `xp.CudaKernelVariants`: Creates and caches one `CudaKernel` per variant key for generated kernel sources (e.g. per dimension and dtype); `compile_all()` compiles given and existing variants.
+- `xp.ctype_of(dtype)`: The C type of a NumPy dtype, e.g. for generating CUDA source.
+- 1D to 3D launches: `CudaKernel` accepts a tuple `block_size`, and calls take `n_threads` as an integer or tuple, or an explicit `grid`, plus a per-call `block`; `launch_shape()` returns the `(grid, block)` of a call. Dynamic shared memory (`shared_mem`) and `stream` are passed through by `Kernel` as well.
+- Compiling at setup: `CudaKernel.compile()` and `is_compiled`, `Kernel.compile()`, and `KernelCatalog.compile_all()`.
+- `Kernel(..., host_options=...)` and `KernelCatalog.from_package(..., host_options=...)`: `PyccelKernel` options (e.g. `object_modules`, `outputs`) for the host kernels, for all kernels or per kernel name; needed for the fallback to find device arrays inside application objects.
+- `xp.local_rank()`: The node-local rank from the MPI launcher's environment (Open MPI, MVAPICH2, Intel MPI/MPICH, PMI, Cray PALS, Slurm, `LOCAL_RANK`), available before `MPI_Init`.
+- `xp.bind_local_device()`: Selects the GPU `local_rank() % device_count()` and creates its context, before `MPI_Init`, for one-rank-per-GPU MPI programs.
+- `xp.synchronize_for_mpi(*arrays)`: Waits for pending work on the current stream before MPI uses device buffers (no-op for host buffers and on the NumPy backend).
+
 ## [0.2.0] - 2026-09-28
 
 ### Changed

README.md

@@ -159,6 +159,22 @@ active CuPy device and `None` on NumPy. `set_device_for_rank(rank)` is a
 round-robin convenience for MPI layouts where local ranks map contiguously to
 GPUs. If your scheduler uses a different mapping, select the device directly.
 
+For MPI programs with one rank per GPU, `bind_local_device()` selects the GPU
+from the node-local rank that the MPI launcher exports (`local_rank()`), so it
+can run before MPI is initialized, as CUDA-aware MPI requires. Before passing
+device buffers to MPI, call `synchronize_for_mpi(*buffers)`: kernels run
+asynchronously, and MPI would otherwise send a buffer a kernel is still
+writing, without an error.
+
+```python
+xp.set_backend("cupy")
+xp.bind_local_device()  # before MPI_Init
+from mpi4py import MPI
+
+xp.synchronize_for_mpi(send, recv)
+MPI.COMM_WORLD.Sendrecv(send, dest, recvbuf=recv, source=source)
+```
+
 CuPy caches released allocations in memory pools. This can make process-level
 GPU memory appear occupied after arrays go out of scope. `free_memory()` asks
 CuPy to release currently free cached blocks; it does not free memory still
@@ -213,6 +229,80 @@ The wrapper can also traverse arrays nested in lists, tuples, dictionaries,
 and selected application objects; see the full [API reference](docs/source/api.md)
 for `object_modules`, `is_array`, aliasing, and output declarations.
 
+## Write CUDA kernels next to host kernels
+
+`CudaKernel` wraps a CUDA C kernel (compiled with NVRTC through
+`cupy.RawKernel`) so that it is called with the same arguments as the host
+kernel it mirrors, plus the number of threads. Arrays are never copied: they
+must be CuPy arrays. The `extern "C" __global__` signature is parsed once and
+every call is checked against it: Python scalars are cast to the declared C
+types, and a wrong argument count, an array of the wrong dtype, or a scalar
+that does not fit its type raises instead of silently producing wrong values.
+
+`Kernel` pairs a host kernel with its CUDA kernel and calls the one matching
+the active backend, so kernels can be ported to CUDA one at a time:
+
+```python
+import cunumpy as xp
+
+AXPY = r"""
+extern "C" __global__
+void axpy(double a, const double* x, double* y, int n) {
+    int i = blockDim.x * blockIdx.x + threadIdx.x;
+    if (i < n) y[i] += a * x[i];
+}
+"""
+
+
+def axpy(a, x, y, n):  # host version, e.g. compiled with Pyccel
+    for i in range(n):
+        y[i] += a * x[i]
+
+
+kernel = xp.Kernel(axpy, xp.CudaKernel(AXPY, "axpy"))
+
+with xp.use_backend("cupy"):
+    x = xp.arange(1000, dtype=xp.float64)
+    y = xp.zeros(1000)
+    kernel(2.0, x, y, 1000, n_threads=1000)  # runs the CUDA kernel
+```
+
+On the CuPy backend, a `Kernel` without CUDA kernel raises
+`NotImplementedError` (or, with `missing_cuda="fallback"`, runs the host kernel
+through `PyccelKernel`, with host copies; `host_options` configure that
+`PyccelKernel`). `KernelCatalog.from_package()` collects kernel pairs from a
+package with one folder per kernel (`name/name_kernels.py` and
+`name/name_cuda.cu`), and `catalog.compile_all()` compiles all CUDA kernels at
+setup.
+
+Groups of arguments can be passed as one: objects implementing
+`__cuda_args__()` (see `CudaArguments`) are flattened into several kernel
+arguments, and `CudaStruct` defines a C struct once (its C `declaration` and
+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.declaration
+    + r"""
+    extern "C" __global__ void scale(Vec v, double a) {
+        int i = blockDim.x * blockIdx.x + threadIdx.x;
+        if (i < v.n) v.data[i] *= a;
+    }""",
+    "scale",
+    structs=[Vec],
+)
+scale(Vec(data=y, n=y.size), 0.5, n_threads=y.size)
+```
+
+Launches can be 1D to 3D (`n_threads=(nx, ny)`, `block_size=(16, 16)`) or use
+an explicit `grid`, with dynamic shared memory (`shared_mem`) and a `stream`.
+C++ function templates are instantiated with `template_args`, and
+`CudaKernelVariants` caches kernels whose source is generated per variant
+(e.g. per dimension and dtype). See the [API reference](docs/source/api.md) for
+details.
+
 ## Pyodide
 
 CuNumpy supports the NumPy backend in Pyodide. It does not provide CuPy/CUDA