Mirror pyccel modules as CUDA namespaces

[max-models]

Stack: part 14 of 14, based on #723 (merge that first), top of the stack. Full order: #709 → #708 → #711 → #712 → #705 → #713 → #714 → #718 → #720 → #721 → #724 → #725 → #723 → #722.

Stack fixes: applied the struphy_cuda::<pyccel module> rule to the CUDA code of the PRs below: linear_vlasov_ampere (#705), vlasov_maxwell (#713), bstar_parallel_3form (#718), eval_spline_mpi_tensor_product_fixed (#724) and the SplineArgs versions of eval_spline_mpi_* (#721). Each of them has using namespace struphy_cuda; and qualified helper calls. The helpers outer, fill_mat, fill_mat_vec and m_v_fill_b_v1_symm (#705/#713) are inside their module namespaces, and m_v_fill_b_v1_symm calls filler_kernels::/evaluation_kernels_3d::. The fill_v1_symm test wrapper is qualified. The "open PRs that need the same treatment" below are therefore covered by this stack. Local re-run on the stacked branch: test_cuda_parity.py, test_kernel_backends.py, test_evaluation_cuda.py, test_device_helpers.py (59 passed, 419 skipped), test_cuda_emulation.py (74 passed; it compiles every .cu/.cuh with the host C++ compiler).

---

Solves the following issue(s):

Closes #672, tracked in #650. The CUDA device helpers now live in C++ namespaces that mirror the pyccel modules they port. A pyccel call linalg_kernels.matrix_inv(...) becomes linalg_kernels::matrix_inv(...) in CUDA, so the 1:1 correspondence is visible at every call site and two helpers with the same name in different modules (det/det_df, f/df, push/pull) can no longer clash.

Core changes:

• One namespace per pyccel module, nested in struphy_cuda as proposed in the issue (struphy_cuda::linalg_kernels::matrix_inv). The namespace is named after the module the header ports:
  • bsplines_kernels, evaluation_kernels_2d, evaluation_kernels_3d (bsplines/)
  • evaluation_kernels, spline_mappings_kernels, transform_kernels (geometry/)
  • one per mapping, named after its pyccel module: cuboid_kernels, orthogonal_kernels, colella_kernels, hollow_cylinder_kernels, powered_elliptic_cylinder_kernels, hollow_torus_kernels, shafranov_shift_cylinder_kernels, shafranov_sqrt_cylinder_kernels, shafranov_dshaped_cylinder_kernels (geometry/domains/<name>/<name>_cuda.cuh)
  • linalg_kernels, filler_kernels, particle_to_mat_kernels, pusher_utilities_kernels
• CUDA-only helpers and constants go into the namespace of the module whose calculation they extract: bsplines_kernels::MAX_SPLINE_DEGREE, evaluation_kernels_3d::SplineScratch, spline_mappings_kernels::spline_index_row/first_plane, transform_kernels::column_norm, pusher_utilities_kernels::reflect_velocity.
• Call sites:
  • Inside a header, calls into another module are qualified (linalg_kernels::det(...) in evaluation_kernels::det_df). Calls within the same module are not, as in pyccel.
  • Every .cu file writes using namespace struphy_cuda; after its includes and qualifies every helper call (evaluation_kernels_3d::get_spans(...), evaluation_kernels::df(...)).
  • The CUDA sources embedded in test_device_helpers.py and geometry/tests/spline_mapping_cases.py use fully qualified names.
• Left global, on purpose:
  • struphy_cuda::pi (geometry/domains/constants_cuda.cuh). It stands in for numpy.pi, not for a struphy module.
  • The generated argument structs MarkerArgs, DerhamArgs, DomainArgs and LocalProjectorsArgs (kernel_arguments/*.cuh, written by cunumpy's write_cuda_header). cunumpy matches the extern "C" __global__ signatures against the struct names, and the CPU emulation (pic/tests/cuda_emulation.py) rebuilds the structs by name, so the generators and the emulation are unchanged.
• evaluation_kernels_3d.cuh: it had three namespace struphy_cuda blocks with an #include in between. It now has one block. eval_spline_mpi was outside the namespace and is now evaluation_kernels_3d::eval_spline_mpi, like its pyccel counterpart.
• Pure refactor: no function body, signature or entry kernel changes. Lines changed by the renaming were rewrapped at 120 columns.

Open PRs that need the same treatment after merge: #705, #713 and #718 add .cu/.cuh code that still calls struphy_cuda::<name>. Whichever of them and this PR merges second needs the rule applied: using namespace struphy_cuda; in new .cu files, and new helpers in their module's namespace. Their branches were not touched.

Model-specific changes:

None.

Documentation changes:

CUDA_STRATEGY.md: a new convention bullet under "Target layout" documents the rule: the namespace is struphy_cuda::<pyccel module>, CUDA-only helpers go into their module, cross-module calls are qualified, .cu files use using namespace struphy_cuda;, and the exceptions are pi and the generated structs.

Local tests (macOS, no GPU; pyccel kernels compiled with GNU/Fortran; GPU tests not run, so nothing was compiled with NVRTC). The CPU emulation compiles every .cu file and every helper header as C++17 with the host compiler. That is the syntax check for the namespaces; it caught one misplaced brace while I was working on this.

• src/struphy/pic/tests/test_cuda_parity.py
• src/struphy/pic/tests/test_cuda_emulation.py
• src/struphy/pic/tests/test_device_helpers.py
• src/struphy/bsplines/tests/test_evaluation_cuda.py
• src/struphy/pic/tests/test_kernel_backends.py

Result: 107 passed, 392 skipped (the GPU-only tests). devel gives the same counts.

One open point for the GPU run: the nested namespace definition (namespace struphy_cuda::bsplines_kernels {) needs C++17. NVRTC uses C++17 by default since CUDA 11, and the emulation compiles with -std=c++17.

🤖 Generated with Claude Code