Cuda pr 8 cuda headers

[max-models]

Solves the following issue(s):

PR 8 of the CUDA strategy (CUDA_STRATEGY.md). No issue closed.

Stack: #649 → #651 → #652 → this PR. Only the commit Shared CUDA header for the kernel argument classes is new relative to #652 (plus a merge of #652).

Core changes:

CUDA kernels used to repeat the full flat signature of the argument classes (26 parameters for markers and domain alone). CuPy does not check kernel signatures, so adding a field to MarkerArguments would have shifted all following arguments of every CUDA kernel without an error. Now each argument class is one C struct, defined once.

• New src/struphy/kernel_arguments/pusher_args.cuh with struct MarkerArgs, DerhamArgs and DomainArgs. Kernels take them by value in place of the pyccel argument classes, e.g. void push_eta_linear(double dt, int stage, MarkerArgs args_markers, DomainArgs args_domain), and read args_markers.markers, args_markers.n_markers, and so on.
• The member names are the attribute names of the pyccel classes. The only CUDA-specific member is MarkerArgs.n_cols, since pyccel kernels take it from markers.shape[1]. Accordingly, CudaMarkerArguments.first_pusher_idx is now first_init_idx, as in MarkerArguments.
• src/struphy/utils/cuda_arguments.py:
  • Each Argument subclass lists its members once, e.g. fields = (("double*", "markers"), ...).
  • struct_dtype() builds a NumPy structured dtype with align=True, which gives C alignment and padding. Pointer members are uint64 device addresses (array.data.ptr).
  • The struct is a numpy.void, which CuPy passes by value. It is packed once in the constructor, so kernel calls do no extra work.
• Scalar members are checked when the struct is packed. A float for an int member raises TypeError, and a value that does not fit raises OverflowError. Before, they arrived in the kernel truncated or wrapped around; plain NumPy assignment also truncates 1.7 to 1. This solves the "scalar types" open question for everything inside the argument classes. Scalars passed directly to a kernel (dt, stage) are still unchecked.
• Because the struct holds device addresses, it is repacked when an argument object is deepcopied or unpickled (__getstate__/__setstate__).
• CudaKernel (src/struphy/utils/kernel_backends.py) compiles with the folder that contains the struphy package on the NVRTC include path, so .cu files can #include "struphy/kernel_arguments/pusher_args.cuh". .cuh files are already package data.

Why structs, and why _pack() and the copy handling are needed:

The underlying problem: CuPy passes kernel arguments positionally and never checks them against the signature. With flat parameter lists, a new field in an argument class shifts every later argument of every CUDA kernel, silently and on the GPU only. Some single, checked definition of each argument layout is therefore necessary. A C struct is the simplest one that also works for device helper functions.

What each piece is for:

• struct_dtype() with align=True: CuPy passes a struct by value only as a NumPy structured scalar, and align=True gives the C padding (the int members are followed by pointers).
• Pointers as uint64 (data.ptr): a struct cannot hold a CuPy array, only its device address.
• Repacking in __setstate__: a direct consequence of storing addresses. The arrays of a deepcopied or unpickled object have new addresses, and without the repack a copy would silently read the original's arrays.
• Packing once in the constructor: so kernel calls do no extra work. Packing at every call would also be fine (well under 1 µs next to a launch of about 70 µs).
• Scalar checks: the only optional part. They are cheap and turn truncation or wrap-around into an exception.

For:

• Silent argument shifts cannot happen, and a test that needs no GPU checks the header against fields.
• CUDA code mirrors the pyccel code: pyccel kernels and helpers take args_domain and read args_domain.kind_map, and CUDA reads exactly the same. This decides it for PR 10. The mapping evaluation, B-spline and boundary helpers take these argument objects, so a __device__ helper takes one const DomainArgs& instead of 12+ parameters forwarded at every call site.
• Short kernel signatures. Adding a field touches fields and the header, not every kernel.
• No runtime cost.

Against:

• Raw pointers bypass CuPy's per-launch checks: a host array or an array on another GPU is no longer rejected at launch. _cupy_array() checks at construction that the array is a C-contiguous CuPy array of the right dtype, but not which device it is on. That will matter with one GPU per MPI rank; adding a device check is a small follow-up.
• If an owner reallocates an array without rebuilding its argument object, the struct points at the old array. The pyccel argument classes have the same issue, since they also hold references, so this is not new.
• About 100 lines of infrastructure plus layout tests before any real CUDA kernel exists. But this choice fixes the signatures of all device helpers in PR 10, so changing it later would mean rewriting them.
• Alignment is subtle. The host sizeof/offsetof check and the NVRTC layout test cover it.

Alternatives considered:

• Flat parameters generated from a C macro (#define MARKER_ARGS double* markers, ...), with get_cuda_args() returning the arrays in fields order. This also prevents silent shifts and keeps CuPy's checks, without pointers or repacking. But every device helper then needs a long parameter list, or a second macro that rebuilds a struct inside the kernel. That is more macro machinery, and the CUDA code no longer mirrors the pyccel helpers.
• Plain flat lists without a single definition: the state before this PR, and the dangerous one.

Tests (src/struphy/pic/tests/test_kernel_backends.py):

• Run everywhere:
  • test_cuda_argument_structs_match_header: parses pusher_args.cuh and compares member names, C types and order with fields.
  • test_cuda_struct_members_are_pyccel_attributes: the member names exist as attributes in pusher_args_kernels.py, except n_cols.
• GPU only:
  • test_cuda_struct_layout: an NVRTC-compiled kernel reports sizeof and the member offsets and sizes, which are compared with the NumPy dtype.
  • test_cuda_struct_scalars_are_checked and test_cuda_struct_follows_copies.
  • The demo kernels push_eta_linear and write_scalars now use the structs, and write_scalars also reads a pointer member of each struct.
• Fixed: several GPU-only tests here and in geometry/tests/test_domain.py still used the values attribute, which get_cuda_args() replaced. They would have failed on a GPU.

The GPU-only tests have not run on a GPU yet. On the host:

• Compiling the header with a C++ compiler gives the same sizeof/offsetof as the NumPy dtypes. x86-64 and arm64 lay these structs out like CUDA.
• The tests that do not launch kernels pass with a strict host stand-in for CuPy.
• test_pusher_accepts_kernel fails with and without this PR (see Cuda pr 7 derham #652). It is unrelated.

Model-specific changes:

None

🤖 Generated with Claude Code