]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Fix up
authorDaniel Arndt <arndtd@ornl.gov>
Tue, 6 Dec 2022 16:23:54 +0000 (11:23 -0500)
committerDaniel Arndt <arndtd@ornl.gov>
Thu, 8 Dec 2022 22:13:15 +0000 (17:13 -0500)
include/deal.II/base/array_view.h
include/deal.II/lac/la_parallel_vector.templates.h
include/deal.II/lac/vector_operations_internal.h
source/lac/CMakeLists.txt
source/lac/vector_memory.cc

index 50053b5330042302b87ebddee7ec3b80fe0a8733..168310404e1716c1c36679d98f43775334e1b14b 100644 (file)
@@ -345,45 +345,6 @@ private:
 //---------------------------------------------------------------------------
 
 
-namespace internal
-{
-  namespace ArrayViewHelper
-  {
-    template <typename MemorySpaceType>
-    inline bool
-    is_in_correct_memory_space(const void *const ptr)
-    {
-#ifndef DEAL_II_COMPILER_CUDA_AWARE
-      (void)ptr;
-      static_assert(std::is_same<MemorySpaceType, MemorySpace::Host>::value,
-                    "If the compiler doesn't understand CUDA code, "
-                    "the only possible memory space is 'MemorySpace::Host'!");
-      return true;
-#else
-      cudaPointerAttributes attributes;
-      const cudaError_t cuda_error = cudaPointerGetAttributes(&attributes, ptr);
-      if (cuda_error != cudaErrorInvalidValue)
-        {
-          AssertCuda(cuda_error);
-          if (std::is_same<MemorySpaceType, MemorySpace::Host>::value)
-            return (attributes.type == cudaMemoryTypeHost) ||
-                   (attributes.type == cudaMemoryTypeUnregistered);
-          else
-            return attributes.type == cudaMemoryTypeDevice;
-        }
-      else
-        {
-          // ignore and reset the error since host pointers produce an error
-          cudaGetLastError();
-          return std::is_same<MemorySpaceType, MemorySpace::Host>::value;
-        }
-#endif
-    }
-  } // namespace ArrayViewHelper
-} // namespace internal
-
-
-
 template <typename ElementType, typename MemorySpaceType>
 inline ArrayView<ElementType, MemorySpaceType>::ArrayView()
   : starting_element(nullptr)
@@ -398,14 +359,7 @@ inline ArrayView<ElementType, MemorySpaceType>::ArrayView(
   const std::size_t n_elements)
   : starting_element(starting_element)
   , n_elements(n_elements)
-{
-  Assert(
-    n_elements == 0 ||
-      internal::ArrayViewHelper::is_in_correct_memory_space<MemorySpaceType>(
-        starting_element),
-    ExcMessage("The memory space indicated by the template parameter "
-               "and the one derived from the pointer value do not match!"));
-}
+{}
 
 
 
index 179a15d8b3da1719f44bbc4905d21fd5161067f3..d4910ca6ef3d6e442e8ad1b4405e78916c2758e4 100644 (file)
@@ -86,9 +86,10 @@ namespace LinearAlgebra
       template <typename Number, typename MemorySpaceType>
       struct la_parallel_vector_templates_functions
       {
-        static_assert(std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
-                        std::is_same<MemorySpaceType, MemorySpace::Device>::value,
-                      "MemorySpace should be Host or CUDA");
+        static_assert(
+          std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
+            std::is_same<MemorySpaceType, MemorySpace::Device>::value,
+          "MemorySpace should be Host or Device");
 
         static void
         resize_val(
@@ -317,24 +318,27 @@ namespace LinearAlgebra
       };
 
       template <typename Number>
-      struct la_parallel_vector_templates_functions<Number,
-                                                    ::dealii::MemorySpace::Device>
+      struct la_parallel_vector_templates_functions<
+        Number,
+        ::dealii::MemorySpace::Device>
       {
         using size_type = types::global_dof_index;
 
         static void
-        resize_val(const types::global_dof_index new_alloc_size,
-                   types::global_dof_index &     allocated_size,
-                   ::dealii::MemorySpace::
-                     MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data,
-                   const MPI_Comm &comm_sm)
+        resize_val(
+          const types::global_dof_index new_alloc_size,
+          types::global_dof_index &     allocated_size,
+          ::dealii::MemorySpace::MemorySpaceData<Number,
+                                                 ::dealii::MemorySpace::Device>
+            &             data,
+          const MPI_Comm &comm_sm)
         {
           (void)comm_sm;
 
           static_assert(
             std::is_same<Number, float>::value ||
               std::is_same<Number, double>::value,
-            "Number should be float or double for CUDA memory space");
+            "Number should be float or double for Device memory space");
 
           if (new_alloc_size > allocated_size)
             {
@@ -378,26 +382,35 @@ namespace LinearAlgebra
           // including ghost entries. this is not really efficient right now
           // because indices are translated twice, once by nth_index_in_set(i)
           // and once for operator() of tmp_vector
-          const IndexSet &       v_stored   = V.get_stored_elements();
-          const size_type        n_elements = v_stored.n_elements();
-         Kokkos::DefaultHostExecutionSpace host_exec;
-         Kokkos::View<size_type*, Kokkos::HostSpace> indices("indices", n_elements);
-         Kokkos::parallel_for(Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec, 0, n_elements),
-                         KOKKOS_LAMBDA(size_type i) { 
-            indices[i] = communication_pattern->global_to_local(
-              v_stored.nth_index_in_set(i));
-           });
-         host_exec.fence();
-         
+          const IndexSet &                  v_stored = V.get_stored_elements();
+          const size_type                   n_elements = v_stored.n_elements();
+          Kokkos::DefaultHostExecutionSpace host_exec;
+          Kokkos::View<size_type *, Kokkos::HostSpace> indices("indices",
+                                                               n_elements);
+          Kokkos::parallel_for(
+            Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec,
+                                                                   0,
+                                                                   n_elements),
+            KOKKOS_LAMBDA(size_type i) {
+              indices[i] = communication_pattern->global_to_local(
+                v_stored.nth_index_in_set(i));
+            });
+          host_exec.fence();
+
           // Move the indices to the device
           ::dealii::MemorySpace::Device::kokkos_space::execution_space exec;
-         Kokkos::View<size_type *, ::dealii::MemorySpace::Device::kokkos_space> indices_dev("indices_dev", n_elements);
-         Kokkos::deep_copy(exec, indices_dev, indices);
-         exec.fence();
-         
+          Kokkos::View<size_type *, ::dealii::MemorySpace::Device::kokkos_space>
+            indices_dev("indices_dev", n_elements);
+          Kokkos::deep_copy(exec, indices_dev, indices);
+          exec.fence();
+
           // Move the data to the device
-         Kokkos::View<Number *, ::dealii::MemorySpace::Device::kokkos_space> V_dev("V_dev", n_elements);
-         Kokkos::deep_copy(exec, V_dev, Kokkos::View<Number*, Kokkos::HostSpace>(V.begin(), n_elements));
+          Kokkos::View<Number *, ::dealii::MemorySpace::Device::kokkos_space>
+            V_dev("V_dev", n_elements);
+          Kokkos::deep_copy(
+            exec,
+            V_dev,
+            Kokkos::View<Number *, Kokkos::HostSpace>(V.begin(), n_elements));
           exec.fence();
 
           // Set the values in tmp_vector
@@ -415,17 +428,22 @@ namespace LinearAlgebra
           // Copy the local elements of tmp_vector to the right place in val
           IndexSet        tmp_index_set  = tmp_vector.locally_owned_elements();
           const size_type tmp_n_elements = tmp_index_set.n_elements();
-         Kokkos::realloc(indices, tmp_n_elements);
-          Kokkos::parallel_for(Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec, 0, n_elements),
-                          KOKKOS_LAMBDA(size_type i) {
-          indices[i] = locally_owned_elem.index_within_set(
-              tmp_index_set.nth_index_in_set(i));
+          Kokkos::realloc(indices, tmp_n_elements);
+          Kokkos::parallel_for(
+            Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec,
+                                                                   0,
+                                                                   n_elements),
+            KOKKOS_LAMBDA(size_type i) {
+              indices[i] = locally_owned_elem.index_within_set(
+                tmp_index_set.nth_index_in_set(i));
             });
           host_exec.fence();
-         Kokkos::realloc(indices_dev, tmp_n_elements);
-         Kokkos::deep_copy(indices_dev, Kokkos::View<size_type*>(indices.data(), tmp_n_elements));
+          Kokkos::realloc(indices_dev, tmp_n_elements);
+          Kokkos::deep_copy(indices_dev,
+                            Kokkos::View<size_type *>(indices.data(),
+                                                      tmp_n_elements));
 
-          if (operation == VectorOperation::add)
+          if (operation == VectorOperation::add)
             Kokkos::parallel_for(
               Kokkos::RangePolicy<
                 ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
@@ -449,13 +467,13 @@ namespace LinearAlgebra
         linfty_norm_local(const ::dealii::MemorySpace::MemorySpaceData<
                             Number,
                             ::dealii::MemorySpace::Device> &data,
-                          const unsigned int              size,
-                          RealType &                      result)
+                          const unsigned int                size,
+                          RealType &                        result)
         {
           static_assert(std::is_same<Number, RealType>::value,
                         "RealType should be the same type as Number");
 
-         typename ::dealii::MemorySpace::Device::kokkos_space::execution_space
+          typename ::dealii::MemorySpace::Device::kokkos_space::execution_space
             exec;
           Kokkos::parallel_reduce(
             Kokkos::RangePolicy<
@@ -929,7 +947,7 @@ namespace LinearAlgebra
       if (data.values_host_buffer.size() > 0)
         Kokkos::deep_copy(Kokkos::subview(data.values_host_buffer, range), 0);
       if (data.values.size() > 0)
-        Kokkos::deep_copy(Kokkos::subview(data.values, range), 0);         
+        Kokkos::deep_copy(Kokkos::subview(data.values, range), 0);
 
       vector_is_ghosted = false;
     }
@@ -1192,11 +1210,12 @@ namespace LinearAlgebra
 
               // The communication is done on the host, so we need to
               // move the data back to the device.
-             auto range = Kokkos::make_pair(
-                partitioner->locally_owned_size(),
-                partitioner->locally_owned_size() + partitioner->n_ghost_indices());
+              auto range = Kokkos::make_pair(partitioner->locally_owned_size(),
+                                             partitioner->locally_owned_size() +
+                                               partitioner->n_ghost_indices());
               Kokkos::deep_copy(Kokkos::subview(data.values, range),
-                            Kokkos::subview(data.values_host_buffer, range));
+                                Kokkos::subview(data.values_host_buffer,
+                                                range));
 
               Kokkos::resize(data.values_host_buffer, 0);
             }
@@ -1764,7 +1783,6 @@ namespace LinearAlgebra
     {
       real_type sum;
 
-
       dealii::internal::VectorOperations::
         functions<Number, Number, MemorySpaceType>::norm_2(
           thread_loop_partitioner,
index 21d08f4413eaa8790c16a07c52840a31121ea33d..452e378ce0084f18112f1ebaccb244f51a20ab48 100644 (file)
@@ -1484,9 +1484,9 @@ namespace internal
         ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace> & /*data*/)
       {
         static_assert(
-          std::is_same<MemorySpace, ::dealii::MemorySpace::CUDA>::value &&
+          std::is_same<MemorySpace, ::dealii::MemorySpace::Device>::value &&
             std::is_same<Number, Number2>::value,
-          "For the CUDA MemorySpace Number and Number2 should be the same type");
+          "For the Device MemorySpace Number and Number2 should be the same type");
       }
 
       static void
@@ -2097,7 +2097,6 @@ namespace internal
           }
       }
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
       template <typename MemorySpace2>
       static void
       import_elements(
@@ -2111,51 +2110,44 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data,
         std::enable_if_t<
-          std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
+          std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
           int> = 0)
       {
         if (operation == VectorOperation::insert)
           {
-            cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
-                                                     v_data.values.data(),
-                                                     size * sizeof(Number),
-                                                     cudaMemcpyDeviceToHost);
-            AssertCuda(cuda_error_code);
+            Kokkos::deep_copy(
+              Kokkos::subview(data.values,
+                              Kokkos::pair<size_type, size_type>(0, size)),
+              Kokkos::subview(v_data.values,
+                              Kokkos::pair<size_type, size_type>(0, size)));
           }
         else
           {
             AssertThrow(false, ExcNotImplemented());
           }
       }
-#endif
     };
 
 
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
     template <typename Number>
-    struct functions<Number, Number, ::dealii::MemorySpace::CUDA>
+    struct functions<Number, Number, ::dealii::MemorySpace::Device>
     {
-      static const int block_size =
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::block_size;
-      static const int chunk_size =
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::chunk_size;
-
       static void
       copy(
         const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
         const size_type size,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
-                                                 v_data.values.data(),
-                                                 size * sizeof(Number),
-                                                 cudaMemcpyDeviceToDevice);
-        AssertCuda(cuda_error_code);
+        Kokkos::deep_copy(
+          Kokkos::subview(data.values,
+                          Kokkos::pair<size_type, size_type>(0, size)),
+          Kokkos::subview(v_data.values,
+                          Kokkos::pair<size_type, size_type>(0, size)));
       }
 
       static void
@@ -2163,13 +2155,13 @@ namespace internal
           const size_type size,
           const Number    s,
           ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                 ::dealii::MemorySpace::CUDA>
+                                                 ::dealii::MemorySpace::Device>
             &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::set<Number>
-          <<<n_blocks, block_size>>>(data.values.data(), s, size);
-        AssertCudaKernel();
+        Kokkos::deep_copy(
+          Kokkos::subview(data.values,
+                          Kokkos::pair<size_type, size_type>(0, size)),
+          s);
       }
 
       static void
@@ -2177,18 +2169,20 @@ namespace internal
         const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
         const size_type size,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
-          <<<n_blocks, block_size>>>(data.values.data(),
-                                     1.,
-                                     v_data.values.data(),
-                                     size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(int i) { data.values(i) += v_data.values(i); });
+        exec.fence();
       }
 
       static void
@@ -2196,18 +2190,20 @@ namespace internal
         const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
         const size_type size,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
-          <<<n_blocks, block_size>>>(data.values.data(),
-                                     -1.,
-                                     v_data.values.data(),
-                                     size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) { data.values(i) -= v_data.values(i); });
+        exec.fence();
       }
 
       static void
@@ -2216,13 +2212,18 @@ namespace internal
         const size_type size,
         Number          a,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add<Number>
-          <<<n_blocks, block_size>>>(data.values.data(), a, size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) { data.values(i) += a; });
+        exec.fence();
       }
 
       static void
@@ -2231,18 +2232,22 @@ namespace internal
         const size_type size,
         const Number    a,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
-          <<<n_blocks, block_size>>>(data.values.data(),
-                                     a,
-                                     v_data.values.data(),
-                                     size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) += a * v_data.values(i);
+          });
+        exec.fence();
       }
 
       static void
@@ -2252,22 +2257,24 @@ namespace internal
         const Number    a,
         const Number    b,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aVbW<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
-                                                    a,
-                                                    v_data.values.data(),
-                                                    b,
-                                                    w_data.values.data(),
-                                                    size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) += a * v_data.values(i) + b * w_data.values(i);
+          });
+        exec.fence();
       }
 
       static void
@@ -2276,16 +2283,22 @@ namespace internal
         const size_type size,
         const Number    x,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(
-            x, data.values.data(), 1., v_data.values.data(), size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) = x * data.values(i) + v_data.values(i);
+          });
+        exec.fence();
       }
 
       static void
@@ -2295,16 +2308,22 @@ namespace internal
         const Number    x,
         const Number    a,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(
-            x, data.values.data(), a, v_data.values.data(), size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) = x * data.values(i) + a * v_data.values(i);
+          });
+        exec.fence();
       }
 
       static void
@@ -2315,23 +2334,25 @@ namespace internal
         const Number    a,
         const Number    b,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(x,
-                                                    data.values.data(),
-                                                    a,
-                                                    v_data.values.data(),
-                                                    b,
-                                                    w_data.values.data(),
-                                                    size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) =
+              x * data.values(i) + a * v_data.values(i) + b * w_data.values(i);
+          });
+        exec.fence();
       }
 
       static void
@@ -2340,13 +2361,18 @@ namespace internal
         const size_type size,
         const Number    factor,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale<Number>
-          <<<n_blocks, block_size>>>(data.values.data(), factor, size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) { data.values(i) *= factor; });
+        exec.fence();
       }
 
       static void
@@ -2354,17 +2380,20 @@ namespace internal
         const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
         const size_type size,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::scale<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
-                                                    v_data.values.data(),
-                                                    size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) { data.values(i) *= v_data.values(i); });
+        exec.fence();
       }
 
       static void
@@ -2373,18 +2402,22 @@ namespace internal
         const size_type size,
         const Number    a,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
-                                                    a,
-                                                    v_data.values.data(),
-                                                    size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) = a * v_data.values(i);
+          });
+        exec.fence();
       }
 
       static void
@@ -2394,63 +2427,50 @@ namespace internal
         const Number    a,
         const Number    b,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
-                                                    a,
-                                                    v_data.values.data(),
-                                                    b,
-                                                    w_data.values.data(),
-                                                    size);
-        AssertCudaKernel();
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_for(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i) {
+            data.values(i) = a * v_data.values(i) + b * w_data.values(i);
+          });
+        exec.fence();
       }
 
       static Number
       dot(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
           const size_type size,
           const ::dealii::MemorySpace::
-            MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+            MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
           ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                 ::dealii::MemorySpace::CUDA>
+                                                 ::dealii::MemorySpace::Device>
             &data)
       {
-        Number *    result_device;
-        cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
-        AssertCuda(error_code);
-        error_code = cudaMemset(result_device, 0, sizeof(Number));
-        AssertCuda(error_code);
-
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::double_vector_reduction<
-          Number,
-          ::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct<Number>>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
-                                                    data.values.data(),
-                                                    v_data.values.data(),
-                                                    static_cast<unsigned int>(
-                                                      size));
-        AssertCudaKernel();
-
-        // Copy the result back to the host
         Number result;
-        error_code = cudaMemcpy(&result,
-                                result_device,
-                                sizeof(Number),
-                                cudaMemcpyDeviceToHost);
-        AssertCuda(error_code);
-        // Free the memory on the device
-        error_code = cudaFree(result_device);
-        AssertCuda(error_code);
 
-        AssertIsFinite(result);
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_reduce(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i, Number & update) {
+            update += data.values(i) * v_data.values(i);
+          },
+          result);
 
+        AssertIsFinite(result);
         return result;
       }
 
@@ -2460,9 +2480,8 @@ namespace internal
                &             thread_loop_partitioner,
              const size_type size,
              real_type &     sum,
-             ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                    ::dealii::MemorySpace::CUDA>
-               &data)
+             ::dealii::MemorySpace::
+               MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
       {
         sum = dot(thread_loop_partitioner, size, data, data);
       }
@@ -2472,32 +2491,23 @@ namespace internal
         const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
         const size_type size,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
-      {
-        Number *    result_device;
-        cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
-        AssertCuda(error_code);
-        error_code = cudaMemset(result_device, 0, sizeof(Number));
-
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
-          Number,
-          ::dealii::LinearAlgebra::CUDAWrappers::kernel::ElemSum<Number>>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
-                                                    data.values.data(),
-                                                    size);
-
-        // Copy the result back to the host
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
+      {
         Number result;
-        error_code = cudaMemcpy(&result,
-                                result_device,
-                                sizeof(Number),
-                                cudaMemcpyDeviceToHost);
-        AssertCuda(error_code);
-        // Free the memory on the device
-        error_code = cudaFree(result_device);
-        AssertCuda(error_code);
 
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_reduce(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i, Number & update) {
+            update += data.values(i);
+          },
+          result);
+
+        AssertIsFinite(result);
         return result;
       }
 
@@ -2508,44 +2518,44 @@ namespace internal
         const size_type size,
         real_type &     sum,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        Number *    result_device;
-        cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
-        AssertCuda(error_code);
-        error_code = cudaMemset(result_device, 0, sizeof(Number));
-
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
-          Number,
-          ::dealii::LinearAlgebra::CUDAWrappers::kernel::L1Norm<Number>>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
-                                                    data.values.data(),
-                                                    size);
-
-        // Copy the result back to the host
-        error_code = cudaMemcpy(&sum,
-                                result_device,
-                                sizeof(Number),
-                                cudaMemcpyDeviceToHost);
-        AssertCuda(error_code);
-        // Free the memory on the device
-        error_code = cudaFree(result_device);
-        AssertCuda(error_code);
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_reduce(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i, Number & update) {
+            update += Kokkos::abs(data.values(i));
+          },
+          sum);
       }
 
       template <typename real_type>
       static void
       norm_p(
         const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
-        const size_type,
-        real_type &,
-        real_type,
+        const size_type size,
+        real_type &     sum,
+        real_type       exp,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA> &)
+                                               ::dealii::MemorySpace::Device>
+          &data)
       {
-        Assert(false, ExcNotImplemented());
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_reduce(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i, Number & update) {
+            update += Kokkos::pow(Kokkos::abs(data.values(i)), exp);
+          },
+          sum);
       }
 
       static Number
@@ -2554,33 +2564,29 @@ namespace internal
         const size_type size,
         const Number    a,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
         const ::dealii::MemorySpace::
-          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+          MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data)
       {
-        Number *    res_d;
-        cudaError_t error_code = cudaMalloc(&res_d, sizeof(Number));
-        AssertCuda(error_code);
-        error_code = cudaMemset(res_d, 0, sizeof(Number));
-        AssertCuda(error_code);
-
-        const int n_blocks = 1 + size / (chunk_size * block_size);
-        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_and_dot<Number>
-          <<<dim3(n_blocks, 1), dim3(block_size)>>>(res_d,
-                                                    data.values.data(),
-                                                    v_data.values.data(),
-                                                    w_data.values.data(),
-                                                    a,
-                                                    size);
-
         Number res;
-        error_code =
-          cudaMemcpy(&res, res_d, sizeof(Number), cudaMemcpyDeviceToHost);
-        AssertCuda(error_code);
-        error_code = cudaFree(res_d);
+
+        auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+          execution_space{};
+        Kokkos::parallel_reduce(
+          Kokkos::RangePolicy<
+            ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+                                                                          0,
+                                                                          size),
+          KOKKOS_LAMBDA(size_type i, Number & update) {
+            data.values(i) += a * v_data.values(i);
+            update +=
+              data.values(i) * Number(numbers::NumberTraits<Number>::conjugate(
+                                 w_data.values(i)));
+          },
+          res);
 
         return res;
       }
@@ -2595,10 +2601,10 @@ namespace internal
         const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
           &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data,
         std::enable_if_t<
-          std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
+          std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
           int> = 0)
       {
         if (operation == VectorOperation::insert)
@@ -2625,7 +2631,7 @@ namespace internal
         const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
           &v_data,
         ::dealii::MemorySpace::MemorySpaceData<Number,
-                                               ::dealii::MemorySpace::CUDA>
+                                               ::dealii::MemorySpace::Device>
           &data,
         std::enable_if_t<
           std::is_same<MemorySpace2, ::dealii::MemorySpace::Host>::value,
@@ -2633,11 +2639,11 @@ namespace internal
       {
         if (operation == VectorOperation::insert)
           {
-            cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
-                                                     v_data.values.data(),
-                                                     size * sizeof(Number),
-                                                     cudaMemcpyHostToDevice);
-            AssertCuda(cuda_error_code);
+            Kokkos::deep_copy(
+              Kokkos::subview(data.values,
+                              Kokkos::pair<size_type, size_type>(0, size)),
+              Kokkos::subview(v_data.values,
+                              Kokkos::pair<size_type, size_type>(0, size)));
           }
         else
           {
@@ -2645,7 +2651,6 @@ namespace internal
           }
       }
     };
-#endif
   } // namespace VectorOperations
 } // namespace internal
 
index b12138d84008a24055189fe0d48fd6713b8b9bb6..24f1ad0fffb7b3b5b2647d18b09c17b53386586d 100644 (file)
@@ -28,6 +28,7 @@ set(_unity_include_src
   scalapack.cc
   la_vector.cc
   la_parallel_vector.cc
+  la_parallel_vector_cuda.cc
   la_parallel_block_vector.cc
   matrix_out.cc
   precondition_block.cc
@@ -169,7 +170,6 @@ if(DEAL_II_WITH_CUDA)
     cuda_solver_direct.cc
     cuda_sparse_matrix.cc
     cuda_vector.cc
-    la_parallel_vector_cuda.cc
     read_write_vector_cuda.cc
   )
 endif()
index 98d71dae83a989b98ea96978360c70c45c70726c..db9ea9dfa2393e56a68d4381575e54d65674061b 100644 (file)
@@ -52,12 +52,12 @@ namespace internal
     release_all_unused_memory()
     {
 #include "vector_memory_release.inst"
-             dealii::GrowingVectorMemory<
-        dealii::LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>::
-        release_unused_memory();
-      dealii::GrowingVectorMemory<
-        dealii::LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>::
-        release_unused_memory();
+      dealii::GrowingVectorMemory<dealii::LinearAlgebra::distributed::Vector<
+        float,
+        MemorySpace::Device>>::release_unused_memory();
+      dealii::GrowingVectorMemory<dealii::LinearAlgebra::distributed::Vector<
+        double,
+        MemorySpace::Device>>::release_unused_memory();
 #ifdef DEAL_II_WITH_CUDA
       release_all_unused_cuda_memory();
 #endif

In the beginning the Universe was created. This has made a lot of people very angry and has been widely regarded as a bad move.

Douglas Adams


Typeset in Trocchi and Trocchi Bold Sans Serif.