By making a temporary variable for accessing the ID view, the Kokkos kernel no longer needs to directly access the particle attribute pdata. This means the kernel doesn't need a copy of pdata and removes the cryptic Cuda compilation warnings of the form
src/Solver/FFTPeriodicPoissonSolver.hpp(193): warning: calling a __host__ function("std::__shared_count<( ::__gnu_cxx::_Lock_policy)2> ::__shared_count") from a __host__ __device__ function("void ippl::ParticleSpatialLayout<double, (unsigned int)3u, ::ippl::UniformCartesian<double, (unsigned int)3u> > ::update< ::ChargedParticles<ippl::ParticleSpatialLayout<double, (unsigned int)3u, ::ippl::UniformCartesian<double, (unsigned int)3u> > > > (T1 &, T1 &)::[lambda(unsigned long, unsigned long &) (instance 1)]::[lambda(unsigned long, unsigned long &) (instance 1)]") is not allowed
where the line reference is completely irrelevant. The relevant __host__ function for this example comes from the shared pointer to the FFT solver instance in ChargedParticles.
vinciguerra_a authorede0e412f0