Argument class for SplineFunction kernel arguments

CUDA_STRATEGY.md

@@ -61,9 +61,9 @@ CUDA kernels can be added one by one. If the code runs on the GPU and needs a ke
 | File | Content |
 |---|---|
 | `cunumpy.kernels`, `cunumpy.arguments`, `cunumpy.cuda` (cunumpy 0.6) | `Kernel`, `KernelCatalog`, `PyccelKernel` and `CudaKernel` in `cunumpy.kernels`; `CudaStructArguments`, `CudaStruct` and `write_cuda_header` in `cunumpy.arguments`; the device runtime (`bind_local_device`, ...) in `cunumpy.cuda`. Struphy uses cunumpy for dispatch and struct packing |
-| `src/struphy/kernel_arguments/` | the argument classes in pairs: the pyccel classes in `pusher_args_kernels.py` / `local_projectors_args_kernels.py` (NumPy backend) and their CUDA versions `Cuda<Name>` in `pusher_args_cuda.py` / `local_projectors_args_cuda.py` (CuPy backend; subclasses of `CudaStructArguments` whose `fields` define the C struct) |
-| `src/struphy/utils/cuda_arguments.py` | `CUDA_STRUCTS`, `CUDA_OPTIONS`, `write_pusher_header()` and `write_local_projectors_header()` |
-| `src/struphy/kernel_arguments/pusher_args.cuh`, `local_projectors_args.cuh` | the C structs `MarkerArgs`, `DerhamArgs`, `DomainArgs` and `LocalProjectorsArgs`, generated from the CUDA classes, using cunumpy array views |
+| `src/struphy/kernel_arguments/` | the argument classes in pairs: the pyccel classes in `pusher_args_kernels.py` / `local_projectors_args_kernels.py` / `spline_args_kernels.py` (NumPy backend) and their CUDA versions `Cuda<Name>` in `pusher_args_cuda.py` / `local_projectors_args_cuda.py` / `spline_args_cuda.py` (CuPy backend; subclasses of `CudaStructArguments` whose `fields` define the C struct) |
+| `src/struphy/utils/cuda_arguments.py` | `CUDA_STRUCTS`, `CUDA_OPTIONS`, `write_pusher_header()`, `write_local_projectors_header()` and `write_spline_header()` |
+| `src/struphy/kernel_arguments/pusher_args.cuh`, `local_projectors_args.cuh`, `spline_args.cuh` | the C structs `MarkerArgs`, `DerhamArgs`, `DomainArgs`, `LocalProjectorsArgs` and `SplineArgs`, generated from the CUDA classes, using cunumpy array views |
 | `src/struphy/geometry/base.py`, `src/struphy/pic/base.py`, `src/struphy/feec/psydac_derham.py` | `Domain.args_domain`, `Particles.args_markers` and `Derham.args_derham` are the pyccel class on NumPy and the CUDA class on CuPy, chosen once at construction; every kernel call goes through a `Kernel` |
 | `src/struphy/*/kernels/` (`pic/pushing`, `pic/accumulation`, `pic/diagnostics`, `pic/sph`, `bsplines`, `geometry`, `feec`, `feec/local_projectors`) | every kernel called from Python with argument objects, one folder each: 44 pusher/evaluation (incl. `reflect`), 16 accumulation, 10 marker diagnostics, 4 SPH evaluation, 3 spline evaluation, 4 geometry, 1 FEEC utility and 8 local projector kernels. Each folder's `__init__.py` declares its `Kernel`; the code imports it (`from struphy.geometry.kernels.kernel_evaluate import kernel_evaluate`). Sixteen have CUDA versions (`linear_vlasov_ampere` the latest), with parity cases in `pic/tests/cuda_parity_cases.py`, among them the four geometry kernels (PR 18) |
 | `src/struphy/geometry/evaluation_kernels.cuh`, `transform_kernels.cuh`, `spline_mappings_kernels.cuh`, `geometry/domains/<name>/<name>_cuda.cuh` | device versions of the mapping helpers for every mapping: `f`/`df` per analytic domain (`kind_map` 10–12, 20–22, 30–32, PR 18) and the spline mappings `spline_3d`, `spline_2d_straight`, `spline_2d_torus` (`kind_map` 0–2, PR 19), the `kind_map` switch, `det_df`, `df_inv`, `g`, `g_inv`, `select_metric_coeff`, `pull`, `push`, `tran` |
@@ -540,6 +540,7 @@ The argument classes come in pairs, one for each backend:
 | `pusher_args_kernels.DerhamArguments` | `pusher_args_cuda.CudaDerhamArguments` (`DerhamArgs`) |
 | `pusher_args_kernels.DomainArguments` | `pusher_args_cuda.CudaDomainArguments` (`DomainArgs`) |
 | `local_projectors_args_kernels.LocalProjectorsArguments` | `local_projectors_args_cuda.CudaLocalProjectorsArguments` (`LocalProjectorsArgs`) |
+| `spline_args_kernels.SplineArguments` | `spline_args_cuda.CudaSplineArguments` (`SplineArgs`) |
 
 The two classes of a pair take the same constructor arguments and have the same attributes; the CUDA struct
 adds only derived members that CUDA pointers cannot carry (`n_markers` is also a pyccel attribute; `nt1`,
@@ -566,6 +567,10 @@ step: CUDA versions of `kernel_evaluate_pic`, `kernel_evaluate` and the pull/pus
 `CudaLocalProjectorsArguments` is not used by any CUDA kernel yet; local projectors are rejected on the
 CuPy backend by `Derham`.
 
+`SplineFunction` creates one `SplineArguments` (`CudaSplineArguments` on CuPy) per component, holding `kind`,
+`pn`, `tn1`, `tn2`, `tn3` and `starts`; the three `eval_spline_mpi_*` entry kernels take it as `args_spline`
+instead of six loose arguments (#678). The device helper `eval_spline_mpi` keeps its flat signature.
+
 ### Kernel folders
 
 Every kernel that Python code calls with argument objects now lives in its own folder, like the pusher and

src/struphy/bsplines/kernels/eval_spline_mpi_markers/eval_spline_mpi_markers_cuda.cu

@@ -1,4 +1,5 @@
 #include "struphy/bsplines/evaluation_kernels_3d.cuh"
+#include "struphy/kernel_arguments/spline_args.cuh"
 
 // Same arguments, in the same order, as the pyccel kernel in eval_spline_mpi_markers_kernels.py. Array views carry
 // dimensions and strides, so no CUDA-only lengths are needed.
@@ -11,18 +12,12 @@
  * @param markers Marker coordinates (Np x 3 or wider, any strides); rows flagged with -1 in the first
  *        column are not on the process domain and are skipped.
  * @param _data Spline coefficients of the current process, any strides.
- * @param kind Kind of 1d basis in each direction (three entries): 0 = N-spline, 1 = D-spline.
- * @param pn Spline degrees of V0 in each direction (three entries, 1 to 8).
- * @param tn1 Knot vector of V0 along the first axis.
- * @param tn2 Knot vector of V0 along the second axis.
- * @param tn3 Knot vector of V0 along the third axis.
- * @param starts Start indices of the splines on the current process (three entries).
+ * @param args_spline Kind of 1d basis (0 = N-spline, 1 = D-spline), spline degrees (1 to 8) and knot vectors of V0
+ *        in each direction, and start indices of the splines on the current process.
  * @param values Output values S_p = S(*markers[p, :]), one per row; skipped rows are left unchanged.
  */
 extern "C" __global__ void eval_spline_mpi_markers(Array2D<double> markers, Array3D<double> _data,
-                                                   const long long* kind, long long* pn, Array1D<double> tn1,
-                                                   Array1D<double> tn2, Array1D<double> tn3, long long* starts,
-                                                   Array1D<double> values) {
+                                                   SplineArgs args_spline, Array1D<double> values) {
     // CUDA-only: the marker row of this thread replaces the pyccel loop variable
     long long ip = (long long)blockDim.x * blockIdx.x + threadIdx.x;
     long long Np = markers.shape[0];
@@ -31,6 +26,6 @@ extern "C" __global__ void eval_spline_mpi_markers(Array2D<double> markers, Arra
     // point not in process domain
     if (markers(ip, 0) == -1.) return;
 
-    values(ip) = eval_spline_mpi(markers(ip, 0), markers(ip, 1), markers(ip, 2), _data, kind, pn, tn1, tn2, tn3,
-                                 starts);
+    values(ip) = eval_spline_mpi(markers(ip, 0), markers(ip, 1), markers(ip, 2), _data, args_spline.kind,
+                                 args_spline.pn, args_spline.tn1, args_spline.tn2, args_spline.tn3, args_spline.starts);
 }

src/struphy/bsplines/kernels/eval_spline_mpi_markers/eval_spline_mpi_markers_kernels.py

@@ -2,18 +2,15 @@
 
 from numpy import shape
 
+import struphy.kernel_arguments.spline_args_kernels as spline_args_kernels  # do not remove; needed to identify dependencies
 from struphy.bsplines.evaluation_kernels_3d import eval_spline_mpi
+from struphy.kernel_arguments.spline_args_kernels import SplineArguments
 
 
 def eval_spline_mpi_markers(
     markers: "float[:,:]",
     _data: "float[:,:,:]",
-    kind: "int[:]",
-    pn: "int[:]",
-    tn1: "float[:]",
-    tn2: "float[:]",
-    tn3: "float[:]",
-    starts: "int[:]",
+    args_spline: "SplineArguments",
     values: "float[:]",
 ):
     """
@@ -28,17 +25,9 @@ def eval_spline_mpi_markers(
         _data : array[float]
             The spline coefficients c_ijk.
 
-        kind : array[int]
-            Kind of 1d basis in each direction: 0 = N-spline, 1 = D-spline.
-
-        pn : array[int]
-            Spline degrees of V0 in each direction.
-
-        tn1, tn2, tn3 : array[float]
-            Knot vectors of V0 in each direction.
-
-        starts : array[float]
-            Start indices of splines on current process.
+        args_spline : SplineArguments
+            Kind of 1d basis, spline degrees and knot vectors of V0, and start indices of the splines on the
+            current process.
 
         values : array[float]
             Return 1D array for spline values S_p = S(*markers[p, :]).
@@ -55,10 +44,10 @@ def eval_spline_mpi_markers(
             markers[ip, 1],
             markers[ip, 2],
             _data,
-            kind,
-            pn,
-            tn1,
-            tn2,
-            tn3,
-            starts,
+            args_spline.kind,
+            args_spline.pn,
+            args_spline.tn1,
+            args_spline.tn2,
+            args_spline.tn3,
+            args_spline.starts,
         )

src/struphy/bsplines/kernels/eval_spline_mpi_matrix/eval_spline_mpi_matrix_cuda.cu

@@ -1,4 +1,5 @@
 #include "struphy/bsplines/evaluation_kernels_3d.cuh"
+#include "struphy/kernel_arguments/spline_args.cuh"
 
 // Same arguments, in the same order, as the pyccel kernel in eval_spline_mpi_matrix_kernels.py. Array views carry
 // dimensions and strides, so no CUDA-only lengths are needed.
@@ -12,18 +13,13 @@
  * @param eta2 Second coordinates of the points, same shape as eta1.
  * @param eta3 Third coordinates of the points, same shape as eta1.
  * @param _data Spline coefficients of the current process, any strides.
- * @param kind Kind of 1d basis in each direction (three entries): 0 = N-spline, 1 = D-spline.
- * @param pn Spline degrees of V0 in each direction (three entries, 1 to 8).
- * @param tn1 Knot vector of V0 along the first axis.
- * @param tn2 Knot vector of V0 along the second axis.
- * @param tn3 Knot vector of V0 along the third axis.
- * @param starts Start indices of the splines on the current process (three entries).
+ * @param args_spline Kind of 1d basis (0 = N-spline, 1 = D-spline), spline degrees (1 to 8) and knot vectors of V0
+ *        in each direction, and start indices of the splines on the current process.
  * @param values Output values, same shape as eta1; entries of flagged points are left unchanged.
  */
 extern "C" __global__ void eval_spline_mpi_matrix(Array3D<double> eta1, Array3D<double> eta2, Array3D<double> eta3,
-                                                  Array3D<double> _data, const long long* kind, long long* pn,
-                                                  Array1D<double> tn1, Array1D<double> tn2, Array1D<double> tn3,
-                                                  long long* starts, Array3D<double> values) {
+                                                  Array3D<double> _data, SplineArgs args_spline,
+                                                  Array3D<double> values) {
     // CUDA-only: flat index of this thread, split into the pyccel loop variables i, j, k
     long long ijk = (long long)blockDim.x * blockIdx.x + threadIdx.x;
     if (ijk >= values.shape[0] * values.shape[1] * values.shape[2]) return;
@@ -34,6 +30,7 @@ extern "C" __global__ void eval_spline_mpi_matrix(Array3D<double> eta1, Array3D<
     // point not in process domain
     if (eta1(i, j, k) == -1. || eta2(i, j, k) == -1. || eta3(i, j, k) == -1.) return;
 
-    values(i, j, k) = eval_spline_mpi(eta1(i, j, k), eta2(i, j, k), eta3(i, j, k), _data, kind, pn, tn1, tn2, tn3,
-                                      starts);
+    values(i, j, k) =
+        eval_spline_mpi(eta1(i, j, k), eta2(i, j, k), eta3(i, j, k), _data, args_spline.kind, args_spline.pn,
+                        args_spline.tn1, args_spline.tn2, args_spline.tn3, args_spline.starts);
 }

src/struphy/bsplines/kernels/eval_spline_mpi_matrix/eval_spline_mpi_matrix_kernels.py

@@ -2,20 +2,17 @@
 
 from numpy import shape
 
+import struphy.kernel_arguments.spline_args_kernels as spline_args_kernels  # do not remove; needed to identify dependencies
 from struphy.bsplines.evaluation_kernels_3d import eval_spline_mpi
+from struphy.kernel_arguments.spline_args_kernels import SplineArguments
 
 
 def eval_spline_mpi_matrix(
     eta1: "float[:,:,:]",
     eta2: "float[:,:,:]",
     eta3: "float[:,:,:]",
     _data: "float[:,:,:]",
-    kind: "int[:]",
-    pn: "int[:]",
-    tn1: "float[:]",
-    tn2: "float[:]",
-    tn3: "float[:]",
-    starts: "int[:]",
+    args_spline: "SplineArguments",
     values: "float[:,:,:]",
 ):
     """
@@ -30,17 +27,9 @@ def eval_spline_mpi_matrix(
         _data : array[float]
             The spline coefficients c_ijk.
 
-        kind : array[int]
-            Kind of 1d basis in each direction: 0 = N-spline, 1 = D-spline.
-
-        pn : array[int]
-            Spline degrees of V0 in each direction.
-
-        tn1, tn2, tn3 : array[float]
-            Knot vectors of V0 in each direction.
-
-        starts : array[float]
-            Start indices of splines on current process.
+        args_spline : SplineArguments
+            Kind of 1d basis, spline degrees and knot vectors of V0, and start indices of the splines on the
+            current process.
 
         values : array[float]
             Return array for spline values S_ijk = S(eta1[i,j,k], eta2[i,j,k], eta3[i,j,k]).
@@ -63,10 +52,10 @@ def eval_spline_mpi_matrix(
                     eta2[i, j, k],
                     eta3[i, j, k],
                     _data,
-                    kind,
-                    pn,
-                    tn1,
-                    tn2,
-                    tn3,
-                    starts,
+                    args_spline.kind,
+                    args_spline.pn,
+                    args_spline.tn1,
+                    args_spline.tn2,
+                    args_spline.tn3,
+                    args_spline.starts,
                 )

src/struphy/bsplines/kernels/eval_spline_mpi_sparse_meshgrid/eval_spline_mpi_sparse_meshgrid_cuda.cu

@@ -1,7 +1,8 @@
 #include "struphy/bsplines/evaluation_kernels_3d.cuh"
+#include "struphy/kernel_arguments/spline_args.cuh"
 
-// Same arguments, in the same order, as the pyccel kernel in eval_spline_mpi_sparse_meshgrid_kernels.py. Array views carry
-// dimensions and strides, so no CUDA-only lengths are needed.
+// Same arguments, in the same order, as the pyccel kernel in eval_spline_mpi_sparse_meshgrid_kernels.py. Array views
+// carry dimensions and strides, so no CUDA-only lengths are needed.
 
 /**
  * Sparse meshgrid evaluation of a distributed spline, as in evaluation_kernels_3d.eval_spline_mpi_sparse_meshgrid.
@@ -12,19 +13,13 @@
  * @param eta2 Second coordinates, shape (1, n2, 1).
  * @param eta3 Third coordinates, shape (1, 1, n3).
  * @param _data Spline coefficients of the current process, any strides.
- * @param kind Kind of 1d basis in each direction (three entries): 0 = N-spline, 1 = D-spline.
- * @param pn Spline degrees of V0 in each direction (three entries, 1 to 8).
- * @param tn1 Knot vector of V0 along the first axis.
- * @param tn2 Knot vector of V0 along the second axis.
- * @param tn3 Knot vector of V0 along the third axis.
- * @param starts Start indices of the splines on the current process (three entries).
+ * @param args_spline Kind of 1d basis (0 = N-spline, 1 = D-spline), spline degrees (1 to 8) and knot vectors of V0
+ *        in each direction, and start indices of the splines on the current process.
  * @param values Output values, shape (n1, n2, n3); entries of flagged points are left unchanged.
  */
 extern "C" __global__ void eval_spline_mpi_sparse_meshgrid(Array3D<double> eta1, Array3D<double> eta2,
                                                            Array3D<double> eta3, Array3D<double> _data,
-                                                           const long long* kind, long long* pn, Array1D<double> tn1,
-                                                           Array1D<double> tn2, Array1D<double> tn3,
-                                                           long long* starts, Array3D<double> values) {
+                                                           SplineArgs args_spline, Array3D<double> values) {
     // CUDA-only: flat index of this thread, split into the pyccel loop variables i, j, k
     long long ijk = (long long)blockDim.x * blockIdx.x + threadIdx.x;
     if (ijk >= values.shape[0] * values.shape[1] * values.shape[2]) return;
@@ -35,6 +30,7 @@ extern "C" __global__ void eval_spline_mpi_sparse_meshgrid(Array3D<double> eta1,
     // point not in process domain
     if (eta1(i, 0, 0) == -1. || eta2(0, j, 0) == -1. || eta3(0, 0, k) == -1.) return;
 
-    values(i, j, k) = eval_spline_mpi(eta1(i, 0, 0), eta2(0, j, 0), eta3(0, 0, k), _data, kind, pn, tn1, tn2, tn3,
-                                      starts);
+    values(i, j, k) =
+        eval_spline_mpi(eta1(i, 0, 0), eta2(0, j, 0), eta3(0, 0, k), _data, args_spline.kind, args_spline.pn,
+                        args_spline.tn1, args_spline.tn2, args_spline.tn3, args_spline.starts);
 }

src/struphy/bsplines/kernels/eval_spline_mpi_sparse_meshgrid/eval_spline_mpi_sparse_meshgrid_kernels.py

@@ -1,19 +1,16 @@
 """Sparse meshgrid evaluation of a tensor-product spline, distributed."""
 
+import struphy.kernel_arguments.spline_args_kernels as spline_args_kernels  # do not remove; needed to identify dependencies
 from struphy.bsplines.evaluation_kernels_3d import eval_spline_mpi
+from struphy.kernel_arguments.spline_args_kernels import SplineArguments
 
 
 def eval_spline_mpi_sparse_meshgrid(
     eta1: "float[:,:,:]",
     eta2: "float[:,:,:]",
     eta3: "float[:,:,:]",
     _data: "float[:,:,:]",
-    kind: "int[:]",
-    pn: "int[:]",
-    tn1: "float[:]",
-    tn2: "float[:]",
-    tn3: "float[:]",
-    starts: "int[:]",
+    args_spline: "SplineArguments",
     values: "float[:,:,:]",
 ):
     """
@@ -28,17 +25,9 @@ def eval_spline_mpi_sparse_meshgrid(
         _data : array[float]
             The spline coefficients c_ijk.
 
-        kind : array[int]
-            Kind of 1d basis in each direction: 0 = N-spline, 1 = D-spline.
-
-        pn : array[int]
-            Spline degrees of V0 in each direction.
-
-        tn1, tn2, tn3 : array[float]
-            Knot vectors of V0 in each direction.
-
-        starts : array[float]
-            Start indices of splines on current process.
+        args_spline : SplineArguments
+            Kind of 1d basis, spline degrees and knot vectors of V0, and start indices of the splines on the
+            current process.
 
         values : array[float]
             Return array for spline values S_ijk = S(eta1[i,0,0], eta2[0,j,0], eta3[0,0,k]).
@@ -65,10 +54,10 @@ def eval_spline_mpi_sparse_meshgrid(
                     eta2[0, j, 0],
                     eta3[0, 0, k],
                     _data,
-                    kind,
-                    pn,
-                    tn1,
-                    tn2,
-                    tn3,
-                    starts,
+                    args_spline.kind,
+                    args_spline.pn,
+                    args_spline.tn1,
+                    args_spline.tn2,
+                    args_spline.tn3,
+                    args_spline.starts,
                 )