Replace for_each with Kokkos::parallel_for

source/MaterialProperty.templates.hh

@@ -221,13 +221,14 @@ void MaterialProperty<dim, MemorySpaceType>::reinit_dofs()
 #ifdef ADAMANTINE_DEBUG
   if constexpr (std::is_same_v<MemorySpaceType, dealii::MemorySpace::Host>)
   {
-    for_each(MemorySpaceType{}, g_n_material_states,
-             [=](int i) mutable
-             {
-               for (unsigned int j = 0; j < _dofs_map.size(); ++j)
-                 _state(i, j) = std::numeric_limits<double>::signaling_NaN();
-               ;
-             });
+    Kokkos::parallel_for(
+        "adamantine::set_state_nan",
+        Kokkos::MDRangePolicy<Kokkos::DefaultHostExecutionSpace,
+                              Kokkos::Rank<2>>(
+            {0, 0}, {g_n_material_states, _dofs_map.size()}),
+        KOKKOS_LAMBDA(int i, int j) {
+          _state(i, j) = std::numeric_limits<double>::signaling_NaN();
+        });
   }
 #endif
 }
@@ -270,8 +271,13 @@ void MaterialProperty<dim, MemorySpaceType>::update(
   double *temperature_average_local = temperature_average.get_values();
 
   bool use_table = _use_table;
-  for_each(
-      MemorySpaceType{}, material_ids_size, KOKKOS_LAMBDA(int i) {
+  using ExecutionSpace = std::conditional_t<
+      std::is_same_v<MemorySpaceType, dealii::MemorySpace::Host>,
+      Kokkos::DefaultHostExecutionSpace, Kokkos::DefaultExecutionSpace>;
+  Kokkos::parallel_for(
+      "adamantine::update_material_properties",
+      Kokkos::RangePolicy<ExecutionSpace>(0, material_ids_size),
+      KOKKOS_LAMBDA(int i) {
         unsigned int constexpr liquid =
             static_cast<unsigned int>(MaterialState::liquid);
         unsigned int constexpr powder =
@@ -658,8 +664,13 @@ void MaterialProperty<dim, MemorySpaceType>::set_state_device(
   auto const powder_state = static_cast<unsigned int>(MaterialState::powder);
   auto const liquid_state = static_cast<unsigned int>(MaterialState::liquid);
   auto const solid_state = static_cast<unsigned int>(MaterialState::solid);
-  for_each(
-      MemorySpaceType{}, cell_i, KOKKOS_LAMBDA(int i) mutable {
+  using ExecutionSpace = std::conditional_t<
+      std::is_same_v<MemorySpaceType, dealii::MemorySpace::Host>,
+      Kokkos::DefaultHostExecutionSpace, Kokkos::DefaultExecutionSpace>;
+  Kokkos::parallel_for(
+      "adamantine::set_state_device",
+      Kokkos::RangePolicy<ExecutionSpace>(0, cell_i),
+      KOKKOS_LAMBDA(int i) mutable {
         double liquid_ratio_sum = 0.;
         double powder_ratio_sum = 0.;
         for (unsigned int q = 0; q < n_q_points; ++q)
@@ -706,10 +717,13 @@ void MaterialProperty<dim, MemorySpaceType>::set_initial_state()
       Kokkos::create_mirror_view_and_copy(memory_space, user_indices_host);
 
   Kokkos::deep_copy(_state, 0.);
-  for_each(
-      MemorySpaceType{}, user_indices.extent(0), KOKKOS_LAMBDA(int i) mutable {
-        _state(user_indices(i), mp_dofs(i)) = 1.;
-      });
+  using ExecutionSpace = std::conditional_t<
+      std::is_same_v<MemorySpaceType, dealii::MemorySpace::Host>,
+      Kokkos::DefaultHostExecutionSpace, Kokkos::DefaultExecutionSpace>;
+  Kokkos::parallel_for(
+      "adamantine::set_initial_state",
+      Kokkos::RangePolicy<ExecutionSpace>(0, user_indices.extent(0)),
+      KOKKOS_LAMBDA(int i) { _state(user_indices(i), mp_dofs(i)) = 1.; });
 }
 
 template <int dim, typename MemorySpaceType>

source/ThermalPhysics.templates.hh

@@ -840,14 +840,18 @@ void ThermalPhysics<dim, fe_degree, MemorySpaceType, QuadratureType>::
   // cells are at the same level than their neighbors.
   rw_solution.reinit(solution.locally_owned_elements());
   rw_solution.import(solution, dealii::VectorOperation::insert);
-  std::for_each(rw_solution.begin(), rw_solution.end(),
-                [&](double &val)
-                {
-                  if (val == std::numeric_limits<double>::infinity())
-                  {
-                    val = new_material_temperature;
-                  }
-                });
+  Kokkos::parallel_for("adamantine::set_new_material_temperature",
+                       Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(
+                           0, rw_solution.locally_owned_size()),
+                       [&](int i)
+                       {
+                         if (rw_solution.local_element(i) ==
+                             std::numeric_limits<double>::infinity())
+                         {
+                           rw_solution.local_element(i) =
+                               new_material_temperature;
+                         }
+                       });
   solution.import(rw_solution, dealii::VectorOperation::insert);
 }
 

source/utils.hh

@@ -8,50 +8,19 @@
 #ifndef UTILS_HH
 #define UTILS_HH
 
-#include <deal.II/base/cuda.h>
-#include <deal.II/base/cuda_size.h>
 #include <deal.II/base/exceptions.h>
 
 #include <cassert>
 #include <cstring>
 #include <exception>
 #include <filesystem>
 #include <iostream>
-#include <numeric>
 #include <stdexcept>
 #include <string>
 #include <utility>
 
 namespace adamantine
 {
-template <typename Functor>
-void for_each(dealii::MemorySpace::Host, unsigned int const size, Functor f)
-{
-  for (unsigned int i = 0; i < size; ++i)
-    f(i);
-}
-
-#ifdef __CUDACC__
-template <typename Functor>
-__global__ void for_each_impl(int size, Functor f)
-{
-  int i = blockDim.x * blockIdx.x + threadIdx.x;
-  if (i < size)
-  {
-    f(i);
-  }
-}
-
-template <typename Functor>
-void for_each(dealii::MemorySpace::Default, unsigned int const size,
-              Functor const &f)
-{
-  const int n_blocks = 1 + size / dealii::CUDAWrappers::block_size;
-  for_each_impl<<<n_blocks, dealii::CUDAWrappers::block_size>>>(size, f);
-  AssertCudaKernel();
-}
-#endif
-
 /**
  * Wait for the file to appear.
  */