Add import() from one LA::d::Vector to another one

diff --git a/include/deal.II/lac/la_parallel_vector.h b/include/deal.II/lac/la_parallel_vector.h
index c2bef479162c3f153fac90e1f3874925ce0ed85b..29ee2ff50af6b705d34b19221d02eda5e89f77e5 100644 (file)
--- 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<Number2, MemorySpace> &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 <typename MemorySpace2>
+      void
+      import(const Vector<Number, MemorySpace2> &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 743aac8780c80c77b861d4175064510f636f84cd..b05ab1ae980070a843fc663eb06bbecb04c07cfc 100644 (file)
--- 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 <typename Number, typename MemorySpaceType>
+    template <typename MemorySpaceType2>
+    void
+    Vector<Number, MemorySpaceType>::import(
+      const Vector<Number, MemorySpaceType2> &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<Number, Number, MemorySpaceType>::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 089ab3b761b6007f76c99fe563ba87758ac7ccfb..877f40d9a02eb7670508377e98cd767a113b5730 100644 (file)
--- a/include/deal.II/lac/vector_operations_internal.h
+++ b/include/deal.II/lac/vector_operations_internal.h
@@ -26,6 +26,7 @@
 
 #include <deal.II/lac/cuda_kernels.h>
 #include <deal.II/lac/cuda_kernels.templates.h>
+#include <deal.II/lac/vector_operation.h>
 
 #include <cstdio>
 #include <cstring>
@@ -1657,6 +1658,18 @@ namespace internal
       {
         return Number();
       }
+
+      template <typename MemorySpace2>
+      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<Number, MemorySpace2>
+          & /*v_data*/,
+        ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace> & /*data*/)
+      {}
     };
 
 
@@ -2002,6 +2015,66 @@ namespace internal
 
         return sum;
       }
+
+      template <typename MemorySpace2>
+      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<Number, MemorySpace2>
+               &v_data,
+             ::dealii::MemorySpace::MemorySpaceData<Number,
+                                                    ::dealii::MemorySpace::Host>
+               &data,
+             typename std::enable_if<
+               std::is_same<MemorySpace2, dealii::MemorySpace::Host>::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 <typename MemorySpace2>
+      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<Number, MemorySpace2>
+               &v_data,
+             ::dealii::MemorySpace::MemorySpaceData<Number,
+                                                    ::dealii::MemorySpace::Host>
+               &data,
+             typename std::enable_if<
+               std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::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 <typename MemorySpace2>
+      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<Number, MemorySpace2>
+               &v_data,
+             ::dealii::MemorySpace::MemorySpaceData<Number,
+                                                    ::dealii::MemorySpace::CUDA>
+               &data,
+             typename std::enable_if<
+               std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::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 <typename MemorySpace2>
+      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<Number, MemorySpace2>
+               &v_data,
+             ::dealii::MemorySpace::MemorySpaceData<Number,
+                                                    ::dealii::MemorySpace::CUDA>
+               &data,
+             typename std::enable_if<
+               std::is_same<MemorySpace2, ::dealii::MemorySpace::Host>::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 352cf2c25c606d3ad333b5ca90c0281f33aa9b20..6f7aa7286f5525359c3c1ebdd57158f884d0d3be 100644 (file)
--- a/source/lac/la_parallel_vector.cu
+++ b/source/lac/la_parallel_vector.cu
@@ -25,6 +25,38 @@ namespace LinearAlgebra
   {
     template class Vector<float, ::dealii::MemorySpace::CUDA>;
     template class Vector<double, ::dealii::MemorySpace::CUDA>;
+    template void
+    Vector<float, ::dealii::MemorySpace::Host>::import<
+      ::dealii::MemorySpace::CUDA>(
+      const Vector<float, ::dealii::MemorySpace::CUDA> &,
+      VectorOperation::values);
+    template void
+    Vector<double, ::dealii::MemorySpace::Host>::import<
+      ::dealii::MemorySpace::CUDA>(
+      const Vector<double, ::dealii::MemorySpace::CUDA> &,
+      VectorOperation::values);
+
+    template void
+    Vector<float, ::dealii::MemorySpace::CUDA>::import<
+      ::dealii::MemorySpace::Host>(
+      const Vector<float, ::dealii::MemorySpace::Host> &,
+      VectorOperation::values);
+    template void
+    Vector<double, ::dealii::MemorySpace::CUDA>::import<
+      ::dealii::MemorySpace::Host>(
+      const Vector<double, ::dealii::MemorySpace::Host> &,
+      VectorOperation::values);
+
+    template void
+    Vector<float, ::dealii::MemorySpace::CUDA>::import<
+      ::dealii::MemorySpace::CUDA>(
+      const Vector<float, ::dealii::MemorySpace::CUDA> &,
+      VectorOperation::values);
+    template void
+    Vector<double, ::dealii::MemorySpace::CUDA>::import<
+      ::dealii::MemorySpace::CUDA>(
+      const Vector<double, ::dealii::MemorySpace::CUDA> &,
+      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 e50f3d68ccf93693adc99246da406d02ee6d93bd..be711c7a551d5e96574671ad408c68e506e5d052 100644 (file)
--- 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<SCALAR, ::dealii::MemorySpace::Host>;
+        template void
+        Vector<SCALAR, ::dealii::MemorySpace::Host>::import<
+          ::dealii::MemorySpace::Host>(
+          const Vector<SCALAR, ::dealii::MemorySpace::Host> &,
+          VectorOperation::values);
       \}
     \}
   }