From 10b0a0dd6974209b1ba582aa56238a3d541341ca Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Tue, 23 Jul 2019 14:53:26 +0000 Subject: [PATCH] Add import() from one LA::d::Vector to another one --- include/deal.II/lac/la_parallel_vector.h | 15 ++ .../lac/la_parallel_vector.templates.h | 20 +++ .../deal.II/lac/vector_operations_internal.h | 131 ++++++++++++++++++ source/lac/la_parallel_vector.cu | 32 +++++ source/lac/la_parallel_vector.inst.in | 5 + 5 files changed, 203 insertions(+) diff --git a/include/deal.II/lac/la_parallel_vector.h b/include/deal.II/lac/la_parallel_vector.h index c2bef47916..29ee2ff50a 100644 --- a/include/deal.II/lac/la_parallel_vector.h +++ b/include/deal.II/lac/la_parallel_vector.h @@ -614,6 +614,21 @@ namespace LinearAlgebra void copy_locally_owned_data_from(const Vector &src); + /** + * Import all the elements present in the distributed vector @p src. + * VectorOperation::values @p operation is used to decide if the elements + * in @p V should be added to the current vector or replace the current + * elements. The main purpose of this function is to get data from one + * memory space, e.g. CUDA, to the other, e.g. the Host. + * + * @note The partitioners of the two distributed vectors need to be the + * same as no MPI communication is performed. + */ + template + void + import(const Vector &src, + VectorOperation::values operation); + //@} /** diff --git a/include/deal.II/lac/la_parallel_vector.templates.h b/include/deal.II/lac/la_parallel_vector.templates.h index 743aac8780..b05ab1ae98 100644 --- a/include/deal.II/lac/la_parallel_vector.templates.h +++ b/include/deal.II/lac/la_parallel_vector.templates.h @@ -727,6 +727,26 @@ namespace LinearAlgebra + template + template + void + Vector::import( + const Vector &src, + VectorOperation::values operation) + { + Assert(src.partitioner.get() != nullptr, ExcNotInitialized()); + Assert(partitioner->locally_owned_range() == + src.partitioner->locally_owned_range(), + ExcMessage("Locally owned indices should be identical.")); + Assert(partitioner->ghost_indices() == src.partitioner->ghost_indices(), + ExcMessage("Ghost indices should be identical.")); + ::dealii::internal::VectorOperations:: + functions::import( + thread_loop_partitioner, allocated_size, operation, src.data, data); + } + + + #ifdef DEAL_II_WITH_PETSC namespace petsc_helpers diff --git a/include/deal.II/lac/vector_operations_internal.h b/include/deal.II/lac/vector_operations_internal.h index 089ab3b761..877f40d9a0 100644 --- a/include/deal.II/lac/vector_operations_internal.h +++ b/include/deal.II/lac/vector_operations_internal.h @@ -26,6 +26,7 @@ #include #include +#include #include #include @@ -1657,6 +1658,18 @@ namespace internal { return Number(); } + + template + static void + import( + const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> & + /*thread_loop_partitioner*/, + const size_type /*size*/, + VectorOperation::values /*operation*/, + const ::dealii::MemorySpace::MemorySpaceData + & /*v_data*/, + ::dealii::MemorySpace::MemorySpaceData & /*data*/) + {} }; @@ -2002,6 +2015,66 @@ namespace internal return sum; } + + template + static void + import(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> + & thread_loop_partitioner, + const size_type size, + VectorOperation::values operation, + const ::dealii::MemorySpace::MemorySpaceData + &v_data, + ::dealii::MemorySpace::MemorySpaceData + &data, + typename std::enable_if< + std::is_same::value, + int>::type = 0) + { + if (operation == VectorOperation::insert) + { + copy(thread_loop_partitioner, size, v_data, data); + } + else if (operation == VectorOperation::add) + { + add_vector(thread_loop_partitioner, size, v_data, data); + } + else + { + AssertThrow(false, ExcNotImplemented()); + } + } + +#ifdef DEAL_II_COMPILER_CUDA_AWARE + template + static void + import(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> + & /*thread_loop_partitioner*/, + const size_type size, + VectorOperation::values operation, + const ::dealii::MemorySpace::MemorySpaceData + &v_data, + ::dealii::MemorySpace::MemorySpaceData + &data, + typename std::enable_if< + std::is_same::value, + int>::type = 0) + { + if (operation == VectorOperation::insert) + { + cudaError_t cuda_error_code = cudaMemcpy(data.values.get(), + v_data.values_dev.get(), + size * sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error_code); + } + else + { + AssertThrow(false, ExcNotImplemented()); + } + } +#endif }; @@ -2542,6 +2615,64 @@ namespace internal return res; } + + template + static void + import(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> + & thread_loop_partitioner, + const size_type size, + VectorOperation::values operation, + const ::dealii::MemorySpace::MemorySpaceData + &v_data, + ::dealii::MemorySpace::MemorySpaceData + &data, + typename std::enable_if< + std::is_same::value, + int>::type = 0) + { + if (operation == VectorOperation::insert) + { + copy(thread_loop_partitioner, size, v_data, data); + } + else if (operation == VectorOperation::add) + { + add_vector(thread_loop_partitioner, size, v_data, data); + } + else + { + AssertThrow(false, ExcNotImplemented()); + } + } + + template + static void + import(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> + & /*thread_loop_partitioner*/, + const size_type size, + VectorOperation::values operation, + const ::dealii::MemorySpace::MemorySpaceData + &v_data, + ::dealii::MemorySpace::MemorySpaceData + &data, + typename std::enable_if< + std::is_same::value, + int>::type = 0) + { + if (operation == VectorOperation::insert) + { + cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.get(), + v_data.values.get(), + size * sizeof(Number), + cudaMemcpyHostToDevice); + AssertCuda(cuda_error_code); + } + else + { + AssertThrow(false, ExcNotImplemented()); + } + } }; #endif } // namespace VectorOperations diff --git a/source/lac/la_parallel_vector.cu b/source/lac/la_parallel_vector.cu index 352cf2c25c..6f7aa7286f 100644 --- a/source/lac/la_parallel_vector.cu +++ b/source/lac/la_parallel_vector.cu @@ -25,6 +25,38 @@ namespace LinearAlgebra { template class Vector; template class Vector; + template void + Vector::import< + ::dealii::MemorySpace::CUDA>( + const Vector &, + VectorOperation::values); + template void + Vector::import< + ::dealii::MemorySpace::CUDA>( + const Vector &, + VectorOperation::values); + + template void + Vector::import< + ::dealii::MemorySpace::Host>( + const Vector &, + VectorOperation::values); + template void + Vector::import< + ::dealii::MemorySpace::Host>( + const Vector &, + VectorOperation::values); + + template void + Vector::import< + ::dealii::MemorySpace::CUDA>( + const Vector &, + VectorOperation::values); + template void + Vector::import< + ::dealii::MemorySpace::CUDA>( + const Vector &, + VectorOperation::values); } // namespace distributed } // namespace LinearAlgebra diff --git a/source/lac/la_parallel_vector.inst.in b/source/lac/la_parallel_vector.inst.in index e50f3d68cc..be711c7a55 100644 --- a/source/lac/la_parallel_vector.inst.in +++ b/source/lac/la_parallel_vector.inst.in @@ -22,6 +22,11 @@ for (SCALAR : REAL_AND_COMPLEX_SCALARS) namespace distributed \{ template class Vector; + template void + Vector::import< + ::dealii::MemorySpace::Host>( + const Vector &, + VectorOperation::values); \} \} } -- 2.39.5