Remove deprecated code usage for Kokkos > 4.7.00

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<lid_t>(&num_send_particles(process_index));
+        Kokkos::atomic_inc<lid_t>(&num_send_particles(process_index));
       }
     };
     parallel_for(count_sending_particles);

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<lid_t>(&elmDegree_d(parent));
+          Kokkos::atomic_inc<lid_t>(&elmDegree_d(parent));
         else // count particles to be deleted
-          Kokkos::atomic_increment<lid_t>(&num_removed_d(0));
+          Kokkos::atomic_inc<lid_t>(&num_removed_d(0));
       }
     };
     Cabana::SimdPolicy<soa_len,execution_space> simd_policy(0, capacity_);

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<lid_t>(&num_send_particles(process_index));
+        Kokkos::atomic_inc<lid_t>(&num_send_particles(process_index));
       }
     };
     parallel_for(count_sending_particles);

particle_structs/src/csr/CSR_rebuild.hpp

@@ -30,17 +30,17 @@ 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
 
     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());
 

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<lid_t>(&num_send_particles(process_index));
+        Kokkos::atomic_inc<lid_t>(&num_send_particles(process_index));
       }
     };
     parallel_for(count_sending_particles);

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<lid_t>(&num_removed_d(0));
+          Kokkos::atomic_inc<lid_t>(&num_removed_d(0));
           active.access(soa,tuple) = false; // delete particles
         }
       }

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<void, Space>::host_mirror_space HostMirrorSpace;
+#else
     typedef typename Kokkos::ViewTraits<void, Space>::HostMirrorSpace HostMirrorSpace;
+#endif
     typedef ParticleStructure<DataTypes, HostMirrorSpace> HostMirror;
     template <typename Space2> using Mirror = ParticleStructure<DataTypes, Space2>;
 
     template <class T> using View = Kokkos::View<T*, device_type>;
     typedef View<lid_t> kkLidView;
     typedef View<gid_t> 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 <std::size_t N> using DataType =
       typename MemberTypeAtIndex<N, DataTypes>::type;

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");

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);

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<lid_t>(&(new_particles_per_row(new_row)));
+        Kokkos::atomic_inc<lid_t>(&(new_particles_per_row(new_row)));
       }
       particle_mask_local(particle_id) = is_particle;
       if (!is_particle)
-        Kokkos::atomic_increment<lid_t>(&(num_holes_per_row(row)));
+        Kokkos::atomic_inc<lid_t>(&(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<lid_t>(&(new_particles_per_row(new_row)));
+        Kokkos::atomic_inc<lid_t>(&(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<lid_t>(&(new_particles_per_elem(new_elem)));
+        Kokkos::atomic_inc<lid_t>(&(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<lid_t>(&(new_particles_per_elem(new_elem)));
+        Kokkos::atomic_inc<lid_t>(&(new_particles_per_elem(new_elem)));
       });
 
     //Reduce the count of particles

particle_structs/src/scs/SellCSigma.h

@@ -407,7 +407,7 @@ void SellCSigma<DataTypes,MemSpace>::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<bool*>::HostMirror particle_mask_host = deviceToHost(particle_mask);
+  Kokkos::View<bool*>::host_mirror_type particle_mask_host = deviceToHost(particle_mask);
 
   std::stringstream ss;
   char buffer[1000];

particle_structs/src/support/psDistributor.hpp

@@ -33,7 +33,7 @@ namespace pumipic {
     typedef Kokkos::View<int*, typename Space::device_type> 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<lid_t, lid_t, typename Space::device_type> MapType;

particle_structs/test/Distribute.cpp

@@ -248,14 +248,14 @@ void gitrm_distribution(int ne, int np, Kokkos::View<int*> ptcls_per_elem,
     auto generator = pool.get_state();
     int index = generator.urand(0,cutoff);
     pool.free_state(generator);
-    Kokkos::atomic_increment<int>(&ptcls_per_elem(index));
+    Kokkos::atomic_inc<int>(&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<int>(&ptcls_per_elem(index));
+    Kokkos::atomic_inc<int>(&ptcls_per_elem(index));
     elem_per_ptcl(ptcls_first + i) = index;
   });
 }

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<lid_t>(&(count(0)));
+      Kokkos::atomic_inc<lid_t>(&(count(0)));
   };
   ps::parallel_for(structure, checkExistence, "check particle existence");
   lid_t c = ps::getLastValue(count);

particle_structs/test/test_types.hpp

@@ -23,5 +23,5 @@ typedef PS::kkGidHostMirror kkGidHost;
 template <class T>
 using KView=Kokkos::View<T*, MemSpace::device_type>;
 template <class T>
-using KViewHost=typename KView<T>::HostMirror;
+using KViewHost=typename KView<T>::host_mirror_type;
 using ps::lid_t;

src/pumipic_kktypes.cpp

@@ -3,30 +3,30 @@
 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<hv.size(); ++i) {
       hv(i) = h[i];
     }
     Kokkos::deep_copy(d,hv);
   }
 
   void deviceToHostLid(kkLidView d, lid_t *h) {
-    kkLidView::HostMirror hv = Kokkos::create_mirror_view(d);
+    kkLidView::host_mirror_type hv = Kokkos::create_mirror_view(d);
     Kokkos::deep_copy(hv,d);
     for(size_t i=0; i<hv.size(); ++i) {
       h[i] = hv(i);
     }
   }
 
   void hostToDeviceFp(kkFpView d, fp_t* h) {
-    kkFpView::HostMirror hv = Kokkos::create_mirror_view(d);
+    kkFpView::host_mirror_type hv = Kokkos::create_mirror_view(d);
     for (size_t i=0; i<hv.size(); ++i)
       hv(i) = h[i];
     Kokkos::deep_copy(d,hv);
   }
 
   void hostToDeviceFp(kkFp3View d, fp_t (*h)[3]) {
-    kkFp3View::HostMirror hv = Kokkos::create_mirror_view(d);
+    kkFp3View::host_mirror_type hv = Kokkos::create_mirror_view(d);
     for (size_t i=0; i<hv.size()/3; ++i) {
       hv(i,0) = h[i][0];
       hv(i,1) = h[i][1];
@@ -36,7 +36,7 @@ namespace pumipic {
   }
 
   void deviceToHostFp(kkFp3View d, fp_t (*h)[3]) {
-    kkFp3View::HostMirror hv = Kokkos::create_mirror_view(d);
+    kkFp3View::host_mirror_type hv = Kokkos::create_mirror_view(d);
     Kokkos::deep_copy(hv,d);
     for(size_t i=0; i<hv.size()/3; ++i) {
       h[i][0] = hv(i,0);

support/SupportKK.h

@@ -36,12 +36,12 @@ namespace pumipic {
     typename std::enable_if<is_specialization<ViewT, View>{}, T>::type;
 
   template <class ViewT> typename
-  IsKokkosView<ViewT>::HostMirror create_mirror_view(ViewT v) {
+  IsKokkosView<ViewT>::host_mirror_type create_mirror_view(ViewT v) {
     return Kokkos::create_mirror_view(v);
   }
   template <class ViewT> typename
-  IsPPView<ViewT>::HostMirror create_mirror_view(ViewT v) {
-    return typename ViewT::HostMirror(Kokkos::create_mirror_view(v.view()));
+  IsPPView<ViewT>::host_mirror_type create_mirror_view(ViewT v) {
+    return typename ViewT::host_mirror_type(Kokkos::create_mirror_view(v.view()));
   }
 
   template <class ViewT, class ViewT2>
@@ -54,36 +54,36 @@ namespace pumipic {
   }
 
   template <class ViewT>
-  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;
   }
 
   template <class ViewT, class T>
   typename std::enable_if<ViewT::rank==1>::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 <class ViewT, class T>
   typename std::enable_if<ViewT::rank==2>::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 <class ViewT, class T>
   typename std::enable_if<ViewT::rank==3>::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)
           hv(i,j,k) = data[i][j][k];
   }
   template <class ViewT, class T>
   typename std::enable_if<ViewT::rank==4>::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,15 +207,15 @@ namespace pumipic {
 
   template <class T, typename Device> struct HostToDevice {
     HostToDevice(Kokkos::View<T*, Device> view, T* data) {
-      typename Kokkos::View<T*, Device>::HostMirror hv = Kokkos::create_mirror_view(view);
+      typename Kokkos::View<T*, Device>::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);
     }
   };
   template <class T, typename Device, std::size_t N> struct HostToDevice<T[N], Device> {
     HostToDevice(Kokkos::View<T*[N], Device> view, T (*data)[N]) {
-      typename Kokkos::View<T*[N], Device>::HostMirror hv = Kokkos::create_mirror_view(view);
+      typename Kokkos::View<T*[N], Device>::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];

support/ViewComm_gpu.hpp

@@ -34,7 +34,7 @@
                        MpiType<BT<ViewType<ViewT> > >::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<BT<ViewType<ViewT> > >::mpitype(),
                        sender, tag, comm, MPI_STATUS_IGNORE);
@@ -83,7 +83,7 @@
                         tag, comm, req);
 #else
     int size_per_entry = BaseType<ViewType<ViewT> >::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<BT<ViewType<ViewT> > >::mpitype(),
                         sender, tag, comm, req);
@@ -135,8 +135,8 @@
     return MPI_Alltoall(send.data(), send_size, MpiType<BT<ViewType<ViewT> > >::mpitype(),
                         recv.data(), recv_size, MpiType<BT<ViewType<ViewT> > >::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<BT<ViewType<ViewT> > >::mpitype(),
                            recv_host.data(), recv_size, MpiType<BT<ViewType<ViewT> > >::mpitype(), comm);
     deep_copy(recv, recv_host);
@@ -154,8 +154,8 @@ IsGPU<ViewSpace<ViewT> > PS_Comm_Ialltoall(ViewT send, int send_size,
                       recv.data(), recv_size, MpiType<BT<ViewType<ViewT> > >::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<BT<ViewType<ViewT> > >::mpitype(),
                           recv_host.data(), recv_size,
@@ -178,8 +178,8 @@ IsGPU<ViewSpace<ViewT> > PS_Comm_Reduce(ViewT send_view, ViewT recv_view, int co
                     MpiType<BT<ViewType<ViewT> > >::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<BT<ViewType<ViewT> > >::mpitype(),
                        op, root, comm);
@@ -200,8 +200,8 @@ IsGPU<ViewSpace<ViewT> > PS_Comm_Allreduce(ViewT send_view, ViewT recv_view, int
   return MPI_Allreduce(send_view.data(), recv_view.data(), count,
                        MpiType<BT<ViewType<ViewT> > >::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<BT<ViewType<ViewT> > >::mpitype(), op, comm);
   deep_copy(recv_view, recv_host);

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<T, typename KView::host_mirror_space, ArrayLayout> HostMirror;
+    typedef View<T, typename KView::host_mirror_space, ArrayLayout> 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) {}