diff --git a/particle_structs/src/cabm/cabm_migrate.hpp b/particle_structs/src/cabm/cabm_migrate.hpp index 75b93d6d..bdfaf83e 100644 --- a/particle_structs/src/cabm/cabm_migrate.hpp +++ b/particle_structs/src/cabm/cabm_migrate.hpp @@ -39,7 +39,7 @@ namespace pumipic { const lid_t process = new_process(particle_id); if (mask && (process != comm_rank)) { const lid_t process_index = dist.index(process); - Kokkos::atomic_increment(&num_send_particles(process_index)); + Kokkos::atomic_inc(&num_send_particles(process_index)); } }; parallel_for(count_sending_particles); diff --git a/particle_structs/src/cabm/cabm_rebuild.hpp b/particle_structs/src/cabm/cabm_rebuild.hpp index 3ff396c8..a75b7e3a 100644 --- a/particle_structs/src/cabm/cabm_rebuild.hpp +++ b/particle_structs/src/cabm/cabm_rebuild.hpp @@ -36,9 +36,9 @@ namespace pumipic { if (active.access(soa,tuple)) { lid_t parent = new_element(soa*soa_len + tuple); if (parent > -1) // count particles to be kept - Kokkos::atomic_increment(&elmDegree_d(parent)); + Kokkos::atomic_inc(&elmDegree_d(parent)); else // count particles to be deleted - Kokkos::atomic_increment(&num_removed_d(0)); + Kokkos::atomic_inc(&num_removed_d(0)); } }; Cabana::SimdPolicy simd_policy(0, capacity_); diff --git a/particle_structs/src/csr/CSR_migrate.hpp b/particle_structs/src/csr/CSR_migrate.hpp index f27863dc..28512b3b 100644 --- a/particle_structs/src/csr/CSR_migrate.hpp +++ b/particle_structs/src/csr/CSR_migrate.hpp @@ -39,7 +39,7 @@ namespace pumipic { const lid_t process = new_process(particle_id); if (mask && (process != comm_rank)) { const lid_t process_index = dist.index(process); - Kokkos::atomic_increment(&num_send_particles(process_index)); + Kokkos::atomic_inc(&num_send_particles(process_index)); } }; parallel_for(count_sending_particles); diff --git a/particle_structs/src/csr/CSR_rebuild.hpp b/particle_structs/src/csr/CSR_rebuild.hpp index 26a9596d..4b769185 100644 --- a/particle_structs/src/csr/CSR_rebuild.hpp +++ b/particle_structs/src/csr/CSR_rebuild.hpp @@ -30,9 +30,9 @@ namespace pumipic { // Fill ptcls per elem for existing ptcls auto count_existing = PS_LAMBDA(const lid_t& elm_id, const lid_t& ptcl_id, const bool& mask) { if (new_element[ptcl_id] > -1) - Kokkos::atomic_increment(&particles_per_element[new_element[ptcl_id]]); + Kokkos::atomic_inc(&particles_per_element[new_element[ptcl_id]]); else - Kokkos::atomic_increment(&num_removed_d(0)); + Kokkos::atomic_inc(&num_removed_d(0)); }; parallel_for(count_existing,"fill particle Per Element existing"); lid_t num_removed = getLastValue(num_removed_d); // save number of removed particles for later @@ -40,7 +40,7 @@ namespace pumipic { Kokkos::parallel_for("fill particlesPerElementNew", new_particle_elements.size(), KOKKOS_LAMBDA(const int& i) { assert(new_particle_elements[i] > -1); - Kokkos::atomic_increment(&particles_per_element[new_particle_elements[i]]); + Kokkos::atomic_inc(&particles_per_element[new_particle_elements[i]]); }); RecordTime("CSR calc ppe", time_ppe.seconds()); diff --git a/particle_structs/src/dps/dps_migrate.hpp b/particle_structs/src/dps/dps_migrate.hpp index d96da6f9..45017dd4 100644 --- a/particle_structs/src/dps/dps_migrate.hpp +++ b/particle_structs/src/dps/dps_migrate.hpp @@ -39,7 +39,7 @@ namespace pumipic { const lid_t process = new_process(particle_id); if (mask && (process != comm_rank)) { const lid_t process_index = dist.index(process); - Kokkos::atomic_increment(&num_send_particles(process_index)); + Kokkos::atomic_inc(&num_send_particles(process_index)); } }; parallel_for(count_sending_particles); diff --git a/particle_structs/src/dps/dps_rebuild.hpp b/particle_structs/src/dps/dps_rebuild.hpp index 734ac21c..cf44f0d3 100644 --- a/particle_structs/src/dps/dps_rebuild.hpp +++ b/particle_structs/src/dps/dps_rebuild.hpp @@ -39,7 +39,7 @@ namespace pumipic { if (parent > -1) // count particles kept and move parentElms_cpy(soa*soa_len + tuple) = parent; else { // count particles deleted and delete - Kokkos::atomic_increment(&num_removed_d(0)); + Kokkos::atomic_inc(&num_removed_d(0)); active.access(soa,tuple) = false; // delete particles } } diff --git a/particle_structs/src/particle_structure.hpp b/particle_structs/src/particle_structure.hpp index 70346b5c..2acb5d93 100644 --- a/particle_structs/src/particle_structure.hpp +++ b/particle_structs/src/particle_structure.hpp @@ -22,15 +22,19 @@ namespace pumipic { typedef typename Space::memory_space memory_space; typedef typename Space::execution_space execution_space; typedef typename Space::device_type device_type; +#if KOKKOS_VERSION >= 40700 + typedef typename Kokkos::ViewTraits::host_mirror_space HostMirrorSpace; +#else typedef typename Kokkos::ViewTraits::HostMirrorSpace HostMirrorSpace; +#endif typedef ParticleStructure HostMirror; template using Mirror = ParticleStructure; template using View = Kokkos::View; typedef View kkLidView; typedef View kkGidView; - typedef typename kkLidView::HostMirror kkLidHostMirror; - typedef typename kkGidView::HostMirror kkGidHostMirror; + typedef typename kkLidView::host_mirror_type kkLidHostMirror; + typedef typename kkGidView::host_mirror_type kkGidHostMirror; template using DataType = typename MemberTypeAtIndex::type; diff --git a/particle_structs/src/ps_for.hpp b/particle_structs/src/ps_for.hpp index 77da60ff..27dec203 100644 --- a/particle_structs/src/ps_for.hpp +++ b/particle_structs/src/ps_for.hpp @@ -67,7 +67,7 @@ namespace pumipic { ViewT ppe("ppe", num_elems+1); auto setPPE = PS_LAMBDA(const lid_t& e, const lid_t& p, const bool& mask) { if (mask) { - Kokkos::atomic_increment(&ppe(e)); + Kokkos::atomic_inc(&ppe(e)); } }; parallel_for(this, setPPE, "setPPE"); diff --git a/particle_structs/src/scs/SCS_migrate.h b/particle_structs/src/scs/SCS_migrate.h index 6793338e..60bc649c 100644 --- a/particle_structs/src/scs/SCS_migrate.h +++ b/particle_structs/src/scs/SCS_migrate.h @@ -31,7 +31,7 @@ namespace pumipic { const lid_t process = new_process(particle_id); if (mask && (process != comm_rank)) { const lid_t process_index = dist.index(process); - Kokkos::atomic_increment(&num_send_particles(process_index)); + Kokkos::atomic_inc(&num_send_particles(process_index)); } }; parallel_for(count_sending_particles); diff --git a/particle_structs/src/scs/SCS_rebuild.h b/particle_structs/src/scs/SCS_rebuild.h index c0ed0881..b502244a 100644 --- a/particle_structs/src/scs/SCS_rebuild.h +++ b/particle_structs/src/scs/SCS_rebuild.h @@ -17,18 +17,18 @@ namespace pumipic { const bool is_moving = is_particle & (new_elem != element_id); if (is_moving && mask) { const lid_t new_row = element_to_row_local(new_elem); - Kokkos::atomic_increment(&(new_particles_per_row(new_row))); + Kokkos::atomic_inc(&(new_particles_per_row(new_row))); } particle_mask_local(particle_id) = is_particle; if (!is_particle) - Kokkos::atomic_increment(&(num_holes_per_row(row))); + Kokkos::atomic_inc(&(num_holes_per_row(row))); }; parallel_for(countNewParticles, "countNewParticles"); // Add new particles to counts Kokkos::parallel_for("reshuffle_count", new_particle_elements.size(), KOKKOS_LAMBDA(const lid_t& i) { const lid_t new_elem = new_particle_elements(i); const lid_t new_row = element_to_row_local(new_elem); - Kokkos::atomic_increment(&(new_particles_per_row(new_row))); + Kokkos::atomic_inc(&(new_particles_per_row(new_row))); }); //Check if the particles will fit in current structure @@ -133,7 +133,7 @@ namespace pumipic { auto countNewParticles = PS_LAMBDA(const lid_t& element_id, const lid_t& particle_id, const bool& mask){ const lid_t new_elem = new_element(particle_id); if (mask && new_elem != -1) - Kokkos::atomic_increment(&(new_particles_per_elem(new_elem))); + Kokkos::atomic_inc(&(new_particles_per_elem(new_elem))); }; parallel_for(countNewParticles, "countNewParticles"); @@ -153,7 +153,7 @@ namespace pumipic { // Add new particles to counts Kokkos::parallel_for("rebuild_count", new_particle_elements.size(), KOKKOS_LAMBDA(const lid_t& i) { const lid_t new_elem = new_particle_elements(i); - Kokkos::atomic_increment(&(new_particles_per_elem(new_elem))); + Kokkos::atomic_inc(&(new_particles_per_elem(new_elem))); }); //Reduce the count of particles diff --git a/particle_structs/src/scs/SellCSigma.h b/particle_structs/src/scs/SellCSigma.h index 9f4f6350..bfacd31c 100644 --- a/particle_structs/src/scs/SellCSigma.h +++ b/particle_structs/src/scs/SellCSigma.h @@ -407,7 +407,7 @@ void SellCSigma::printFormat(const char* prefix) const { kkGidHostMirror element_to_gid_host = deviceToHost(element_to_gid); kkLidHostMirror row_to_element_host = deviceToHost(row_to_element); kkLidHostMirror offsets_host = deviceToHost(offsets); - Kokkos::View::HostMirror particle_mask_host = deviceToHost(particle_mask); + Kokkos::View::host_mirror_type particle_mask_host = deviceToHost(particle_mask); std::stringstream ss; char buffer[1000]; diff --git a/particle_structs/src/support/psDistributor.hpp b/particle_structs/src/support/psDistributor.hpp index a1dd1d3c..a2628df0 100644 --- a/particle_structs/src/support/psDistributor.hpp +++ b/particle_structs/src/support/psDistributor.hpp @@ -33,7 +33,7 @@ namespace pumipic { typedef Kokkos::View IndexView; //List of ranks on the device IndexView ranks_d; - typename IndexView::HostMirror ranks_h; + typename IndexView::host_mirror_type ranks_h; //Unordered map from rank to index on device typedef Kokkos::UnorderedMap MapType; diff --git a/particle_structs/test/Distribute.cpp b/particle_structs/test/Distribute.cpp index 2880f506..fc92377c 100644 --- a/particle_structs/test/Distribute.cpp +++ b/particle_structs/test/Distribute.cpp @@ -248,14 +248,14 @@ void gitrm_distribution(int ne, int np, Kokkos::View ptcls_per_elem, auto generator = pool.get_state(); int index = generator.urand(0,cutoff); pool.free_state(generator); - Kokkos::atomic_increment(&ptcls_per_elem(index)); + Kokkos::atomic_inc(&ptcls_per_elem(index)); elem_per_ptcl(i) = index; }); Kokkos::parallel_for(ptcls_second, KOKKOS_LAMBDA(const int i) { auto generator = pool.get_state(); int index = generator.urand(cutoff,ne); pool.free_state(generator); - Kokkos::atomic_increment(&ptcls_per_elem(index)); + Kokkos::atomic_inc(&ptcls_per_elem(index)); elem_per_ptcl(ptcls_first + i) = index; }); } diff --git a/particle_structs/test/test_constructor.cpp b/particle_structs/test/test_constructor.cpp index 666bde5a..7f56df71 100644 --- a/particle_structs/test/test_constructor.cpp +++ b/particle_structs/test/test_constructor.cpp @@ -39,7 +39,7 @@ int testParticleExistence(const char* name, PS* structure, lid_t num_ptcls) { kkLidView count("count", 1); auto checkExistence = PS_LAMBDA(const lid_t& e, const lid_t& p, const bool& mask) { if (mask) - Kokkos::atomic_increment(&(count(0))); + Kokkos::atomic_inc(&(count(0))); }; ps::parallel_for(structure, checkExistence, "check particle existence"); lid_t c = ps::getLastValue(count); diff --git a/particle_structs/test/test_types.hpp b/particle_structs/test/test_types.hpp index ca052470..e0a9a0b3 100644 --- a/particle_structs/test/test_types.hpp +++ b/particle_structs/test/test_types.hpp @@ -23,5 +23,5 @@ typedef PS::kkGidHostMirror kkGidHost; template using KView=Kokkos::View; template -using KViewHost=typename KView::HostMirror; +using KViewHost=typename KView::host_mirror_type; using ps::lid_t; diff --git a/src/pumipic_kktypes.cpp b/src/pumipic_kktypes.cpp index 5420aa23..2fc43486 100644 --- a/src/pumipic_kktypes.cpp +++ b/src/pumipic_kktypes.cpp @@ -3,7 +3,7 @@ namespace pumipic { void hostToDeviceLid(kkLidView d, lid_t *h) { - kkLidView::HostMirror hv = Kokkos::create_mirror_view(d); + kkLidView::host_mirror_type hv = Kokkos::create_mirror_view(d); for (size_t i=0; i{}, T>::type; template typename - IsKokkosView::HostMirror create_mirror_view(ViewT v) { + IsKokkosView::host_mirror_type create_mirror_view(ViewT v) { return Kokkos::create_mirror_view(v); } template typename - IsPPView::HostMirror create_mirror_view(ViewT v) { - return typename ViewT::HostMirror(Kokkos::create_mirror_view(v.view())); + IsPPView::host_mirror_type create_mirror_view(ViewT v) { + return typename ViewT::host_mirror_type(Kokkos::create_mirror_view(v.view())); } template @@ -54,7 +54,7 @@ namespace pumipic { } template - typename ViewT::HostMirror deviceToHost(ViewT view) { + typename ViewT::host_mirror_type deviceToHost(ViewT view) { auto hv = create_mirror_view(view); deep_copy(hv, view); return hv; @@ -62,20 +62,20 @@ namespace pumipic { template typename std::enable_if::type - hostToDevice(typename ViewT::HostMirror hv, ViewT, T* data) { + hostToDevice(typename ViewT::host_mirror_type hv, ViewT, T* data) { for (size_t i = 0; i < hv.extent(0); ++i) hv(i) = data[i]; } template typename std::enable_if::type - hostToDevice(typename ViewT::HostMirror hv, ViewT, T* data) { + hostToDevice(typename ViewT::host_mirror_type hv, ViewT, T* data) { for (size_t i = 0; i < hv.extent(0); ++i) for (size_t j = 0; j < hv.extent(1); ++j) hv(i,j) = data[i][j]; } template typename std::enable_if::type - hostToDevice(typename ViewT::HostMirror hv, ViewT, T* data) { + hostToDevice(typename ViewT::host_mirror_type hv, ViewT, T* data) { for (size_t i = 0; i < hv.extent(0); ++i) for (size_t j = 0; j < hv.extent(1); ++j) for (size_t k = 0; k < hv.extent(2); ++k) @@ -83,7 +83,7 @@ namespace pumipic { } template typename std::enable_if::type - hostToDevice(typename ViewT::HostMirror hv, ViewT, T* data) { + hostToDevice(typename ViewT::host_mirror_type hv, ViewT, T* data) { for (size_t i = 0; i < hv.extent(0); ++i) for (size_t j = 0; j < hv.extent(1); ++j) for (size_t k = 0; k < hv.extent(2); ++k) @@ -207,7 +207,7 @@ namespace pumipic { template struct HostToDevice { HostToDevice(Kokkos::View view, T* data) { - typename Kokkos::View::HostMirror hv = Kokkos::create_mirror_view(view); + typename Kokkos::View::host_mirror_type hv = Kokkos::create_mirror_view(view); for (size_t i = 0; i < hv.size(); ++i) hv(i) = data[i]; Kokkos::deep_copy(view, hv); @@ -215,7 +215,7 @@ namespace pumipic { }; template struct HostToDevice { HostToDevice(Kokkos::View view, T (*data)[N]) { - typename Kokkos::View::HostMirror hv = Kokkos::create_mirror_view(view); + typename Kokkos::View::host_mirror_type hv = Kokkos::create_mirror_view(view); for (size_t i = 0; i < hv.extent(0); ++i) for (size_t j = 0; j < N; ++j) hv(i,j) = data[i][j]; diff --git a/support/ViewComm_gpu.hpp b/support/ViewComm_gpu.hpp index 43c11f50..0d125077 100644 --- a/support/ViewComm_gpu.hpp +++ b/support/ViewComm_gpu.hpp @@ -34,7 +34,7 @@ MpiType > >::mpitype(), sender, tag, comm, MPI_STATUS_IGNORE); #else - typename ViewT::HostMirror view_host = create_mirror_view(new_view); + typename ViewT::host_mirror_type view_host = create_mirror_view(new_view); int ret = MPI_Recv(view_host.data(), view_host.size(), MpiType > >::mpitype(), sender, tag, comm, MPI_STATUS_IGNORE); @@ -83,7 +83,7 @@ tag, comm, req); #else int size_per_entry = BaseType >::size; - typename ViewT::HostMirror view_host = create_mirror_view(new_view); + typename ViewT::host_mirror_type view_host = create_mirror_view(new_view); int ret = MPI_Irecv(view_host.data(), size * size_per_entry, MpiType > >::mpitype(), sender, tag, comm, req); @@ -135,8 +135,8 @@ return MPI_Alltoall(send.data(), send_size, MpiType > >::mpitype(), recv.data(), recv_size, MpiType > >::mpitype(), comm); #else - typename ViewT::HostMirror send_host = deviceToHost(send); - typename ViewT::HostMirror recv_host = create_mirror_view(recv); + typename ViewT::host_mirror_type send_host = deviceToHost(send); + typename ViewT::host_mirror_type recv_host = create_mirror_view(recv); int ret = MPI_Alltoall(send_host.data(), send_size, MpiType > >::mpitype(), recv_host.data(), recv_size, MpiType > >::mpitype(), comm); deep_copy(recv, recv_host); @@ -154,8 +154,8 @@ IsGPU > PS_Comm_Ialltoall(ViewT send, int send_size, recv.data(), recv_size, MpiType > >::mpitype(), comm, request); #else - typename ViewT::HostMirror send_host = deviceToHost(send); - typename ViewT::HostMirror recv_host = create_mirror_view(recv); + typename ViewT::host_mirror_type send_host = deviceToHost(send); + typename ViewT::host_mirror_type recv_host = create_mirror_view(recv); int ret = MPI_Ialltoall(send_host.data(), send_size, MpiType > >::mpitype(), recv_host.data(), recv_size, @@ -178,8 +178,8 @@ IsGPU > PS_Comm_Reduce(ViewT send_view, ViewT recv_view, int co MpiType > >::mpitype(), op, root, comm); #else - typename ViewT::HostMirror send_host = deviceToHost(send_view); - typename ViewT::HostMirror recv_host = create_mirror_view(recv_view); + typename ViewT::host_mirror_type send_host = deviceToHost(send_view); + typename ViewT::host_mirror_type recv_host = create_mirror_view(recv_view); int ret = MPI_Reduce(send_host.data(), recv_host.data(), count, MpiType > >::mpitype(), op, root, comm); @@ -200,8 +200,8 @@ IsGPU > PS_Comm_Allreduce(ViewT send_view, ViewT recv_view, int return MPI_Allreduce(send_view.data(), recv_view.data(), count, MpiType > >::mpitype(),op, comm); #else - typename ViewT::HostMirror send_host = deviceToHost(send_view); - typename ViewT::HostMirror recv_host = create_mirror_view(recv_view); + typename ViewT::host_mirror_type send_host = deviceToHost(send_view); + typename ViewT::host_mirror_type recv_host = create_mirror_view(recv_view); int ret = MPI_Allreduce(send_host.data(), recv_host.data(), count, MpiType > >::mpitype(), op, comm); deep_copy(recv_view, recv_host); diff --git a/support/ppView.h b/support/ppView.h index a784f9d0..7c488ff4 100644 --- a/support/ppView.h +++ b/support/ppView.h @@ -13,7 +13,8 @@ namespace pumipic { typedef typename KView::device_type device_type; typedef typename KView::data_type data_type; typedef typename KView::value_type value_type; - typedef View HostMirror; + typedef View host_mirror_type; + [[deprecated("Use host_mirror_type instead")]] typedef host_mirror_type HostMirror; View() : view_() {} View(lid_t size) : view_("ppView", size) {} View(std::string name, lid_t size) : view_(name, size) {}