Fix Cuda host/device function warnings

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.

src/Particle/ParticleBase.hpp

@@ -112,10 +112,13 @@ namespace ippl {
         }
 
         // set the unique ID value for these new particles
+        auto pIDs = ID.getView();
+        auto nextID = this->nextID_m;
+        auto numNodes = this->numNodes_m;
         Kokkos::parallel_for("ParticleBase<PLayout, Properties...>::create(size_t)",
                              Kokkos::RangePolicy(localNum_m, nLocal),
-                             KOKKOS_CLASS_LAMBDA(const std::int64_t i) {
-                                 ID(i) = this->nextID_m + this->numNodes_m * i;
+                             KOKKOS_LAMBDA(const std::int64_t i) {
+                                 pIDs(i) = nextID + numNodes * i;
                              });
         nextID_m += numNodes_m * (nLocal - localNum_m);