Fixissue/#415

[aaadelmann]

## Summary

This PR add mixed-value scatter and fixes CUDA compilation issues.

The affected API is `ParticleAttrib::scatter`, in particular cases where:
- a custom `Kokkos::RangePolicy` is used,
- an optional `hash_type` index map is supplied,
- the particle attribute value type differs from the field value type, for example `ParticleAttrib<float> -> Field<double>`.

## What Changed

### CUDA-safe scatter dispatch

`ParticleAttrib::scatter` now performs the hash/no-hash choice on the host and dispatches to a compile-time implementation:

```cpp
particleAttribScatterImpl<true>(...)
particleAttribScatterImpl<false>(...)

The actual Kokkos kernel body was moved out of the private ParticleAttrib member scope into a namespace-scope helper:

ippl::detail::particleAttribScatterImpl(...)

This avoids NVCC errors for extended __host__ __device__ lambdas whose enclosing parent function has private/protected class access.

The mapped-index selection is also handled through a small KOKKOS_INLINE_FUNCTION helper:

ippl::detail::scatterMappedIndex<UseHashView>(idx, hashView)

This avoids CUDA’s restriction around first-capturing a variable inside an if constexpr context.

Mixed value type scatter

The scatter implementation now supports different particle attribute and field value types, for example:

ParticleAttrib<float> q;
Field<double, Dim, Mesh, Centering> rho;

scatter(q, rho, positions);

The scattered particle value is passed to the field scatter operation using the particle attribute’s value type, while the field accumulation happens in the field’s value type.

Regression tests

unit_tests/Particle/GatherScatterTest.cpp now includes explicit coverage for:

• ParticleAttrib<float> -> Field<double> plain scatter
• ParticleAttrib<float> -> Field<double> hashed scatter
• dimensions 1, 2, and 3 through the existing typed test suite

The test uses a conservation tolerance based on the weaker precision of the scatter attribute and field value type. This avoids applying double-precision tolerances to a float-valued particle source while still scaling the tolerance with the conserved charge norm.

Motivation

The original hashed scatter refactor triggered CUDA compilation failures on GH200/NVCC:

error: An extended __host__ __device__ lambda cannot first-capture variable in constexpr-if context

and then:

error: The enclosing parent function ("scatterImpl") for an extended __host__ __device__ lambda cannot have private or protected access within its class

These were CUDA-specific restrictions around extended lambdas. The final implementation keeps the same host-side dispatch semantics, but places the kernel-bearing function in namespace scope and avoids first-capture inside if constexpr.

Validation

Tested locally with a serial debug build:

cmake --build build-fixissue-415-merge --target GatherScatterTest -j 8
ctest --test-dir build-fixissue-415-merge --output-on-failure -R '^GatherScatterTest$'
mpirun -np 2 build-fixissue-415-merge/unit_tests/Particle/GatherScatterTest --gtest_filter='*MixedValueType*'
mpirun -np 4 build-fixissue-415-merge/unit_tests/Particle/GatherScatterTest --gtest_filter='*MixedValueType*'

All tests passed.

The CUDA build was validated externally through the project CI/CD matrix after the scatter implementation changes.

[aaadelmann]

closes #415

[srikrrish]

Looks good to me. But currently there is one conflict in the CMakeLists which must be resolved and CI needs to be invoked before merging

[aaadelmann]

You already approved in the past, here I needed to fix some CUDA related issues and added an other test

[aaadelmann]

cscs-ci run cscs-ci-gh200, cscs-ci-mi300, cscs-ci-openmp