Use CUNUMPY_BACKEND for environment variable and calculate n_threads …

.gitlab-ci.yml

@@ -1,7 +1,7 @@
 variables:
   CUDA_IMAGE: "gitlab-registry.mpcdf.mpg.de/mpcdf/ci-module-image/nvhpcsdk_26:2026"
   PIP_DISABLE_PIP_VERSION_CHECK: "1"
-  ARRAY_BACKEND: "cupy"
+  CUNUMPY_BACKEND: "cupy"
   CUNUMPY_REQUIRE_CUDA: "1"
 
 stages:
@@ -42,6 +42,7 @@ gpu_sanitizers:
       compute-sanitizer --tool "$SANITIZER" --error-exitcode 1
       --target-processes all python3 -m pytest -q
       tests/unit/test_cuda_collectives.py
+      tests/unit/test_automatic_launch.py
       tests/unit/test_cuda_launch_limits.py
       tests/unit/test_reusable_streams.py
       tests/unit/test_staging.py

CHANGELOG.md

@@ -8,6 +8,9 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 ## [Unreleased]
 
 ### Added
+- CUDA launches infer thread counts from the first array by default, including
+  arrays in argument objects. 1D blocks use rows; multidimensional blocks use
+  matching leading shape axes. Explicit sizes and callbacks override inference.
 - Per-device eager CUDA compilation, compiler log streams and explicit `recompile()`.
 - Launch validation against actual device/kernel dimensions, thread limits and
   static plus dynamic shared memory, with cached per-device opt-in state.
@@ -57,6 +60,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 - Support for Python 3.8 and 3.9 (both end-of-life); `cunumpy` now requires Python 3.10 or newer.
 
 ### Changed
+- Renamed the startup backend selector from `ARRAY_BACKEND` to `CUNUMPY_BACKEND`;
+  update batch-job environments. CuNumpy-owned environment options use `CUNUMPY_*`.
 - `Kernel(..., dispatch="arrays")` calls the host kernel directly for host arguments, without `PyccelKernel`'s check for device arrays to convert, which converted nothing but ran while CuPy was the active backend (unless the `PyccelKernel` was built with `use_cupy=True`).
 - `Kernel.check_signature()` (and `KernelCatalog.check_signatures()`) also compares the parameters of every host implementation (numba, NumPy; nothing is compiled) and of a `CompiledHostKernel`'s fallback with those of the host kernel.
 - `KernelCatalog.from_package(..., compile_host=..., host_fallback=...)` builds `HostImplementations` instead of `CompiledHostKernel`s: `host_fallback` becomes the `"numpy"` implementation, and `<name>_numba.py`/`<name>_numpy.py` files in a kernel folder are picked up.
@@ -75,7 +80,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 - `cunumpy/morton.cuh` and `xp.morton_keys`, `morton_encode`, `morton_decode`, `morton_scales`, `MAX_MORTON_LEVELS`: Morton (Z-order) keys of 2D and 3D points (`uint64`, up to 32 and 21 bits per axis), the same in a kernel and on the host (bit for bit), for sorting particles along a space-filling curve and building quadtrees and octrees from sorted keys.
 - `xp.sort_by_key(keys, *arrays)`: one stable argsort of `keys` applied to every array; returns the sorted keys, the order and the sorted arrays.
 - `Kernel.from_folder(package, ...)`: the kernel of one kernel folder, with the options of `KernelCatalog.from_package` (`host_suffix`, `compile_host`, `dispatch`, `include_dirs`, CUDA options...). Every version in the folder is an implementation: `<name><host_suffix>.py` (`"pyccel"`, compiled with `compile_host`, and `"python"`, uncompiled), `<name>_numba.py`, `<name>_numpy.py` and `<name>_cuda.cu`. A folder's own `__init__.py` can declare `kernel = xp.Kernel.from_folder(__name__, ...)`, so code imports the kernel from where it is written; `from_package` now builds each kernel with it. `Kernel.implementations` lists them, `Kernel.selected(device=False)` tells which one a call runs.
-- `xp.HostImplementations` and `xp.HOST_IMPLEMENTATIONS`: the host implementations of a kernel, loaded on first use; a call runs the default, the first available of pyccel, numba and NumPy (the uncompiled Python version, with a warning, if none is), or the one chosen with `xp.set_kernel_implementation(name)` / `with xp.use_kernel_implementation(name):` / `CUNUMPY_KERNEL_IMPLEMENTATION=name` (read at import), like `set_backend`/`use_backend`/`ARRAY_BACKEND`; a chosen implementation that a kernel lacks or cannot load raises instead of running another. `xp.get_kernel_implementation()` reads the setting.
+- `xp.HostImplementations` and `xp.HOST_IMPLEMENTATIONS`: the host implementations of a kernel, loaded on first use; a call runs the default, the first available of pyccel, numba and NumPy (the uncompiled Python version, with a warning, if none is), or the one chosen with `xp.set_kernel_implementation(name)` / `with xp.use_kernel_implementation(name):` / `CUNUMPY_KERNEL_IMPLEMENTATION=name` (read at import), like `set_backend`/`use_backend`/`CUNUMPY_BACKEND`; a chosen implementation that a kernel lacks or cannot load raises instead of running another. `xp.get_kernel_implementation()` reads the setting.
 - `xp.as_kernel_array(value, like, dtype=None)` and `xp.kernel_output(out, like, dtype=None)`: bring the arguments of a `dispatch="arrays"` kernel to the side of the main array `like` (CuPy or NumPy, C-contiguous, `dtype`), without a copy when they already fit; `kernel_output` yields the buffer the kernel writes and copies it back into `out` if it had to be converted.
 - `CompiledHostKernel.fallback` returns the fallback.
 - `cunumpy/random.cuh` and `xp.philox_uniform`, `philox_uniform2`, `philox_normal`, `philox_normal2`, `philox4x32_10`: counter-based random numbers (Philox4x32-10, passing the Random123 known-answer tests) as a pure function of `(seed, stream, counter)`, the same in a kernel and on the host (uniform numbers bit for bit, normal numbers up to the last bits of the math functions), so kernels that draw random numbers can be compared with their host versions.

README.md

@@ -55,7 +55,7 @@ explanation and examples.
 
 ## Choose a backend
 
-CuNumpy starts with NumPy unless `ARRAY_BACKEND=cupy` is set before import.
+CuNumpy starts with NumPy unless `CUNUMPY_BACKEND=cupy` is set before import.
 You can also choose at runtime:
 
 ```python
@@ -303,7 +303,10 @@ for `object_modules`, `is_array`, aliasing, and output declarations.
 
 `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
+kernel it mirrors. Thread counts default to the first array's leading shape
+axes: one thread per row for 1D blocks, matching axes for 2D/3D blocks. Explicit
+`n_threads`, `grid`, or a custom `n_threads_from` controls the launch when needed.
+Arrays are never copied: they
 must be C-contiguous 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
@@ -335,7 +338,7 @@ kernel = xp.kernels.Kernel(axpy, xp.cuda.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
+    kernel(2.0, x, y, 1000)  # infer n_threads = x.shape[0], run the CUDA kernel
 ```
 
 On the CuPy backend, a `Kernel` without CUDA kernel raises

docs/source/api.md

@@ -93,7 +93,7 @@ availability and failure reason, dependency versions, and current device/runtime
 information when available. Inspection errors are included in the result. It
 does not change the backend or initialize MPI.
 
-The initial backend is NumPy unless `ARRAY_BACKEND=cupy` is set before CuNumpy
+The initial backend is NumPy unless `CUNUMPY_BACKEND=cupy` is set before CuNumpy
 is imported. Other values of this environment variable result in the NumPy
 default.
 
@@ -870,6 +870,8 @@ xp.cuda.CudaKernel(
     template_args=None,
     check_signature=True,
     debug=None,
+    n_threads_from="auto",
+    check_finite=False,
 )
 xp.cuda.CudaKernel.from_file(path, name=None, *, suffix="_cuda.cu", **kwargs)
 xp.cuda.CudaKernel.all_from_file(path, **kwargs)
@@ -991,20 +993,36 @@ shape is given either by `n_threads` or by `grid`:
   subtracting static storage opts in through `max_dynamic_shared_size_bytes`.
   Block/grid dimensions and device/kernel thread limits are also checked.
 
-Without `n_threads` and `grid`, a launch uses `n_threads_from(args)` if the
-kernel has one (constructor argument and settable property): a function of the
-argument tuple, or `"first_array"` for the length of the first array argument
-(its first axis), i.e. one thread per marker:
+Without `n_threads` and `grid`, the default `n_threads_from="auto"` infers the
+launch from the first array argument, skipping scalars and zero-dimensional
+arrays. Arrays inside `CudaArguments`, `CudaStructArguments`, and struct values
+are searched in argument/field order. The effective block shape determines
+how many leading array axes are used: a 1D block uses `shape[0]` (one thread per
+row/particle), a 2D block uses `shape[:2]`, and a 3D block uses `shape[:3]`. These
+axes map to CUDA x, y, z. A `block=` override changes the inference dimensions.
+Missing arrays or insufficient array dimensions raise with instructions to
+give an explicit size.
 
 ```python
-push = xp.cuda.CudaKernel.from_file("push_cuda.cu", n_threads_from="first_array")
+push = xp.cuda.CudaKernel.from_file("push_cuda.cu")
 push(positions, velocities, e_field, dt)  # n_threads = positions.shape[0]
 ```
 
+Explicit `n_threads` or `grid` always takes precedence. Set `n_threads_from`
+(constructor argument or settable property) to a callable such as
+`lambda args: args[0].size` for flattened element kernels, or
+`lambda args: args[0].shape[::-1]` for kernels whose x index follows columns.
+`"first_array"` always uses the first axis; None disables inference and requires
+explicit launch sizes. The same defaults apply through `kernels.Kernel` and
+`kernel_testing.assert_kernels_agree`, and in CPU emulation.
+
 Nothing is launched if the grid has a zero dimension (e.g. `n_threads=0`).
-`kernel.launch_shape(n_threads=None, *, grid=None, block=None)` returns the
+`kernel.launch_shape(n_threads=None, *, grid=None, block=None, args=None)` returns the
 `(grid, block)` a call would use, e.g. to size a per-block output:
 
+Supply `args=(...)` to inspect an automatically inferred launch without running
+the kernel, e.g. `kernel.launch_shape(args=(positions, velocities, e_field, dt))`.
+
 ```python
 BLOCK_SUM = r"""
 extern "C" __global__ void block_sum(const double* x, double* out, int n) {
@@ -1521,8 +1539,9 @@ kernel(*args, n_threads=None, grid=None, block=None, shared_mem=0, stream=None)
 ```
 
 calls the kernel of the active backend. The launch arguments are passed to the
-CUDA kernel (`n_threads` or `grid` is required there) and ignored by the host
-kernel. Arguments implementing `KernelArguments` are replaced by their
+CUDA kernel and ignored by the host kernel. Omitted sizes use the CUDA kernel's
+shape-based default or its configured callback. Arguments implementing
+`KernelArguments` are replaced by their
 `__host_args__()` on the host path and flattened via `__cuda_args__()` on the
 CUDA path. `kernel.compile()` compiles the CUDA kernel now and returns whether
 there is one.
@@ -1571,8 +1590,8 @@ Properties: `name`, `host_kernel`, `cuda_kernel`, `has_cuda`, `missing_cuda`,
 * `Kernel.host_parameters()` falls back to the `__pyccel__/<module>.pyi` stub
   for a pyccel-compiled host function, so `check_signature()` works for
   compiled kernels.
-* `Kernel.__call__` needs no `n_threads` when the CUDA kernel has
-  `n_threads_from`.
+* `Kernel.__call__` infers thread counts by default; custom `n_threads_from`
+  callbacks are supported, and None requires explicit sizes.
 
 ## `kernels.KernelCatalog`
 

docs/source/best-practices.md

@@ -7,7 +7,7 @@ A condensed checklist. Each item links to the guide with the reasoning.
 * Import as `import cunumpy as xp` and always call `xp.<name>(...)`; never
   `from cunumpy import <array function>`. ([Backend-agnostic
   code](guides/portable-code.md))
-* Choose the backend once, in the entry point, with `ARRAY_BACKEND` or
+* Choose the backend once, in the entry point, with `CUNUMPY_BACKEND` or
   `set_backend()`. Library code never calls `set_backend()`. ([Choosing a
   backend](guides/backends.md))
 * Read back `xp.get_backend()` after requesting CuPy; log it with

docs/source/guides/backends.md

@@ -6,17 +6,21 @@ where newly created arrays live.
 
 ## Select the backend at start-up
 
-The backend is NumPy unless the environment variable `ARRAY_BACKEND=cupy` is
+The backend is NumPy unless the environment variable `CUNUMPY_BACKEND=cupy` is
 set when CuNumpy is first imported:
 
 ```bash
-ARRAY_BACKEND=cupy python simulate.py
+CUNUMPY_BACKEND=cupy python simulate.py
 ```
 
 This is the least intrusive option for scripts and batch jobs: the code does
 not change, and a job script decides whether it runs on a GPU. The variable is
 read once, at import; setting it later in `os.environ` has no effect.
 
+The startup setting is named `CUNUMPY_BACKEND`; migrate existing job scripts
+from `ARRAY_BACKEND`. CuNumpy-owned environment options share the `CUNUMPY_`
+prefix; the complete list is in [Installation](../installation.md).
+
 To choose from inside the program, for example from a command-line flag, call
 `set_backend()` once, early, before arrays are created:
 

docs/source/guides/gpu-devices.md

@@ -30,7 +30,7 @@ The usual alternative is to restrict visibility from outside the process,
 which also works for libraries that do not know about CuNumpy:
 
 ```bash
-CUDA_VISIBLE_DEVICES=1 ARRAY_BACKEND=cupy python simulate.py
+CUDA_VISIBLE_DEVICES=1 CUNUMPY_BACKEND=cupy python simulate.py
 ```
 
 For MPI programs with one rank per GPU, use `bind_local_device()` instead

docs/source/installation.md

@@ -67,8 +67,16 @@ Tests that need a GPU are skipped automatically where CuPy is not functional.
 
 | Variable | Effect |
 | --- | --- |
-| `ARRAY_BACKEND=cupy` | start with the CuPy backend instead of NumPy (read once, at import) |
+| `CUNUMPY_BACKEND=cupy` | start with the CuPy backend instead of NumPy (read once, at import) |
 | `CUNUMPY_CUDA_DEBUG=1` | enable [CUDA debug mode](kernels/debugging.md) for all kernels |
+| `CUNUMPY_KERNEL_IMPLEMENTATION=numpy` | choose the host kernel implementation (read at import) |
+| `CUNUMPY_MPI=1` / `0` | require MPI / use serial MPI regardless of launcher detection |
+| `CUNUMPY_FAKE_CUPY=1` | install the strict CPU stand-in for CuPy for tests |
+| `CUNUMPY_REQUIRE_CUDA=1` | require a real usable GPU when starting the test suite (CI guard) |
+
+Use `CUNUMPY_BACKEND` in job scripts; the former `ARRAY_BACKEND` setting is no
+longer read. Standard toolchain/device variables such as `CXX` and
+`CUDA_VISIBLE_DEVICES` keep their standard meanings.
 
 MPI launchers also export node-local rank variables (`OMPI_COMM_WORLD_LOCAL_RANK`,
 `SLURM_LOCALID`, ...), which `xp.mpi.local_rank()` reads to pick a GPU per process.

docs/source/kernels/cuda-kernel.md

@@ -85,6 +85,24 @@ correct; the kernel then behaves like a raw `RawKernel`.
 kernel(*args, n_threads=None, grid=None, block=None, shared_mem=0, stream=None)
 ```
 
+Omit both `n_threads` and `grid` to infer the launch from the first array's shape.
+The default `n_threads_from="auto"` uses `shape[0]` for a 1D block (one thread per
+row/particle), `shape[:2]` for a 2D block, and `shape[:3]` for a 3D block. Axes map
+to CUDA x, y, z. Scalar arguments and zero-dimensional arrays are skipped;
+arrays inside supported argument objects are searched in field/argument order.
+Per-call block overrides determine the inference dimensions too.
+
+```python
+axpy(2.0, x, y, x.size)  # infer n_threads = x.shape[0]
+```
+
+Explicit sizes override inference. Set `n_threads_from=lambda args: args[0].size`
+for a flattened element kernel, or another callback for a different axis order
+or logical work count. `n_threads_from=None` requires explicit sizes;
+`"first_array"` always selects only the first axis. `launch_shape(args=(...))`
+inspects the inferred grid/block without compilation or a launch. CPU emulation
+and host/CUDA parity tests use the same defaults.
+
 * **1D**: `n_threads=n` with the default `block_size=128` (set per kernel with
   `CudaKernel(..., block_size=256)`).
 * **2D/3D**: `n_threads=(nx, ny)` and a tuple block, e.g.

docs/source/kernels/dispatch.md

@@ -26,14 +26,15 @@ def axpy_host(a, x, y, n):
 
 axpy = xp.kernels.Kernel(axpy_host, xp.cuda.CudaKernel(AXPY, "axpy"), name="axpy")
 
-axpy(2.0, x, y, x.size, n_threads=x.size)
+axpy(2.0, x, y, x.size)  # infer n_threads = x.shape[0]
 ```
 
 * On the **NumPy backend** the host kernel is called with the positional
   arguments; `n_threads`, `grid`, `block`, `shared_mem` and `stream` are
   ignored.
-* On the **CuPy backend** the CUDA kernel is launched; `n_threads` (or `grid`)
-  is required.
+* On the **CuPy backend** the CUDA kernel is launched; thread counts default to
+  the first array's leading shape axes. Explicit `n_threads` or `grid` overrides
+  this choice; `n_threads_from=None` requires an explicit size.
 * A plain host function is wrapped in a [`PyccelKernel`](pyccel-kernel.md);
   pass `host_options={"outputs": (2,)}` to configure that wrapper, or pass a
   `PyccelKernel` you built yourself.
@@ -179,7 +180,7 @@ xp.kernels.set_kernel_implementation(None)  # back to the default
 ```
 
 or `CUNUMPY_KERNEL_IMPLEMENTATION=numpy` for a whole run (read at import, like
-`ARRAY_BACKEND`). A chosen implementation that a kernel does not have, or
+`CUNUMPY_BACKEND`). A chosen implementation that a kernel does not have, or
 cannot load, raises `LookupError` instead of running another one: a benchmark
 of numba never silently measures NumPy. `kernel.implementations` lists the
 implementations, `kernel.selected()` names the one a call with host arrays runs

docs/source/kernels/overview.md

@@ -24,7 +24,7 @@ unchanged.
 
 ## Which one do I need?
 
-* **"I just want my existing code to run with `ARRAY_BACKEND=cupy`."** Wrap the
+* **"I just want my existing code to run with `CUNUMPY_BACKEND=cupy`."** Wrap the
   host kernels in `PyccelKernel`. Everything works, but every kernel call
   copies its arrays to the host and back. This is a correct starting point, not
   a fast one.

docs/source/kernels/testing.md

@@ -274,7 +274,7 @@ transfer counting, the backend branches of a simulation) runs only with CuPy
 present. For CI machines without a GPU, cunumpy ships a strict stand-in:
 
 ```bash
-CUNUMPY_FAKE_CUPY=1 ARRAY_BACKEND=cupy pytest tests/
+CUNUMPY_FAKE_CUPY=1 CUNUMPY_BACKEND=cupy pytest tests/
 ```
 
 Its arrays live in host memory but are not NumPy arrays: `numpy.asarray(a)`
@@ -328,7 +328,7 @@ offsets.
   are reported as skipped, and `emulate_cuda_kernel` tests check the CUDA
   kernels' arithmetic (the runner needs a C++ compiler, which Linux images
   have).
-* Run it a second time with `CUNUMPY_FAKE_CUPY=1 ARRAY_BACKEND=cupy`, so the
+* Run it a second time with `CUNUMPY_FAKE_CUPY=1 CUNUMPY_BACKEND=cupy`, so the
   CuPy code paths are exercised on the CPU runner too (kernel launches are
   skipped).
 * Run the same suite on a GPU runner, optionally with `CUNUMPY_CUDA_DEBUG=1` so

docs/source/pyodide.md

@@ -31,7 +31,7 @@ await pyodide.runPythonAsync(`
 `);
 ```
 
-NumPy is selected by default when `ARRAY_BACKEND` is unset. Leave it unset or set
+NumPy is selected by default when `CUNUMPY_BACKEND` is unset. Leave it unset or set
 it to `numpy` before importing cuNumPy. `to_numpy` and `to_cunumpy` preserve
 existing NumPy arrays, including views. `synchronize` and `set_device` are no-ops
 on this backend.

docs/source/quickstart.md

@@ -23,7 +23,7 @@ results are ordinary NumPy (or CuPy) arrays.
 Select CuPy before the arrays are created, either from the shell:
 
 ```bash
-ARRAY_BACKEND=cupy python my_script.py
+CUNUMPY_BACKEND=cupy python my_script.py
 ```
 
 or in the program:

docs/source/troubleshooting.md

@@ -9,8 +9,8 @@ the real error. Typical causes: no CuPy installed, a CuPy wheel for a different
 CUDA major version, no GPU visible (`CUDA_VISIBLE_DEVICES` empty, a login node
 without GPUs), or a driver too old for the CUDA runtime.
 
-**`ARRAY_BACKEND=cupy` has no effect.** The variable is read once, when
-CuNumpy is first imported. Setting `os.environ["ARRAY_BACKEND"]` after the
+**`CUNUMPY_BACKEND=cupy` has no effect.** The variable is read once, when
+CuNumpy is first imported. Setting `os.environ["CUNUMPY_BACKEND"]` after the
 import does nothing; use `xp.set_backend()`.
 
 **Code still runs on the CPU after `set_backend("cupy")`.** Look for