From: Bruno Turcksin <bruno.turcksin@gmail.com>
Date: Tue, 21 Mar 2023 15:09:59 +0000 (-0400)
Subject: Move constrained_dofs to Kokkos::View
X-Git-Tag: v9.5.0-rc1~349^2~7
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=2e338966c21de429a7444d0f0cf55c4bfe46740a;p=dealii.git

Move constrained_dofs to Kokkos::View
---

diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h
index 108274852b..b3d1238624 100644
--- a/include/deal.II/matrix_free/cuda_matrix_free.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.h
@@ -536,9 +536,10 @@ namespace CUDAWrappers
       JxW;
 
     /**
-     * Pointer to the constrained degrees of freedom.
+     * Kokkos::View to the constrained degrees of freedom.
      */
-    types::global_dof_index *constrained_dofs;
+    Kokkos::View<types::global_dof_index *, MemorySpace::Default::kokkos_space>
+      constrained_dofs;
 
     /**
      * Mask deciding where constraints are set on a given cell.
diff --git a/include/deal.II/matrix_free/cuda_matrix_free.templates.h b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
index 7ebf79bd91..5dd210b3f1 100644
--- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
@@ -607,7 +607,6 @@ namespace CUDAWrappers
   MatrixFree<dim, Number>::MatrixFree()
     : my_id(-1)
     , n_dofs(0)
-    , constrained_dofs(nullptr)
     , padding_length(0)
     , dof_handler(nullptr)
   {}
@@ -723,8 +722,6 @@ namespace CUDAWrappers
       Utilities::CUDA::free(constraint_mask_color_ptr);
     constraint_mask.clear();
 
-    Utilities::CUDA::free(constrained_dofs);
-
     internal::used_objects[my_id].store(false);
     my_id = -1;
   }
@@ -744,7 +741,7 @@ namespace CUDAWrappers
            ExcMessage("src and dst vectors have different size."));
     // FIXME When using C++17, we can use KOKKOS_CLASS_LAMBDA and this
     // work-around can be removed.
-    types::global_dof_index *constr_dofs = constrained_dofs;
+    auto               constr_dofs = constrained_dofs;
     const unsigned int size = internal::VectorLocalSize<VectorType>::get(dst);
     const Number *     src_ptr = src.get_values();
     Number *           dst_ptr = dst.get_values();
@@ -775,7 +772,7 @@ namespace CUDAWrappers
     Number *dst_ptr = dst.get_values();
     // FIXME When using C++17, we can use KOKKOS_CLASS_LAMBDA and this
     // work-around can be removed.
-    types::global_dof_index *constr_dofs = constrained_dofs;
+    auto constr_dofs = constrained_dofs;
     // When working with distributed vectors, the constrained dofs are
     // computed for ghosted vectors but we want to set the values of the
     // constrained dofs of non-ghosted vectors.
@@ -1144,17 +1141,17 @@ namespace CUDAWrappers
               }
           }
 
-        cuda_error = cudaMalloc(&constrained_dofs,
-                                n_constrained_dofs *
-                                  sizeof(dealii::types::global_dof_index));
-        AssertCuda(cuda_error);
+        constrained_dofs = Kokkos::View<types::global_dof_index *,
+                                        MemorySpace::Default::kokkos_space>(
+          Kokkos::view_alloc("constrained_dofs", Kokkos::WithoutInitializing),
+          n_constrained_dofs);
 
-        cuda_error = cudaMemcpy(constrained_dofs,
-                                constrained_dofs_host.data(),
-                                n_constrained_dofs *
-                                  sizeof(dealii::types::global_dof_index),
-                                cudaMemcpyHostToDevice);
-        AssertCuda(cuda_error);
+        Kokkos::View<types::global_dof_index *,
+                     MemorySpace::Default::kokkos_space,
+                     Kokkos::MemoryTraits<Kokkos::Unmanaged>>
+          constrained_dofs_host_view(constrained_dofs_host.data(),
+                                     constrained_dofs_host.size());
+        Kokkos::deep_copy(constrained_dofs, constrained_dofs_host_view);
       }
   }