]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Merge remote-tracking branch 'upstream/master' into kokkos_la_d_vector_kokkos
authorDaniel Arndt <arndtd@ornl.gov>
Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
committerDaniel Arndt <arndtd@ornl.gov>
Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
26 files changed:
1  2 
include/deal.II/base/memory_space_data.h
include/deal.II/lac/la_parallel_vector.templates.h
include/deal.II/lac/read_write_vector.templates.h
include/deal.II/lac/vector_operations_internal.h
source/lac/la_parallel_vector_cuda.cc
source/lac/read_write_vector.cc
source/lac/vector_memory.cc
tests/lac/parallel_vector_01.cc
tests/lac/parallel_vector_02.cc
tests/lac/parallel_vector_03.cc
tests/lac/parallel_vector_03a.cc
tests/lac/parallel_vector_04.cc
tests/lac/parallel_vector_05.cc
tests/lac/parallel_vector_06.cc
tests/lac/parallel_vector_08.cc
tests/lac/parallel_vector_10.cc
tests/lac/parallel_vector_11.cc
tests/lac/parallel_vector_12.cc
tests/lac/parallel_vector_13.cc
tests/lac/parallel_vector_14.cc
tests/lac/parallel_vector_15.cc
tests/lac/parallel_vector_16.cc
tests/lac/parallel_vector_21.cc
tests/lac/parallel_vector_22.cc
tests/lac/parallel_vector_23.cc
tests/lac/parallel_vector_25.cc

index 564d7afccbe57b1032a8dda9f62e536a09908da7,19c615a628710c2d21fe6922e0df86812cb8ce03..f30bbb32a1a9067679b68ff9c471035dfef5d4fb
@@@ -38,9 -38,10 +38,10 @@@ namespace MemorySpac
  {
    /**
     * Structure which stores data on the host or the device depending on the
-    * template parameter @p MemorySpace. The data is copied into the structure
-    * which then owns the data and will release the memory when the destructor is
-    * called.
+    * template parameter @p MemorySpace. Valid choices are MemorySpace::Host,
 -   * MemorySpace::Device, and MemorySpace::CUDA (if CUDA was enabled in
++   * MemorySpace::Default, and MemorySpace::CUDA (if CUDA was enabled in
+    * deal.II). The data is copied into the structure which then owns the data
+    * and will release the memory when the destructor is called.
     */
    template <typename T, typename MemorySpace>
    struct MemorySpaceData
index e310f207ec4810fae56b6c6d0b8f55934b27b084,21b101d18b6dbdb8390532ef5d5bb5e153d69a80..f0060403e5547e4f35ddc74b58458a686b4ed0cf
@@@ -86,10 -86,9 +86,10 @@@ namespace LinearAlgebr
        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::CUDA>::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");
++            std::is_same<MemorySpaceType, MemorySpace::Default>::value,
++          "MemorySpace should be Host or Default");
  
          static void
          resize_val(
          }
        };
  
 -#ifdef DEAL_II_COMPILER_CUDA_AWARE
        template <typename Number>
 -      struct la_parallel_vector_templates_functions<Number,
 -                                                    ::dealii::MemorySpace::CUDA>
 +      struct la_parallel_vector_templates_functions<
 +        Number,
-         ::dealii::MemorySpace::Device>
++        ::dealii::MemorySpace::Default>
        {
          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::CUDA> &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>
++                                                 ::dealii::MemorySpace::Default>
 +            &             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 Device memory space");
 -            "Number should be float or double for CUDA memory space");
++            "Number should be float or double for Default memory space");
  
            if (new_alloc_size > allocated_size)
              {
            const ReadWriteVector<Number> &V,
            VectorOperation::values        operation,
            std::shared_ptr<const Utilities::MPI::Partitioner>
 -                          communication_pattern,
 +            &             communication_pattern,
            const IndexSet &locally_owned_elem,
            ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
 -                                                 ::dealii::MemorySpace::CUDA>
++                                                 ::dealii::MemorySpace::Default>
              &data)
          {
            Assert(
                "Only VectorOperation::add and VectorOperation::insert are allowed"));
  
            ::dealii::LinearAlgebra::distributed::
-             Vector<Number, ::dealii::MemorySpace::Device>
 -            Vector<Number, ::dealii::MemorySpace::CUDA>
++            Vector<Number, ::dealii::MemorySpace::Default>
                tmp_vector(communication_pattern);
  
            // fill entries from ReadWriteVector into the distributed vector,
            // 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();
 -          std::vector<size_type> indices(n_elements);
 -          for (size_type i = 0; i < n_elements; ++i)
 -            indices[i] = communication_pattern->global_to_local(
 -              v_stored.nth_index_in_set(i));
 +          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>
 -          size_type *indices_dev;
 -          ::dealii::Utilities::CUDA::malloc(indices_dev, n_elements);
 -          ::dealii::Utilities::CUDA::copy_to_dev(indices, indices_dev);
++          ::dealii::MemorySpace::Default::kokkos_space::execution_space exec;
++          Kokkos::View<size_type *,
++                       ::dealii::MemorySpace::Default::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>
 -          Number *V_dev;
 -          ::dealii::Utilities::CUDA::malloc(V_dev, n_elements);
 -          cudaError_t cuda_error_code = cudaMemcpy(V_dev,
 -                                                   V.begin(),
 -                                                   n_elements * sizeof(Number),
 -                                                   cudaMemcpyHostToDevice);
 -          AssertCuda(cuda_error_code);
++          Kokkos::View<Number *, ::dealii::MemorySpace::Default::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
 -          const int n_blocks =
 -            1 + n_elements / (::dealii::CUDAWrappers::chunk_size *
 -                              ::dealii::CUDAWrappers::block_size);
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<Number>
 -            <<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
 -              indices_dev, tmp_vector.begin(), V_dev, n_elements);
 +          Kokkos::parallel_for(
 +            Kokkos::RangePolicy<
-               ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++              ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +              exec, 0, n_elements),
 +            KOKKOS_LAMBDA(size_type i) {
 +              tmp_vector(indices_dev(i)) = V_dev(i);
 +            });
 +          exec.fence();
  
            tmp_vector.compress(operation);
  
            // 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();
 -          indices.resize(tmp_n_elements);
 -          for (size_type i = 0; i < tmp_n_elements; ++i)
 -            indices[i] = locally_owned_elem.index_within_set(
 -              tmp_index_set.nth_index_in_set(i));
 -          ::dealii::Utilities::CUDA::free(indices_dev);
 -          ::dealii::Utilities::CUDA::malloc(indices_dev, tmp_n_elements);
 -          ::dealii::Utilities::CUDA::copy_to_dev(indices, indices_dev);
 +          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));
  
            if (operation == VectorOperation::add)
 -            ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated<
 -              Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
 -              indices_dev,
 -              data.values.data(),
 -              tmp_vector.begin(),
 -              tmp_n_elements);
 +            Kokkos::parallel_for(
 +              Kokkos::RangePolicy<
-                 ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++                ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +                exec, 0, n_elements),
 +              KOKKOS_LAMBDA(size_type i) {
 +                data.values(indices_dev(i)) += tmp_vector(i);
 +              });
            else
 -            ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<
 -              Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
 -              indices_dev,
 -              data.values.data(),
 -              tmp_vector.begin(),
 -              tmp_n_elements);
 -
 -          ::dealii::Utilities::CUDA::free(indices_dev);
 -          ::dealii::Utilities::CUDA::free(V_dev);
 +            Kokkos::parallel_for(
 +              Kokkos::RangePolicy<
-                 ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++                ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +                exec, 0, n_elements),
 +              KOKKOS_LAMBDA(size_type i) {
 +                data.values(indices_dev(i)) = tmp_vector(i);
 +              });
 +          exec.fence();
          }
  
          template <typename RealType>
          static void
          linfty_norm_local(const ::dealii::MemorySpace::MemorySpaceData<
                              Number,
-                             ::dealii::MemorySpace::Device> &data,
-                           const unsigned int                size,
-                           RealType &                        result)
 -                            ::dealii::MemorySpace::CUDA> &data,
 -                          const unsigned int              size,
 -                          RealType &                      result)
++                            ::dealii::MemorySpace::Default> &data,
++                          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
 -          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 / (::dealii::CUDAWrappers::chunk_size *
 -                                           ::dealii::CUDAWrappers::block_size);
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
 -            Number,
 -            ::dealii::LinearAlgebra::CUDAWrappers::kernel::LInfty<Number>>
 -            <<<dim3(n_blocks, 1), dim3(::dealii::CUDAWrappers::block_size)>>>(
 -              result_device, data.values.data(), size);
 -
 -          // Copy the result back to the host
 -          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);
++          typename ::dealii::MemorySpace::Default::kokkos_space::execution_space
 +            exec;
 +          Kokkos::parallel_reduce(
 +            Kokkos::RangePolicy<
-               ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++              ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +              exec, 0, size),
 +            KOKKOS_LAMBDA(size_type i, RealType & update) {
 +#if KOKKOS_VERSION < 30400
 +              update = fmax(update, fabs(data.values(i)));
 +#elif KOKKOS_VERSION < 30700
 +              update = Kokkos::Experimental::fmax(
 +                update, Kokkos::Experimental::fabs(data.values(i)));
 +#else
 +              update = Kokkos::fmax(update, Kokkos::abs(data.values(i)));
 +#endif
 +            },
 +            Kokkos::Max<RealType, Kokkos::HostSpace>(result));
          }
        };
 -#endif
      } // namespace internal
  
  
        // allocate import_data in case it is not set up yet
        if (partitioner->n_import_indices() > 0)
          {
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -          if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-           if (std::is_same<MemorySpaceType, dealii::MemorySpace::Device>::value)
++          if (std::is_same<MemorySpaceType,
++                           dealii::MemorySpace::Default>::value)
              {
                if (import_data.values_host_buffer.size() == 0)
                  Kokkos::resize(import_data.values_host_buffer,
              }
          }
  
 -#  if defined DEAL_II_COMPILER_CUDA_AWARE && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -      if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-       if (std::is_same<MemorySpaceType, dealii::MemorySpace::Device>::value)
++      if (std::is_same<MemorySpaceType, dealii::MemorySpace::Default>::value)
          {
            // Move the data to the host and then move it back to the
            // device. We use values to store the elements because the function
  
        // make this function thread safe
        std::lock_guard<std::mutex> lock(mutex);
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -      if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-       if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++      if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
          {
            Assert(partitioner->n_import_indices() == 0 ||
                     import_data.values_host_buffer.size() != 0,
                  partitioner->n_ghost_indices()),
                compress_requests);
          }
--
  #else
        (void)operation;
  #endif
        // allocate import_data in case it is not set up yet
        if (partitioner->n_import_indices() > 0)
          {
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -          if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-           if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++          if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
              {
                if (import_data.values_host_buffer.size() == 0)
                  Kokkos::resize(import_data.values_host_buffer,
              }
          }
  
 -#  if defined DEAL_II_COMPILER_CUDA_AWARE && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -      if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-       if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++      if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
          {
            // Move the data to the host and then move it back to the
            // device. We use values to store the elements because the function
            // make this function thread safe
            std::lock_guard<std::mutex> lock(mutex);
  
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -          if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-           if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++          if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
              {
                partitioner->export_to_ghosted_array_finish(
                  ArrayView<Number, MemorySpace::Host>(
index 08135c2d1033a47e14328c7a60b3e90085e9ca66,d1b9d6c070666daf6b5d31af1fe114ee7b066489..c9a8f3347075f6b0871d5309ef28585aec54327a
@@@ -109,8 -109,8 +109,8 @@@ namespace LinearAlgebr
  
          static_assert(
            std::is_same<MemorySpace, ::dealii::MemorySpace::Host>::value ||
-             std::is_same<MemorySpace, ::dealii::MemorySpace::Device>::value,
-           "MemorySpace should be Host or Device");
 -            std::is_same<MemorySpace, ::dealii::MemorySpace::CUDA>::value,
 -          "MemorySpace should be Host or CUDA");
++            std::is_same<MemorySpace, ::dealii::MemorySpace::Default>::value,
++          "MemorySpace should be Host or Default");
        }
      };
  
  
  
  
 -#ifdef DEAL_II_COMPILER_CUDA_AWARE
      template <typename Number>
-     struct read_write_vector_functions<Number, ::dealii::MemorySpace::Device>
 -    struct read_write_vector_functions<Number, ::dealii::MemorySpace::CUDA>
++    struct read_write_vector_functions<Number, ::dealii::MemorySpace::Default>
      {
        using size_type = types::global_dof_index;
  
  
          const unsigned int n_elements =
            communication_pattern->locally_owned_size();
 -        cudaError_t cuda_error_code = cudaMemcpy(tmp_vector.begin(),
 -                                                 values,
 -                                                 n_elements * sizeof(Number),
 -                                                 cudaMemcpyDeviceToHost);
 -        AssertCuda(cuda_error_code);
 +        Kokkos::deep_copy(
 +          Kokkos::View<Number *, Kokkos::HostSpace>(tmp_vector.begin(),
 +                                                    n_elements),
 +          Kokkos::View<const Number *,
-                        ::dealii::MemorySpace::Device::kokkos_space>(
++                       ::dealii::MemorySpace::Default::kokkos_space>(
 +            values, n_elements));
          tmp_vector.update_ghost_values();
  
          const IndexSet &stored = rw_vector.get_stored_elements();
index 4f2a2b475accbf1350e8daf7fac748432d0a02f7,21d08f4413eaa8790c16a07c52840a31121ea33d..92300d124e6a69438f580111b196f36f073319cf
@@@ -1484,9 -1484,9 +1484,9 @@@ namespace interna
          ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace> & /*data*/)
        {
          static_assert(
-           std::is_same<MemorySpace, ::dealii::MemorySpace::Device>::value &&
 -          std::is_same<MemorySpace, ::dealii::MemorySpace::CUDA>::value &&
++          std::is_same<MemorySpace, ::dealii::MemorySpace::Default>::value &&
              std::is_same<Number, Number2>::value,
-           "For the Device MemorySpace Number and Number2 should be the same type");
 -          "For the CUDA MemorySpace Number and Number2 should be the same type");
++          "For the Default MemorySpace Number and Number2 should be the same type");
        }
  
        static void
                                                 ::dealii::MemorySpace::Host>
            &data,
          std::enable_if_t<
-           std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
 -          std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
++          std::is_same<MemorySpace2, ::dealii::MemorySpace::Default>::value,
            int> = 0)
        {
          if (operation == VectorOperation::insert)
  
  
  
 -#ifdef DEAL_II_COMPILER_CUDA_AWARE
      template <typename Number>
-     struct functions<Number, Number, ::dealii::MemorySpace::Device>
 -    struct functions<Number, Number, ::dealii::MemorySpace::CUDA>
++    struct functions<Number, Number, ::dealii::MemorySpace::Default>
      {
 -      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::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &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
            const size_type size,
            const Number    s,
            ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
 -                                                 ::dealii::MemorySpace::CUDA>
++                                                 ::dealii::MemorySpace::Default>
              &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
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(int i) { data.values(i) += v_data.values(i); });
 +        exec.fence();
        }
  
        static void
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) -= v_data.values(i); });
 +        exec.fence();
        }
  
        static void
          const size_type size,
          Number          a,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) += a; });
 +        exec.fence();
        }
  
        static void
          const size_type size,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) += a * v_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
          const Number    a,
          const Number    b,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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
          const size_type size,
          const Number    x,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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
          const Number    x,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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
          const Number    a,
          const Number    b,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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
          const size_type size,
          const Number    factor,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) *= factor; });
 +        exec.fence();
        }
  
        static void
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) *= v_data.values(i); });
 +        exec.fence();
        }
  
        static void
          const size_type size,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) = a * v_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
          const Number    a,
          const Number    b,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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::Device> &v_data,
 -            MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++            MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
            ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
 -                                                 ::dealii::MemorySpace::CUDA>
++                                                 ::dealii::MemorySpace::Default>
              &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);
  
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        AssertIsFinite(result);
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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;
        }
  
                 &             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)
++               MemorySpaceData<Number, ::dealii::MemorySpace::Default> &data)
        {
          sum = dot(thread_loop_partitioner, size, data, data);
        }
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
 -          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::Default> &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::
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +            update += data.values(i);
 +          },
 +          result);
 +
 +        AssertIsFinite(result);
          return result;
        }
  
          const size_type size,
          real_type &     sum,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        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::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +#if KOKKOS_VERSION < 30400
 +            update += fabs(data.values(i));
 +#elif KOKKOS_VERSION < 30700
 +            update += Kokkos::Experimental::fabs(data.values(i));
 +#else
 +            update += Kokkos::fabs(data.values(i));
 +#endif
 +          },
 +          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::Device>
 -                                               ::dealii::MemorySpace::CUDA> &)
++                                               ::dealii::MemorySpace::Default>
 +          &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        Assert(false, ExcNotImplemented());
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +#if KOKKOS_VERSION < 30400
 +            update += pow(fabs(data.values(i)), exp);
 +#elif KOKKOS_VERSION < 30700
 +            update += Kokkos::Experimental::pow(
 +              Kokkos::Experimental::fabs(data.values(i)), exp);
 +#else
 +            update += Kokkos::pow(Kokkos::fabs(data.values(i)), exp);
 +#endif
 +          },
 +          sum);
        }
  
        static Number
          const size_type size,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &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::
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::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;
        }
          const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
            &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data,
          std::enable_if_t<
-           std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
 -          std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
++          std::is_same<MemorySpace2, ::dealii::MemorySpace::Default>::value,
            int> = 0)
        {
          if (operation == VectorOperation::insert)
          const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
            &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data,
          std::enable_if_t<
            std::is_same<MemorySpace2, ::dealii::MemorySpace::Host>::value,
index 76dc6c01ebb9dc9e8e2e9a88881b0f0e2f0d4e98,e4ebd6c6f51318305df4254ae38d6a9d1034adad..91fc69a62564ff1f224675623c254fe7b8e53e4a
@@@ -24,50 -24,40 +24,50 @@@ namespace LinearAlgebr
  {
    namespace distributed
    {
-     template class Vector<float, ::dealii::MemorySpace::Device>;
-     template class Vector<double, ::dealii::MemorySpace::Device>;
 -    template class Vector<float, ::dealii::MemorySpace::CUDA>;
 -    template class Vector<double, ::dealii::MemorySpace::CUDA>;
++    template class Vector<float, ::dealii::MemorySpace::Default>;
++    template class Vector<double, ::dealii::MemorySpace::Default>;
 +
      template void
      Vector<float, ::dealii::MemorySpace::Host>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<float, ::dealii::MemorySpace::CUDA> &,
++      ::dealii::MemorySpace::Default>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
      template void
      Vector<double, ::dealii::MemorySpace::Host>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<double, ::dealii::MemorySpace::CUDA> &,
++      ::dealii::MemorySpace::Default>(
++      const Vector<double, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
  
      template void
-     Vector<float, ::dealii::MemorySpace::Device>::import<
 -    Vector<float, ::dealii::MemorySpace::CUDA>::import<
++    Vector<float, ::dealii::MemorySpace::Default>::import<
        ::dealii::MemorySpace::Host>(
        const Vector<float, ::dealii::MemorySpace::Host> &,
        VectorOperation::values);
      template void
-     Vector<double, ::dealii::MemorySpace::Device>::import<
 -    Vector<double, ::dealii::MemorySpace::CUDA>::import<
++    Vector<double, ::dealii::MemorySpace::Default>::import<
        ::dealii::MemorySpace::Host>(
        const Vector<double, ::dealii::MemorySpace::Host> &,
        VectorOperation::values);
  
      template void
-     Vector<float, ::dealii::MemorySpace::Device>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
 -    Vector<float, ::dealii::MemorySpace::CUDA>::import<
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<float, ::dealii::MemorySpace::CUDA> &,
++    Vector<float, ::dealii::MemorySpace::Default>::import<
++      ::dealii::MemorySpace::Default>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
      template void
-     Vector<double, ::dealii::MemorySpace::Device>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
 -    Vector<double, ::dealii::MemorySpace::CUDA>::import<
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<double, ::dealii::MemorySpace::CUDA> &,
++    Vector<double, ::dealii::MemorySpace::Default>::import<
++      ::dealii::MemorySpace::Default>(
++      const Vector<double, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
-     Vector<float, ::dealii::MemorySpace::Device>::reinit<float>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
 +
 +    template void
-     Vector<double, ::dealii::MemorySpace::Device>::reinit<double>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
++    Vector<float, ::dealii::MemorySpace::Default>::reinit<float>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
 +      const bool);
 +    template void
++    Vector<double, ::dealii::MemorySpace::Default>::reinit<double>(
++      const Vector<double, ::dealii::MemorySpace::Default> &,
 +      const bool);
    } // namespace distributed
  } // namespace LinearAlgebra
  
index a6dc6febd2d640a7d22797542ae0fc66167cea56,f1ca4fee7ce39a71acb04e44e265a0c18c9ecf93..23ced9c5bf7cfd126c0d147f822e08adb65505a7
@@@ -46,22 -46,12 +46,22 @@@ namespace LinearAlgebr
      const distributed::Vector<float, ::dealii::MemorySpace::Host> &,
      VectorOperation::values,
      const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
-     const distributed::Vector<float, ::dealii::MemorySpace::Device> &,
 +  template void
 +  ReadWriteVector<float>::import(
++    const distributed::Vector<float, ::dealii::MemorySpace::Default> &,
 +    VectorOperation::values,
 +    const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
  
    template void
    ReadWriteVector<double>::import(
      const distributed::Vector<double, ::dealii::MemorySpace::Host> &,
      VectorOperation::values,
      const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
-     const distributed::Vector<double, ::dealii::MemorySpace::Device> &,
 +  template void
 +  ReadWriteVector<double>::import(
++    const distributed::Vector<double, ::dealii::MemorySpace::Default> &,
 +    VectorOperation::values,
 +    const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
  #  ifdef DEAL_II_WITH_COMPLEX_VALUES
    template void
    ReadWriteVector<std::complex<float>>::import(
index db9ea9dfa2393e56a68d4381575e54d65674061b,461c1cf50c28440efbc24a35a712f2d5bede00bd..5ab7a4c8d683f1628067929248c7d7d745d90df6
  DEAL_II_NAMESPACE_OPEN
  
  #include "vector_memory.inst"
-   LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>;
 +template class VectorMemory<
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<float, MemorySpace::Default>>;
 +template class VectorMemory<
-   LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>;
 +template class GrowingVectorMemory<
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<float, MemorySpace::Default>>;
 +template class GrowingVectorMemory<
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>;
  
  namespace internal
  {
      release_all_unused_memory()
      {
  #include "vector_memory_release.inst"
-         MemorySpace::Device>>::release_unused_memory();
 +      dealii::GrowingVectorMemory<dealii::LinearAlgebra::distributed::Vector<
 +        float,
-         MemorySpace::Device>>::release_unused_memory();
++        MemorySpace::Default>>::release_unused_memory();
 +      dealii::GrowingVectorMemory<dealii::LinearAlgebra::distributed::Vector<
 +        double,
++        MemorySpace::Default>>::release_unused_memory();
  #ifdef DEAL_II_WITH_CUDA
        release_all_unused_cuda_memory();
  #endif
index 60dfbc9f2ba23d16a6296710c5b05533790eba07,5021d1aa16d082e8255c60e1cd5a1903632c403f..cad97db078eefb049544df65efc50f5d1903b8cc
@@@ -44,7 -44,7 +44,7 @@@ test(
    IndexSet local_owned(numproc * 2);
    local_owned.add_range(myid * 2, myid * 2 + 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_owned, MPI_COMM_WORLD);
  
    // set local values
index 891f0ee834464afcb866396ef7bd685d252a6d6b,e7664f4e1c158273ffc55d4dc84686e57a854fcf..79e9922ceed1b29cd3658c67bf1013c4a785f0e0
@@@ -47,7 -47,7 +47,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
index 5d686847c82b2a8850f4bc164927d2e5f53cc569,0e34c7ed54694455ca6dd1bc49f5300b8ffeb928..33145afb4ebffa541adeeffb017a01f51882e36f
@@@ -47,7 -47,7 +47,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
index a7ff1fc723e86edceebf3651ced811b1db33bb7d,47f510ab7e9ff903e856c6c6b14d772190a5aaae..190dc37fc1f9854a9d34c2e6d25f233280a171e5
@@@ -80,7 -80,7 +80,7 @@@ test(
        local_relevant.add_range(4, 7);
      }
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
    v = 0.;
  
index a3cb0938547cebf80dc2de56fc4bc92b0875c3e2,e928ed32b176e04ce59ba699ff288d43ea77bc7e..08f2e81776f27ed4bb4cdcf200c413a4782fd013
@@@ -48,7 -54,7 +48,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
      {
        unsigned int local_index = partitioner->global_to_local(1);
        double *     values_dev  = v.get_values();
-       Kokkos::deep_copy(Kokkos::View<double, MemorySpace::Device::kokkos_space>(
-                           values_dev + local_index),
-                         7);
 -      set_value<<<1, 1>>>(values_dev, local_index, 7);
++      Kokkos::deep_copy(
++        Kokkos::View<double, MemorySpace::Default::kokkos_space>(values_dev +
++                                                                 local_index),
++        7);
      }
  
 -  unsigned int        allocated_size = local_relevant.n_elements();
 -  std::vector<double> v_host(allocated_size);
 -  Utilities::CUDA::copy_to_host(v.get_values(), v_host);
 +  unsigned int allocated_size = local_relevant.n_elements();
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v_device(
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v_device(
 +    v.get_values(), allocated_size);
 +  Kokkos::View<double *, Kokkos::HostSpace> v_host("v_host", allocated_size);
 +  Kokkos::deep_copy(v_host, v_device);
  
    AssertThrow(v_host[partitioner->global_to_local(myid * 2)] == myid * 4.0,
                ExcInternalError());
index 9cf4631e2f1944025de6a85f6600bbd2d89fc09e,1519bd97e6934fe48c0cd5a7031c88858d1854ac..005e4e1c17b0b0e04946b5c4d90bfd630140182e
@@@ -47,7 -47,7 +47,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
index 01b1432a569719aad33ff65a022ed62fd2ca364d,bfe06fffaeff66fe648fadde5dd1cd007411da0d..804248706cc63dacda3543c75febd6572576822f
@@@ -48,7 -48,7 +48,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_owned, MPI_COMM_WORLD);
  
    // set local values
    {
      const double norm_sqr = v.l2_norm() * v.l2_norm();
      AssertThrow(std::fabs(v * v - norm_sqr) < 1e-15, ExcInternalError());
-     LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -    LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++    LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
      v2 = v;
      AssertThrow(std::fabs(v2 * v - norm_sqr) < 1e-15, ExcInternalError());
  
      bool allzero = v.all_zero();
      if (myid == 0)
        deallog << " v==0 ? " << allzero << std::endl;
-     LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -    LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++    LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
      v2.reinit(v);
      allzero = v2.all_zero();
      if (myid == 0)
index c33c5dadf0625b01d2d1b08484a8b81f5caa8c4c,f0c4716232a51f064edf2acc1f9437e95ed7e7b6..b9c4b6cd97d5ef537807daab05d54627da99e002
@@@ -66,9 -66,9 +66,9 @@@ test(
    // v has ghosts, w has none. set some entries
    // on w, copy into v and check if they are
    // there
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(
      local_owned, local_owned, MPI_COMM_WORLD);
  
    // set a few of the local elements
index 0e13b1b21f8045946bfbe3cfd2d366f23da59298,6b250c3292583de315dd4c7b0a7e9b377a23f60f..715a646fbd19d3c5c42f4cdfe59f3309f3b103aa
@@@ -47,9 -47,9 +47,9 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v);
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v);
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(v);
  
    // set local values and check them
    LinearAlgebra::ReadWriteVector<double> rw_vector(local_owned);
index 2185db298096f9a910df24555c0fb50558bcfc4b,a46a73426e26b5ce9090e66b957e0d108f841382..df0285c862ca1c60dca7e3babd23daaee65d9235
@@@ -55,10 -55,10 +55,10 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_index(2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
    AssertDimension(static_cast<unsigned int>(actual_local_size), v.local_size());
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v), x(v),
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), x(v),
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(v), x(v),
      y(v);
  
    // set local elements
index 0cdea76c30a9b1f61101eb549a02f3f7ce2d5071,9103d9ef77dcd1566b902328ef2990ef0a768f2d..97f83361f80fb5a5c73bdc7054e95b0e794b417c
@@@ -61,7 -61,7 +61,7 @@@ test(
    if (numproc > 2)
      local_relevant0.add_index(8);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v0(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v0(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v0(
      local_owned0, local_relevant0, MPI_COMM_WORLD);
  
    // vector1: local size 4
@@@ -83,7 -83,7 +83,7 @@@
        local_relevant1.add_index(10);
      }
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v1(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v1(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v1(
      local_owned1, local_relevant1, MPI_COMM_WORLD);
  
    v0 = 1;
      deallog << "Ghost values after re-set OK" << std::endl;
  
    // swap with an empty vector
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
    v2.swap(v0);
    AssertDimension(v0.size(), 0);
    AssertDimension(v2.size(), global_size1);
index eafc5552b5dbcf567d9ace4d17dcea3d297ab7f5,c157244982da4a3c4c47cc23a8ebd893243a2fa4..5249faa317161bec491f429cabf6808c35d7d2e6
@@@ -49,7 -49,7 +49,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
index 4ff6f1bbeca5763974121611a2ff50f74b049bf9,c3fab28bade4d73c1f4aaa827e022c14c18434ea..2d568e8cc15852df7cae636908d6df92d293d5e9
@@@ -55,7 -55,7 +55,7 @@@ test(
    // and once where they have not
    for (unsigned int run = 0; run < 2; ++run)
      {
-       LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++      LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
          local_owned, local_relevant, MPI_COMM_WORLD);
  
        // set local values
@@@ -68,8 -68,7 +68,8 @@@
  
        v.import(rw_vector, VectorOperation::insert);
  
-       LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v),
 -      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), u(v);
++      LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(v),
 +        u(v);
        u = 0;
  
        v *= 2.0;
@@@ -91,7 -90,7 +91,7 @@@
          }
  
        // copy vector content to non-ghosted vectors, manually created.
-       LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v_dist(
 -      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v_dist(
++      LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v_dist(
          local_owned, MPI_COMM_WORLD),
          w_dist(v_dist), u_dist(v_dist);
  
index 2246d321a23a03a57f8c7de0b06b914fabeb2dd2,050c320fedf489e4390c06341505db575470b5c8..4233a9d88fee5a6d879a7ccd15c6229cd110aa23
@@@ -51,7 -51,7 +51,7 @@@ test(
    if (numproc > 1)
      local_relevant.add_range(3, 4);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values
@@@ -67,7 -67,7 +67,7 @@@
    if (myid == 0)
      deallog << "v has ghost elements: " << v.has_ghost_elements() << std::endl;
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w, x;
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w, x;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w, x;
    w = v;
    if (myid == 0)
      deallog << "w has ghost elements: " << w.has_ghost_elements() << std::endl;
index 21e0f5cc2358566024abb528da9d51ab88095a90,68e71da7662fcca71d70be2062a81094c9d4dcd7..23670af24175933187ad5ba87ce7463756c66bc1
@@@ -52,14 -58,11 +52,14 @@@ test(
    local_relevant.add_range(min_index + 38, min_index + 40);
    local_relevant.add_range(min_index + 41, min_index + 43);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v_view(
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v_view(
 +    v.get_values(), local_relevant.n_elements());
  
 -  deallog << "Local range of proc 0: " << v.local_range().first << " "
 -          << v.local_range().second << std::endl;
 +  deallog << "Local range of proc 0: "
 +          << v.get_partitioner()->local_range().first << " "
 +          << v.get_partitioner()->local_range().second << std::endl;
  
    // set local values
    for (types::global_dof_index i = 0; i < local_size; ++i)
index ac37704758e27c6755b01e6ff5d8893592820c4b,99f867327d2b3cb838488d4760c3ebb0a4a99361..ddedd449b9038eaafe16bb383aff993ed5970e32
@@@ -41,23 -47,17 +41,23 @@@ test(
    ghost_set.add_index(0);
    ghost_set.add_index(2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      locally_owned, ghost_set, MPI_COMM_WORLD);
  
    // create vector without actually setting the entries since they will be
    // overwritten soon anyway
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
    v2.reinit(v, true);
  
    // set locally owned range of v2 manually
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v2_view(
 -  set<<<1, v2.local_size()>>>(v2.get_values());
 -
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v2_view(
 +    v2.get_values(), v2.local_size());
 +  Kokkos::deep_copy(v2_view, 1.);
 +
 +  Kokkos::parallel_for(
 +    v2.local_size(), KOKKOS_LAMBDA(int i) {
 +      KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d: %f\n", i, v2_view(i));
 +    });
    // add entries to ghost values
    // Because of limitation in import, the IndexSet of the ReadWriteVector needs
    // to have the local elements.
index 1697623d1118672a0d3967ebebba4a4413748f71,6456af6e9b6fbf96175d13701bd3a2c16d0ec4e4..cfde38daf6b4621555fe6de7bafa9b36f0d1b392
@@@ -57,10 -64,8 +57,10 @@@ test(
    local_relevant.add_range(1, 2);
  
    // create vector
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v_device(
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v_device(
 +    v.get_values(), local_relevant.n_elements());
    const auto &partitioner = v.get_partitioner();
  
    // set local values
index 16bcfa90e6a4895c816976478309c401be9deca4,a9bb9c5d9ce18bff86554a2dc23f900bd1233da2..064e50e71256d520196408fca644995eb5ccbc5d
@@@ -58,7 -65,7 +58,7 @@@ test(
    local_relevant.add_range(1, 2);
  
    // create vector
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
    const auto &partitioner = v.get_partitioner();
  
index f4ae375c78960cb98543e306fa70edd84455f677,88f3ca0f5eaa4225550be4d58dbe5e10049e48a2..91858b67e81b0d07f4a77486a1127de88628fd2e
@@@ -53,15 -53,15 +53,15 @@@ test(
    vec_ref.compress(VectorOperation::insert);
  
    auto partitioner = vec_ref.get_partitioner();
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> vec_dev(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> vec_dev(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> vec_dev(
      partitioner);
    LinearAlgebra::distributed::Vector<double, MemorySpace::Host> vec_host(
      partitioner);
  
-   // Assignment from Host to Device
 -  // Assignment from Host to CUDA
++  // Assignment from Host to Default
    vec_dev.import(vec_ref, VectorOperation::insert);
  
-   // Assignment from Device to Host
 -  // Assignment from CUDA to Host
++  // Assignment from Default to Host
    vec_host.import(vec_dev, VectorOperation::insert);
  
    for (unsigned int i = 0; i < ghost_size; ++i)

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.