From cfd7c36fbb4b4cbf456ab26669a9fd9198382813 Mon Sep 17 00:00:00 2001
From: Daniel Arndt <arndtd@ornl.gov>
Date: Wed, 7 Dec 2022 22:07:47 +0000
Subject: [PATCH] Fix tests

---
 tests/lac/parallel_vector_01.cc  |  2 +-
 tests/lac/parallel_vector_02.cc  |  2 +-
 tests/lac/parallel_vector_03.cc  |  2 +-
 tests/lac/parallel_vector_03a.cc |  2 +-
 tests/lac/parallel_vector_05.cc  |  2 +-
 tests/lac/parallel_vector_06.cc  |  6 +++---
 tests/lac/parallel_vector_08.cc  |  4 ++--
 tests/lac/parallel_vector_10.cc  |  4 ++--
 tests/lac/parallel_vector_11.cc  |  4 ++--
 tests/lac/parallel_vector_12.cc  |  6 +++---
 tests/lac/parallel_vector_13.cc  |  2 +-
 tests/lac/parallel_vector_14.cc  |  6 +++---
 tests/lac/parallel_vector_15.cc  |  4 ++--
 tests/lac/parallel_vector_16.cc  | 33 +++++++++-----------------------
 tests/lac/parallel_vector_21.cc  |  4 ++--
 tests/lac/parallel_vector_25.cc  |  6 +++---
 tests/lac/parallel_vector_26.cc  |  2 +-
 17 files changed, 38 insertions(+), 53 deletions(-)

diff --git a/tests/lac/parallel_vector_01.cc b/tests/lac/parallel_vector_01.cc
index 3a6b82053d..a0923b67af 100644
--- a/tests/lac/parallel_vector_01.cc
+++ b/tests/lac/parallel_vector_01.cc
@@ -44,7 +44,7 @@ test()
   IndexSet local_owned(numproc * 2);
   local_owned.add_range(myid * 2, myid * 2 + 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_owned, MPI_COMM_WORLD);
 
   // set local values
diff --git a/tests/lac/parallel_vector_02.cc b/tests/lac/parallel_vector_02.cc
index afd643c1eb..891f0ee834 100644
--- a/tests/lac/parallel_vector_02.cc
+++ b/tests/lac/parallel_vector_02.cc
@@ -47,7 +47,7 @@ test()
   local_relevant = local_owned;
   local_relevant.add_range(1, 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
 
   // set local values and check them
diff --git a/tests/lac/parallel_vector_03.cc b/tests/lac/parallel_vector_03.cc
index 5465361153..5d686847c8 100644
--- a/tests/lac/parallel_vector_03.cc
+++ b/tests/lac/parallel_vector_03.cc
@@ -47,7 +47,7 @@ test()
   local_relevant = local_owned;
   local_relevant.add_range(1, 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
 
   // set local values and check them
diff --git a/tests/lac/parallel_vector_03a.cc b/tests/lac/parallel_vector_03a.cc
index ed1ebfaa63..a7ff1fc723 100644
--- a/tests/lac/parallel_vector_03a.cc
+++ b/tests/lac/parallel_vector_03a.cc
@@ -80,7 +80,7 @@ test()
       local_relevant.add_range(4, 7);
     }
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
   v = 0.;
 
diff --git a/tests/lac/parallel_vector_05.cc b/tests/lac/parallel_vector_05.cc
index 7c42d01f40..9cf4631e2f 100644
--- a/tests/lac/parallel_vector_05.cc
+++ b/tests/lac/parallel_vector_05.cc
@@ -47,7 +47,7 @@ test()
   local_relevant = local_owned;
   local_relevant.add_range(1, 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
 
   // set local values and check them
diff --git a/tests/lac/parallel_vector_06.cc b/tests/lac/parallel_vector_06.cc
index 3ac6c1bab8..01b1432a56 100644
--- a/tests/lac/parallel_vector_06.cc
+++ b/tests/lac/parallel_vector_06.cc
@@ -48,7 +48,7 @@ test()
   local_relevant = local_owned;
   local_relevant.add_range(1, 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_owned, MPI_COMM_WORLD);
 
   // set local values
@@ -102,7 +102,7 @@ test()
   {
     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::CUDA> v2;
+    LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
     v2 = v;
     AssertThrow(std::fabs(v2 * v - norm_sqr) < 1e-15, ExcInternalError());
 
@@ -116,7 +116,7 @@ test()
     bool allzero = v.all_zero();
     if (myid == 0)
       deallog << " v==0 ? " << allzero << std::endl;
-    LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
+    LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
     v2.reinit(v);
     allzero = v2.all_zero();
     if (myid == 0)
diff --git a/tests/lac/parallel_vector_08.cc b/tests/lac/parallel_vector_08.cc
index 75da428ee5..c33c5dadf0 100644
--- a/tests/lac/parallel_vector_08.cc
+++ b/tests/lac/parallel_vector_08.cc
@@ -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::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(
     local_owned, local_owned, MPI_COMM_WORLD);
 
   // set a few of the local elements
diff --git a/tests/lac/parallel_vector_10.cc b/tests/lac/parallel_vector_10.cc
index 3ef59387db..0e13b1b21f 100644
--- a/tests/lac/parallel_vector_10.cc
+++ b/tests/lac/parallel_vector_10.cc
@@ -47,9 +47,9 @@ test()
   local_relevant = local_owned;
   local_relevant.add_range(1, 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v);
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v);
 
   // set local values and check them
   LinearAlgebra::ReadWriteVector<double> rw_vector(local_owned);
diff --git a/tests/lac/parallel_vector_11.cc b/tests/lac/parallel_vector_11.cc
index 592f1a0d42..2185db2980 100644
--- a/tests/lac/parallel_vector_11.cc
+++ b/tests/lac/parallel_vector_11.cc
@@ -55,10 +55,10 @@ test()
   local_relevant = local_owned;
   local_relevant.add_index(2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
   AssertDimension(static_cast<unsigned int>(actual_local_size), v.local_size());
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), x(v),
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v), x(v),
     y(v);
 
   // set local elements
diff --git a/tests/lac/parallel_vector_12.cc b/tests/lac/parallel_vector_12.cc
index 9236cc0520..0cdea76c30 100644
--- a/tests/lac/parallel_vector_12.cc
+++ b/tests/lac/parallel_vector_12.cc
@@ -61,7 +61,7 @@ test()
   if (numproc > 2)
     local_relevant0.add_index(8);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v0(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v0(
     local_owned0, local_relevant0, MPI_COMM_WORLD);
 
   // vector1: local size 4
@@ -83,7 +83,7 @@ test()
       local_relevant1.add_index(10);
     }
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v1(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v1(
     local_owned1, local_relevant1, MPI_COMM_WORLD);
 
   v0 = 1;
@@ -179,7 +179,7 @@ test()
     deallog << "Ghost values after re-set OK" << std::endl;
 
   // swap with an empty vector
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
   v2.swap(v0);
   AssertDimension(v0.size(), 0);
   AssertDimension(v2.size(), global_size1);
diff --git a/tests/lac/parallel_vector_13.cc b/tests/lac/parallel_vector_13.cc
index 5c4a9d2b64..eafc5552b5 100644
--- a/tests/lac/parallel_vector_13.cc
+++ b/tests/lac/parallel_vector_13.cc
@@ -49,7 +49,7 @@ test()
   local_relevant = local_owned;
   local_relevant.add_range(1, 2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
 
   // set local values and check them
diff --git a/tests/lac/parallel_vector_14.cc b/tests/lac/parallel_vector_14.cc
index c0e05a2901..fbb594bef1 100644
--- a/tests/lac/parallel_vector_14.cc
+++ b/tests/lac/parallel_vector_14.cc
@@ -55,7 +55,7 @@ test()
   // and once where they have not
   for (unsigned int run = 0; run < 2; ++run)
     {
-      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+      LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
         local_owned, local_relevant, MPI_COMM_WORLD);
 
       // set local values
@@ -68,7 +68,7 @@ test()
 
       v.import(rw_vector, VectorOperation::insert);
 
-      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), u(v);
+      LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v), u(v);
       u = 0;
 
       v *= 2.0;
@@ -90,7 +90,7 @@ test()
         }
 
       // copy vector content to non-ghosted vectors, manually created.
-      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v_dist(
+      LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v_dist(
         local_owned, MPI_COMM_WORLD),
         w_dist(v_dist), u_dist(v_dist);
 
diff --git a/tests/lac/parallel_vector_15.cc b/tests/lac/parallel_vector_15.cc
index ec64b62a28..2246d321a2 100644
--- a/tests/lac/parallel_vector_15.cc
+++ b/tests/lac/parallel_vector_15.cc
@@ -51,7 +51,7 @@ test()
   if (numproc > 1)
     local_relevant.add_range(3, 4);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
 
   // set local values
@@ -67,7 +67,7 @@ test()
   if (myid == 0)
     deallog << "v has ghost elements: " << v.has_ghost_elements() << std::endl;
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w, x;
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w, x;
   w = v;
   if (myid == 0)
     deallog << "w has ghost elements: " << w.has_ghost_elements() << std::endl;
diff --git a/tests/lac/parallel_vector_16.cc b/tests/lac/parallel_vector_16.cc
index 7ffc1ecbbc..b59379f967 100644
--- a/tests/lac/parallel_vector_16.cc
+++ b/tests/lac/parallel_vector_16.cc
@@ -30,12 +30,6 @@
 #include "../tests.h"
 
 
-__global__ void
-set_value(double *values_dev, unsigned int index, double val)
-{
-  values_dev[index] = val;
-}
-
 void
 test()
 {
@@ -58,17 +52,17 @@ 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::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
     local_owned, local_relevant, MPI_COMM_WORLD);
+  Kokkos::View<double*, MemorySpace::Device::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)
     {
-      double *values_dev = v.get_values();
-      set_value<<<1, 1>>>(values_dev, i, min_index + myid * local_size + i);
+	    Kokkos::deep_copy(Kokkos::subview(v_view, i), min_index + myid * local_size + i);
     }
 
   deallog << "vector norm: " << v.l2_norm() << std::endl;
@@ -79,20 +73,11 @@ test()
   v.print(deallog.get_file_stream(), 12, false, false);
 
   v.zero_out_ghost_values();
-  double *    values_dev  = v.get_values();
   const auto &partitioner = v.get_partitioner();
-  set_value<<<1, 1>>>(values_dev,
-                      partitioner->global_to_local(min_index + 38),
-                      min_index);
-  set_value<<<1, 1>>>(values_dev,
-                      partitioner->global_to_local(min_index + 39),
-                      min_index * 2);
-  set_value<<<1, 1>>>(values_dev,
-                      partitioner->global_to_local(min_index + 41),
-                      min_index + 7);
-  set_value<<<1, 1>>>(values_dev,
-                      partitioner->global_to_local(min_index + 42),
-                      -static_cast<double>(min_index));
+  Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 38)), min_index);
+  Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 39)), min_index*2);
+  Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 41)), min_index+7);
+  Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 42)), -static_cast<double>(min_index));
   v.compress(VectorOperation::add);
   v.update_ghost_values();
   v.print(deallog.get_file_stream(), 12, false, false);
diff --git a/tests/lac/parallel_vector_21.cc b/tests/lac/parallel_vector_21.cc
index a2eed7634c..3a1890530f 100644
--- a/tests/lac/parallel_vector_21.cc
+++ b/tests/lac/parallel_vector_21.cc
@@ -41,12 +41,12 @@ test()
   ghost_set.add_index(0);
   ghost_set.add_index(2);
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> 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::CUDA> v2;
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
   v2.reinit(v, true);
 
   // set locally owned range of v2 manually
diff --git a/tests/lac/parallel_vector_25.cc b/tests/lac/parallel_vector_25.cc
index d081dc558c..f4ae375c78 100644
--- a/tests/lac/parallel_vector_25.cc
+++ b/tests/lac/parallel_vector_25.cc
@@ -53,15 +53,15 @@ test()
   vec_ref.compress(VectorOperation::insert);
 
   auto partitioner = vec_ref.get_partitioner();
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> vec_dev(
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Device> vec_dev(
     partitioner);
   LinearAlgebra::distributed::Vector<double, MemorySpace::Host> vec_host(
     partitioner);
 
-  // Assignment from Host to CUDA
+  // Assignment from Host to Device
   vec_dev.import(vec_ref, VectorOperation::insert);
 
-  // Assignment from CUDA to Host
+  // Assignment from Device to Host
   vec_host.import(vec_dev, VectorOperation::insert);
 
   for (unsigned int i = 0; i < ghost_size; ++i)
diff --git a/tests/lac/parallel_vector_26.cc b/tests/lac/parallel_vector_26.cc
index aea7f2006e..6d6543f71a 100644
--- a/tests/lac/parallel_vector_26.cc
+++ b/tests/lac/parallel_vector_26.cc
@@ -15,7 +15,7 @@
 
 
 // Check for a bug where compress cannot be called for a LA::d::Vector<Host> in
-// a CUDA file
+// a Device file
 
 #include <deal.II/base/index_set.h>
 #include <deal.II/base/utilities.h>
-- 
2.39.5