From 2e338966c21de429a7444d0f0cf55c4bfe46740a Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Tue, 21 Mar 2023 11:09:59 -0400 Subject: [PATCH] Move constrained_dofs to Kokkos::View --- .../deal.II/matrix_free/cuda_matrix_free.h | 5 ++-- .../matrix_free/cuda_matrix_free.templates.h | 27 +++++++++---------- 2 files changed, 15 insertions(+), 17 deletions(-) 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 + 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::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::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( + 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> + constrained_dofs_host_view(constrained_dofs_host.data(), + constrained_dofs_host.size()); + Kokkos::deep_copy(constrained_dofs, constrained_dofs_host_view); } } -- 2.39.5