]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Move CUDA tests
authorDaniel Arndt <arndtd@ornl.gov>
Fri, 30 Dec 2022 17:00:15 +0000 (18:00 +0100)
committerDaniel Arndt <arndtd@ornl.gov>
Sat, 31 Dec 2022 17:40:59 +0000 (18:40 +0100)
20 files changed:
tests/base/array_view_access_data.cc [moved from tests/cuda/array_view_access_data.cc with 67% similarity]
tests/base/array_view_access_data.debug.output [moved from tests/cuda/array_view_access_data.debug.output with 100% similarity]
tests/base/kokkos_point.cc [moved from tests/cuda/cuda_point.cc with 53% similarity]
tests/base/kokkos_point.output [moved from tests/cuda/cuda_point.output with 100% similarity]
tests/base/kokkos_tensor_01.cc [new file with mode: 0644]
tests/base/kokkos_tensor_01.output [moved from tests/cuda/cuda_tensor_01.output with 100% similarity]
tests/base/kokkos_tensor_02.cc [new file with mode: 0644]
tests/base/kokkos_tensor_02.output [moved from tests/cuda/cuda_tensor_02.output with 100% similarity]
tests/cuda/array_view_wrong_memory.cc [deleted file]
tests/cuda/array_view_wrong_memory.debug.output [deleted file]
tests/cuda/cuda_tensor_01.cc [deleted file]
tests/cuda/cuda_tensor_02.cc [deleted file]
tests/lac/affine_constraints_set_zero.cc [moved from tests/cuda/affine_constraints_set_zero.cc with 74% similarity]
tests/lac/affine_constraints_set_zero.mpirun=2.output [moved from tests/cuda/affine_constraints_set_zero.mpirun=2.output with 100% similarity]
tests/lac/vector_memory_01.cc [moved from tests/cuda/vector_memory_01.cc with 100% similarity]
tests/lac/vector_memory_01.output [moved from tests/cuda/vector_memory_01.output with 100% similarity]
tests/lac/vector_memory_02.cc [moved from tests/cuda/vector_memory_02.cc with 100% similarity]
tests/lac/vector_memory_02.debug.output [moved from tests/cuda/vector_memory_02.debug.output with 100% similarity]
tests/lac/vector_reinit_04.cc [moved from tests/cuda/vector_reinit_01.cc with 99% similarity]
tests/lac/vector_reinit_04.output [moved from tests/cuda/vector_reinit_01.output with 100% similarity]

similarity index 67%
rename from tests/cuda/array_view_access_data.cc
rename to tests/base/array_view_access_data.cc
index d8f708db8352da42bf7fa05d4062acae6116edce..ea48eba17533c32aa36619fa15d603e6177c2dd4 100644 (file)
@@ -14,8 +14,8 @@
 // ---------------------------------------------------------------------
 
 
-// check that we detect that accessing CUDA memory in an ArrayView object
-// is not allowed.
+// check that we detect that accessing memory in MemorySpace::Default using an
+// ArrayView object is not allowed.
 
 #include <deal.II/base/array_view.h>
 
@@ -28,16 +28,14 @@ main(int argc, char **argv)
 
   initlog();
 
-  init_cuda();
-
-  std::unique_ptr<unsigned int[], void (*)(unsigned int *)> dummy_cuda(
-    Utilities::CUDA::allocate_device_data<unsigned int>(2),
-    Utilities::CUDA::delete_device_data<unsigned int>);
+  Kokkos::ScopeGuard                                               guard;
+  Kokkos::View<unsigned int *, MemorySpace::Default::kokkos_space> dummy(
+    "dummy", 2);
 
   try
     {
-      ArrayView<unsigned int, MemorySpace::CUDA> view(dummy_cuda.get(), 2);
-      const auto                                 dummy = view[0];
+      ArrayView<unsigned int, MemorySpace::Default> view(dummy.data(), 2);
+      const auto                                    dummy = view[0];
     }
   catch (const ExceptionBase &exc)
     {
similarity index 53%
rename from tests/cuda/cuda_point.cc
rename to tests/base/kokkos_point.cc
index 871cf723182fe149a5caa6772bac8d8e9b7138c0..093069fc00c91fe725eedd0f75f5145fcdd0cc74 100644 (file)
 
 #include "../tests.h"
 
-template <int dim, typename Number>
-__global__ void
-miscellaneous_kernel(Number check[16])
-{
-  Point<dim, Number> p_1;
-  check[0] = p_1.norm_square();
-  Point<dim, Number> p_2(Tensor<1, dim, Number>{});
-  check[1] = p_2.norm_square();
-  if (dim == 1)
-    {
-      Point<dim, Number> p(1.);
-      check[2] = p.norm_square();
-    }
-  if (dim == 2)
-    {
-      Point<dim, Number> p(.6, .8);
-      check[2] = p.norm_square();
-    }
-  if (dim == 3)
-    {
-      Point<dim, Number> p(.48, .64, .6);
-      check[2] = p.norm_square();
-    }
-
-  auto p_3 = Point<dim, Number>::unit_vector(0);
-  check[3] = p_3.norm_square();
-
-  auto entry_1 = p_1(0);
-  check[4]     = entry_1;
-  p_1(0)       = Number{1.};
-  check[5]     = p_1.norm_square();
-  auto p_4     = p_1 + Tensor<1, dim, Number>{};
-  check[6]     = p_4.norm_square();
-  auto p_5     = p_1 - Tensor<1, dim, Number>{};
-  check[7]     = p_5.norm_square();
-  auto t_1     = p_1 - p_2;
-  check[8]     = t_1.norm_square();
-  auto p_6     = -p_3;
-  check[9]     = p_6.norm_square();
-  auto p_7     = p_4 / 2.;
-  check[10]    = p_7.norm_square();
-  auto p_8     = p_7 * 5.;
-  check[11]    = p_8.norm_square();
-
-  auto s_1  = p_1 * t_1;
-  check[12] = s_1;
-  auto s_2  = p_2.square();
-  check[13] = s_2;
-  auto s_3  = p_3.distance(p_5);
-  check[14] = s_3;
-  auto s_4  = p_4.distance_square(p_1);
-  check[15] = s_4;
-}
-
 template <int dim, typename Number>
 void
 test_gpu()
 {
-  Number *           check;
   const unsigned int n_tests = 16;
 
-  auto cuda_error = cudaMalloc(&check, n_tests * sizeof(Number));
-  AssertCuda(cuda_error);
+  Kokkos::View<Number *, MemorySpace::Default::kokkos_space> check("check",
+                                                                   n_tests);
 
   // Miscellaneous
-  miscellaneous_kernel<dim, Number><<<1, 1>>>(check);
-  // Check that the kernel was launched correctly
-  AssertCuda(cudaPeekAtLastError());
-  // Check that there was no problem during the execution of the kernel
-  AssertCuda(cudaDeviceSynchronize());
-
-  std::vector<Number> check_host(n_tests);
-
-  cuda_error = cudaMemcpy(check_host.data(),
-                          check,
-                          n_tests * sizeof(Number),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
+  using ExecutionSpace = MemorySpace::Default::kokkos_space::execution_space;
+  ExecutionSpace exec;
+  Kokkos::parallel_for(
+    Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+      Point<dim, Number> p_1;
+      check[0] = p_1.norm_square();
+      Point<dim, Number> p_2(Tensor<1, dim, Number>{});
+      check[1] = p_2.norm_square();
+      if (dim == 1)
+        {
+          Point<dim, Number> p(1.);
+          check[2] = p.norm_square();
+        }
+      if (dim == 2)
+        {
+          Point<dim, Number> p(.6, .8);
+          check[2] = p.norm_square();
+        }
+      if (dim == 3)
+        {
+          Point<dim, Number> p(.48, .64, .6);
+          check[2] = p.norm_square();
+        }
+
+      auto p_3 = Point<dim, Number>::unit_vector(0);
+      check[3] = p_3.norm_square();
+
+      auto entry_1 = p_1(0);
+      check[4]     = entry_1;
+      p_1(0)       = Number{1.};
+      check[5]     = p_1.norm_square();
+      auto p_4     = p_1 + Tensor<1, dim, Number>{};
+      check[6]     = p_4.norm_square();
+      auto p_5     = p_1 - Tensor<1, dim, Number>{};
+      check[7]     = p_5.norm_square();
+      auto t_1     = p_1 - p_2;
+      check[8]     = t_1.norm_square();
+      auto p_6     = -p_3;
+      check[9]     = p_6.norm_square();
+      auto p_7     = p_4 / 2.;
+      check[10]    = p_7.norm_square();
+      auto p_8     = p_7 * 5.;
+      check[11]    = p_8.norm_square();
+
+      auto s_1  = p_1 * t_1;
+      check[12] = s_1;
+      auto s_2  = p_2.square();
+      check[13] = s_2;
+      auto s_3  = p_3.distance(p_5);
+      check[14] = s_3;
+      auto s_4  = p_4.distance_square(p_1);
+      check[15] = s_4;
+    });
+
+  exec.fence();
+
+  auto check_host =
+    Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, check);
 
   const double tolerance = 1.e-8;
   AssertThrow(std::abs(check_host[0] - 0.) < tolerance, ExcInternalError());
@@ -116,9 +106,6 @@ test_gpu()
   AssertThrow(std::abs(check_host[14] - 0.) < tolerance, ExcInternalError());
   AssertThrow(std::abs(check_host[15] - 0.) < tolerance, ExcInternalError());
 
-  cuda_error = cudaFree(check);
-  AssertCuda(cuda_error);
-
   deallog << "OK" << std::endl;
 }
 
@@ -127,7 +114,7 @@ main()
 {
   initlog();
 
-  init_cuda();
+  Kokkos::initialize();
 
   test_gpu<1, double>();
   test_gpu<2, double>();
@@ -135,4 +122,6 @@ main()
   test_gpu<1, float>();
   test_gpu<2, float>();
   test_gpu<3, float>();
+
+  Kokkos::finalize();
 }
diff --git a/tests/base/kokkos_tensor_01.cc b/tests/base/kokkos_tensor_01.cc
new file mode 100644 (file)
index 0000000..db13557
--- /dev/null
@@ -0,0 +1,91 @@
+// ---------------------------------------------------------------------
+//
+// Copyright (C) 2016 by the deal.II authors
+//
+// This file is part of the deal.II library.
+//
+// The deal.II library is free software; you can use it, redistribute
+// it, and/or modify it under the terms of the GNU Lesser General
+// Public License as published by the Free Software Foundation; either
+// version 2.1 of the License, or (at your option) any later version.
+// The full text of the license can be found in the file LICENSE.md at
+// the top level directory of deal.II.
+//
+// ---------------------------------------------------------------------
+
+// Test operator[], norm and norm_square of cuda_tensor.
+
+#include <deal.II/base/tensor.h>
+
+#include "../tests.h"
+
+void
+test_cpu()
+{
+  double             a[3][3] = {{1, 2, 3}, {4, 5, 6}, {7, 8, 9}};
+  const unsigned int dim     = 3;
+  Tensor<2, dim>     t;
+  for (unsigned int i = 0; i < dim; ++i)
+    for (unsigned int j = 0; j < dim; ++j)
+      t[i][j] = a[i][j];
+
+  deallog.push("values");
+  for (unsigned int i = 0; i < dim; ++i)
+    for (unsigned int j = 0; j < dim; ++j)
+      deallog << t[i][j] << std::endl;
+  deallog.pop();
+
+  deallog << "norm: " << t.norm() << std::endl;
+  deallog << "norm_square: " << t.norm_square() << std::endl;
+}
+
+void
+test_gpu()
+{
+  const unsigned int                                       dim = 3;
+  Kokkos::View<double, MemorySpace::Default::kokkos_space> norm_dev("norm_dev");
+  double                                                   norm_host;
+  Kokkos::View<double, MemorySpace::Default::kokkos_space> norm_square_dev(
+    "norm_square_dev");
+  double norm_square_host;
+  Kokkos::View<Tensor<2, dim>, MemorySpace::Default::kokkos_space> t_dev(
+    "t_dev");
+
+  using ExecutionSpace = MemorySpace::Default::kokkos_space::execution_space;
+  ExecutionSpace exec;
+
+  // Launch the kernels.
+  Kokkos::parallel_for(
+    Kokkos::MDRangePolicy<ExecutionSpace, Kokkos::Rank<2>>(exec,
+                                                           {{0, 0}},
+                                                           {{dim, dim}}),
+    KOKKOS_LAMBDA(int i, int j) { t_dev()[i][j] = j + i * dim + 1.; });
+  Kokkos::parallel_for(
+    Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+      norm_dev()        = t_dev().norm();
+      norm_square_dev() = t_dev().norm_square();
+    });
+  exec.fence();
+
+  // Copy the result to the host
+  Kokkos::deep_copy(norm_host, norm_dev);
+  Kokkos::deep_copy(norm_square_host, norm_square_dev);
+
+  // Output result
+  deallog << "norm GPU: " << norm_host << std::endl;
+  deallog << "norm_square GPU: " << norm_square_host << std::endl;
+}
+
+int
+main()
+{
+  initlog();
+
+  Kokkos::initialize();
+
+  test_cpu();
+
+  test_gpu();
+
+  Kokkos::finalize();
+}
diff --git a/tests/base/kokkos_tensor_02.cc b/tests/base/kokkos_tensor_02.cc
new file mode 100644 (file)
index 0000000..7581ecb
--- /dev/null
@@ -0,0 +1,220 @@
+// ---------------------------------------------------------------------
+//
+// Copyright (C) 2016 by the deal.II authors
+//
+// This file is part of the deal.II library.
+//
+// The deal.II library is free software; you can use it, redistribute
+// it, and/or modify it under the terms of the GNU Lesser General
+// Public License as published by the Free Software Foundation; either
+// version 2.1 of the License, or (at your option) any later version.
+// The full text of the license can be found in the file LICENSE.md at
+// the top level directory of deal.II.
+//
+// ---------------------------------------------------------------------
+
+// Test operator[] and norm_square of cuda_tensor.
+
+#include <deal.II/base/tensor.h>
+
+#include "../tests.h"
+
+template <int rank, int dim, typename Number>
+struct InitFunctor;
+
+template <int dim, typename Number>
+struct InitFunctor<0, dim, Number>
+{
+  Kokkos::View<Tensor<0, dim, Number>, MemorySpace::Default::kokkos_space> t;
+
+  KOKKOS_FUNCTION void
+  operator()(int k) const
+  {
+    t() = 1.;
+  }
+};
+
+template <int dim, typename Number>
+struct InitFunctor<1, dim, Number>
+{
+  Kokkos::View<Tensor<1, dim, Number>, MemorySpace::Default::kokkos_space> t;
+
+  KOKKOS_FUNCTION void
+  operator()(int k) const
+  {
+    t()[k] = k + 1.;
+  }
+};
+
+template <int dim, typename Number>
+struct InitFunctor<2, dim, Number>
+{
+  Kokkos::View<Tensor<2, dim, Number>, MemorySpace::Default::kokkos_space> t;
+
+  KOKKOS_FUNCTION void
+  operator()(int k) const
+  {
+    int i     = k / dim;
+    int j     = k % dim;
+    t()[i][j] = k + 1.;
+  }
+};
+
+
+template <int rank, int dim, typename Number>
+void
+test_gpu()
+{
+  const double tolerance = 1.e-8;
+
+  Kokkos::View<Tensor<rank, dim, Number>, MemorySpace::Default::kokkos_space>
+    t_dev("t_dev");
+  Kokkos::View<Tensor<rank, dim, Number>, MemorySpace::Default::kokkos_space>
+    t1_dev("t1_dev");
+  Kokkos::View<Tensor<rank, dim, Number>, MemorySpace::Default::kokkos_space>
+    t2_dev("t2_dev");
+
+  Tensor<rank, dim, Number> t_host;
+  Tensor<rank, dim, Number> t1_host;
+  Tensor<rank, dim, Number> t2_host;
+
+  Tensor<rank, dim, Number> reference_host;
+
+  using ExecutionSpace = MemorySpace::Default::kokkos_space::execution_space;
+  ExecutionSpace exec;
+
+  // Initialize
+  Kokkos::parallel_for(Kokkos::RangePolicy<ExecutionSpace>(
+                         exec, 0, Utilities::fixed_power<dim>(rank)),
+                       InitFunctor<rank, dim, Number>{t_dev});
+  Kokkos::deep_copy(reference_host, t_dev);
+
+  // Test multiplication.
+  Kokkos::parallel_for(
+    Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+      t1_dev() = t_dev() * Number(2.);
+      t2_dev() = Number(2.) * t_dev();
+      t_dev() *= 2.;
+    });
+  Kokkos::deep_copy(t_host, t_dev);
+  Kokkos::deep_copy(t1_host, t1_dev);
+  Kokkos::deep_copy(t2_host, t2_dev);
+
+  reference_host *= 2;
+  AssertThrow((t_host - reference_host).norm() < tolerance, ExcInternalError());
+  AssertThrow((t1_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+  AssertThrow((t2_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+
+  deallog << "multiplication OK" << std::endl;
+
+  // Test division.
+  Kokkos::parallel_for(
+    Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+      t1_dev() = t_dev() / Number(2.);
+      t_dev() /= 2.;
+      t2_dev() = t1_dev();
+    });
+  Kokkos::deep_copy(t_host, t_dev);
+  Kokkos::deep_copy(t1_host, t1_dev);
+
+  reference_host /= 2.;
+  AssertThrow((t_host - reference_host).norm() < tolerance, ExcInternalError());
+  AssertThrow((t1_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+
+  deallog << "division OK" << std::endl;
+
+  // Test summation
+  Kokkos::parallel_for(
+    Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+      t2_dev() += t_dev();
+      t1_dev() = t1_dev() + t_dev();
+    });
+  Kokkos::deep_copy(t1_host, t1_dev);
+  Kokkos::deep_copy(t2_host, t2_dev);
+
+  reference_host *= 2.;
+  AssertThrow((t1_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+  AssertThrow((t2_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+
+
+  // Test subtraction
+  Kokkos::parallel_for(
+    Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+      t2_dev() -= t_dev();
+      t1_dev() = t1_dev() - t_dev();
+    });
+  Kokkos::deep_copy(t1_host, t1_dev);
+  Kokkos::deep_copy(t2_host, t2_dev);
+
+  reference_host /= 2.;
+  AssertThrow((t1_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+  AssertThrow((t2_host - reference_host).norm() < tolerance,
+              ExcInternalError());
+
+  // Miscellaneous
+  {
+    Kokkos::View<Number, MemorySpace::Default::kokkos_space> check_1("check_1");
+    Kokkos::View<Number, MemorySpace::Default::kokkos_space> check_2("check_2");
+    Kokkos::View<Number, MemorySpace::Default::kokkos_space> check_3("check_3");
+    Kokkos::View<Number, MemorySpace::Default::kokkos_space> check_4("check_4");
+    Kokkos::View<Number, MemorySpace::Default::kokkos_space> check_5("check_5");
+
+    Kokkos::parallel_for(
+      Kokkos::RangePolicy<ExecutionSpace>(exec, 0, 1), KOKKOS_LAMBDA(int) {
+        // constructors
+        typename Tensor<rank, dim, Number>::array_type array{};
+        Tensor<rank, dim, Number>                      dummy_1(array);
+        check_1() = dummy_1.norm_square();
+        Tensor<rank, dim, Number> dummy_2;
+        check_2()                         = dummy_2.norm_square();
+        Tensor<rank, dim, Number> dummy_3 = dummy_2;
+        check_3()                         = dummy_3.norm_square();
+
+        // access
+        Tensor<rank + 1, dim, Number>   initializer_1;
+        const Tensor<rank, dim, Number> dummy_5 = initializer_1[0];
+        check_4()                               = dummy_5.norm_square();
+
+        // assignment
+        dummy_2   = dummy_3;
+        check_5() = dummy_2.norm_square();
+      });
+
+    Number check_1_host, check_2_host, check_3_host, check_4_host, check_5_host;
+
+    Kokkos::deep_copy(check_1_host, check_1);
+    Kokkos::deep_copy(check_2_host, check_2);
+    Kokkos::deep_copy(check_3_host, check_3);
+    Kokkos::deep_copy(check_4_host, check_4);
+    Kokkos::deep_copy(check_5_host, check_5);
+
+    AssertThrow(std::abs(check_1_host) < tolerance, ExcInternalError());
+    AssertThrow(std::abs(check_2_host) < tolerance, ExcInternalError());
+    AssertThrow(std::abs(check_3_host) < tolerance, ExcInternalError());
+    AssertThrow(std::abs(check_4_host) < tolerance, ExcInternalError());
+    AssertThrow(std::abs(check_5_host) < tolerance, ExcInternalError());
+  }
+}
+
+int
+main()
+{
+  initlog();
+
+  Kokkos::initialize();
+
+  test_gpu<0, 3, double>();
+  test_gpu<1, 3, double>();
+  test_gpu<2, 3, double>();
+  test_gpu<0, 3, float>();
+  test_gpu<1, 3, float>();
+  test_gpu<2, 3, float>();
+
+  Kokkos::finalize();
+}
diff --git a/tests/cuda/array_view_wrong_memory.cc b/tests/cuda/array_view_wrong_memory.cc
deleted file mode 100644 (file)
index a8a58e8..0000000
+++ /dev/null
@@ -1,70 +0,0 @@
-// ---------------------------------------------------------------------
-//
-// Copyright (C) 2018 by the deal.II authors
-//
-// This file is part of the deal.II library.
-//
-// The deal.II library is free software; you can use it, redistribute
-// it, and/or modify it under the terms of the GNU Lesser General
-// Public License as published by the Free Software Foundation; either
-// version 2.1 of the License, or (at your option) any later version.
-// The full text of the license can be found in the file LICENSE.md at
-// the top level directory of deal.II.
-//
-// ---------------------------------------------------------------------
-
-
-// check that we detect creating ArrayView objects with the wrong memory space.
-
-#include <deal.II/base/array_view.h>
-
-#include "../tests.h"
-
-int
-main(int argc, char **argv)
-{
-  deal_II_exceptions::disable_abort_on_exception();
-
-  initlog();
-
-  init_cuda();
-
-  std::vector<unsigned int>                                 dummy_host(2);
-  std::unique_ptr<unsigned int[], void (*)(unsigned int *)> dummy_cuda(
-    Utilities::CUDA::allocate_device_data<unsigned int>(2),
-    Utilities::CUDA::delete_device_data<unsigned int>);
-
-  deallog << "Testing host ArrayView with host memory" << std::endl;
-  ArrayView<unsigned int, MemorySpace::Host> view_1(dummy_host);
-
-  deallog << "Testing device ArrayView with host memory" << std::endl;
-  try
-    {
-      ArrayView<unsigned int, MemorySpace::CUDA> view_2(dummy_host);
-    }
-  catch (const ExceptionBase &exc)
-    {
-      deallog << exc.what() << std::endl;
-    }
-
-  deallog << "Testing host ArrayView with device memory" << std::endl;
-  try
-    {
-      ArrayView<unsigned int, MemorySpace::Host> view_3(dummy_cuda.get(), 2);
-    }
-  catch (const ExceptionBase &exc)
-    {
-      deallog << exc.what() << std::endl;
-    }
-
-  deallog << "Testing device ArrayView with device memory" << std::endl;
-  ArrayView<unsigned int, MemorySpace::CUDA> view_4(dummy_cuda.get(), 2);
-
-  deallog << "Testing host ArrayView to a nullptr with length 0" << std::endl;
-  ArrayView<unsigned int, MemorySpace::Host> view_5(nullptr, 0);
-
-  deallog << "Testing device ArrayView to a nullptr with length 0" << std::endl;
-  ArrayView<unsigned int, MemorySpace::CUDA> view_6(nullptr, 0);
-
-  return 0;
-}
diff --git a/tests/cuda/array_view_wrong_memory.debug.output b/tests/cuda/array_view_wrong_memory.debug.output
deleted file mode 100644 (file)
index 4659b11..0000000
+++ /dev/null
@@ -1,29 +0,0 @@
-
-DEAL::Testing host ArrayView with host memory
-DEAL::Testing device ArrayView with host memory
-DEAL::
---------------------------------------------------------
-An error occurred in file <array_view.h> in function
-    dealii::ArrayView<ElementType, MemorySpace>::ArrayView(dealii::ArrayView<ElementType, MemorySpace>::value_type*, std::size_t) [with ElementType = unsigned int; MemorySpaceType = dealii::MemorySpace::Default; dealii::ArrayView<ElementType, MemorySpace>::value_type = unsigned int; std::size_t = long unsigned int]
-The violated condition was: 
-    n_elements == 0 || internal::ArrayViewHelper::is_in_correct_memory_space<MemorySpaceType>( starting_element)
-Additional information: 
-    The memory space indicated by the template parameter and the one
-    derived from the pointer value do not match!
---------------------------------------------------------
-
-DEAL::Testing host ArrayView with device memory
-DEAL::
---------------------------------------------------------
-An error occurred in file <array_view.h> in function
-    dealii::ArrayView<ElementType, MemorySpace>::ArrayView(dealii::ArrayView<ElementType, MemorySpace>::value_type*, std::size_t) [with ElementType = unsigned int; MemorySpaceType = dealii::MemorySpace::Host; dealii::ArrayView<ElementType, MemorySpace>::value_type = unsigned int; std::size_t = long unsigned int]
-The violated condition was: 
-    n_elements == 0 || internal::ArrayViewHelper::is_in_correct_memory_space<MemorySpaceType>( starting_element)
-Additional information: 
-    The memory space indicated by the template parameter and the one
-    derived from the pointer value do not match!
---------------------------------------------------------
-
-DEAL::Testing device ArrayView with device memory
-DEAL::Testing host ArrayView to a nullptr with length 0
-DEAL::Testing device ArrayView to a nullptr with length 0
diff --git a/tests/cuda/cuda_tensor_01.cc b/tests/cuda/cuda_tensor_01.cc
deleted file mode 100644 (file)
index 68e1f27..0000000
+++ /dev/null
@@ -1,117 +0,0 @@
-// ---------------------------------------------------------------------
-//
-// Copyright (C) 2016 by the deal.II authors
-//
-// This file is part of the deal.II library.
-//
-// The deal.II library is free software; you can use it, redistribute
-// it, and/or modify it under the terms of the GNU Lesser General
-// Public License as published by the Free Software Foundation; either
-// version 2.1 of the License, or (at your option) any later version.
-// The full text of the license can be found in the file LICENSE.md at
-// the top level directory of deal.II.
-//
-// ---------------------------------------------------------------------
-
-// Test operator[], norm and norm_square of cuda_tensor.
-
-#include <deal.II/base/tensor.h>
-
-#include "../tests.h"
-
-void
-test_cpu()
-{
-  double             a[3][3] = {{1, 2, 3}, {4, 5, 6}, {7, 8, 9}};
-  const unsigned int dim     = 3;
-  Tensor<2, dim>     t;
-  for (unsigned int i = 0; i < dim; ++i)
-    for (unsigned int j = 0; j < dim; ++j)
-      t[i][j] = a[i][j];
-
-  deallog.push("values");
-  for (unsigned int i = 0; i < dim; ++i)
-    for (unsigned int j = 0; j < dim; ++j)
-      deallog << t[i][j] << std::endl;
-  deallog.pop();
-
-  deallog << "norm: " << t.norm() << std::endl;
-  deallog << "norm_square: " << t.norm_square() << std::endl;
-}
-
-__global__ void
-init_kernel(Tensor<2, 3> *t, const unsigned int N)
-{
-  const unsigned int i = threadIdx.y;
-  const unsigned int j = threadIdx.x;
-  if ((i < N) && (j < N))
-    (*t)[i][j] = j + i * N + 1.;
-}
-
-__global__ void
-norm_kernel(Tensor<2, 3> *t, double *norm, double *norm_square)
-{
-  if (threadIdx.x == 0)
-    {
-      *norm        = t->norm();
-      *norm_square = t->norm_square();
-    }
-}
-
-void
-test_gpu()
-{
-  const unsigned int dim = 3;
-  double *           norm_dev;
-  double             norm_host;
-  double *           norm_square_dev;
-  double             norm_square_host;
-  Tensor<2, dim> *   t_dev;
-
-  // Allocate objects on the device
-  cudaError_t cuda_error = cudaMalloc(&t_dev, sizeof(Tensor<2, dim>));
-  AssertCuda(cuda_error);
-  cuda_error = cudaMalloc(&norm_dev, sizeof(double));
-  AssertCuda(cuda_error);
-  cuda_error = cudaMalloc(&norm_square_dev, sizeof(double));
-  AssertCuda(cuda_error);
-
-  // Launch the kernels.
-  dim3 block_dim(dim, dim);
-  init_kernel<<<1, block_dim>>>(t_dev, dim);
-  norm_kernel<<<1, 1>>>(t_dev, norm_dev, norm_square_dev);
-
-  // Copy the result to the device
-  cuda_error =
-    cudaMemcpy(&norm_host, norm_dev, sizeof(double), cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-  cuda_error = cudaMemcpy(&norm_square_host,
-                          norm_square_dev,
-                          sizeof(double),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-
-  // Free memory
-  cuda_error = cudaFree(t_dev);
-  AssertCuda(cuda_error);
-  cuda_error = cudaFree(norm_dev);
-  AssertCuda(cuda_error);
-  cuda_error = cudaFree(norm_square_dev);
-  AssertCuda(cuda_error);
-
-  // Output result
-  deallog << "norm GPU: " << norm_host << std::endl;
-  deallog << "norm_square GPU: " << norm_square_host << std::endl;
-}
-
-int
-main()
-{
-  initlog();
-
-  init_cuda();
-
-  test_cpu();
-
-  test_gpu();
-}
diff --git a/tests/cuda/cuda_tensor_02.cc b/tests/cuda/cuda_tensor_02.cc
deleted file mode 100644 (file)
index e3a5226..0000000
+++ /dev/null
@@ -1,329 +0,0 @@
-// ---------------------------------------------------------------------
-//
-// Copyright (C) 2016 by the deal.II authors
-//
-// This file is part of the deal.II library.
-//
-// The deal.II library is free software; you can use it, redistribute
-// it, and/or modify it under the terms of the GNU Lesser General
-// Public License as published by the Free Software Foundation; either
-// version 2.1 of the License, or (at your option) any later version.
-// The full text of the license can be found in the file LICENSE.md at
-// the top level directory of deal.II.
-//
-// ---------------------------------------------------------------------
-
-// Test operator[] and norm_square of cuda_tensor.
-
-#include <deal.II/base/tensor.h>
-
-#include "../tests.h"
-
-template <int rank, int dim, typename Number>
-__global__ void
-miscellaneous_kernel(Number *check_1,
-                     Number *check_2,
-                     Number *check_3,
-                     Number *check_4,
-                     Number *check_5)
-{
-  // constructors
-  typename Tensor<rank, dim, Number>::array_type array{};
-  Tensor<rank, dim, Number>                      dummy_1(array);
-  *check_1 = dummy_1.norm_square();
-  Tensor<rank, dim, Number> dummy_2;
-  *check_2                          = dummy_2.norm_square();
-  Tensor<rank, dim, Number> dummy_3 = dummy_2;
-  *check_3                          = dummy_3.norm_square();
-
-  // access
-  Tensor<rank + 1, dim, Number>   initializer_1;
-  const Tensor<rank, dim, Number> dummy_5 = initializer_1[0];
-  *check_4                                = dummy_5.norm_square();
-
-  // assignment
-  dummy_2  = dummy_3;
-  *check_5 = dummy_2.norm_square();
-}
-
-template <int rank, int dim, typename Number>
-__global__ void
-summation_kernel(Tensor<rank, dim, Number> *t,
-                 Tensor<rank, dim, Number> *t1,
-                 Tensor<rank, dim, Number> *t2)
-{
-  *t2 += *t;
-  *t1 = *t1 + *t;
-}
-
-template <int rank, int dim, typename Number>
-__global__ void
-subtraction_kernel(Tensor<rank, dim, Number> *t,
-                   Tensor<rank, dim, Number> *t1,
-                   Tensor<rank, dim, Number> *t2)
-{
-  *t2 -= *t;
-  *t1 = *t1 - *t;
-}
-
-template <int rank, int dim, typename Number>
-__global__ void
-multiplication_kernel(Tensor<rank, dim, Number> *t,
-                      Tensor<rank, dim, Number> *t1,
-                      Tensor<rank, dim, Number> *t2)
-{
-  *t1 = *t * Number(2.);
-  *t2 = Number(2.) * *t;
-  *t *= 2.;
-}
-
-template <int rank, int dim, typename Number>
-__global__ void
-division_kernel(Tensor<rank, dim, Number> *t,
-                Tensor<rank, dim, Number> *t1,
-                Tensor<rank, dim, Number> *t2)
-{
-  *t1 = *t / Number(2.);
-  *t /= 2.;
-  *t2 = *t1;
-}
-
-template <int dim, typename Number>
-__global__ void
-init_kernel(Tensor<0, dim, Number> *t)
-{
-  if (threadIdx.x == 0)
-    *t = 1.;
-}
-
-template <int dim, typename Number>
-__global__ void
-init_kernel(Tensor<1, dim, Number> *t)
-{
-  const unsigned int i = threadIdx.x;
-  if (i < dim)
-    (*t)[i] = i + 1.;
-}
-
-template <int dim, typename Number>
-__global__ void
-init_kernel(Tensor<2, dim, Number> *t)
-{
-  const unsigned int i = threadIdx.y;
-  const unsigned int j = threadIdx.x;
-  if ((i < dim) && (j < dim))
-    (*t)[i][j] = j + i * dim + 1.;
-}
-
-
-template <int rank, int dim, typename Number>
-void
-test_gpu()
-{
-  const double tolerance = 1.e-8;
-
-  Tensor<rank, dim, Number> *t_dev;
-  Tensor<rank, dim, Number> *t1_dev;
-  Tensor<rank, dim, Number> *t2_dev;
-
-  Tensor<rank, dim, Number> t_host;
-  Tensor<rank, dim, Number> t1_host;
-  Tensor<rank, dim, Number> t2_host;
-
-  Tensor<rank, dim, Number> reference_host;
-
-  // Allocate objects on the device
-  cudaError_t cuda_error =
-    cudaMalloc(&t_dev, sizeof(Tensor<rank, dim, Number>));
-  AssertCuda(cuda_error);
-  cuda_error = cudaMalloc(&t1_dev, sizeof(Tensor<rank, dim, Number>));
-  AssertCuda(cuda_error);
-  cuda_error = cudaMalloc(&t2_dev, sizeof(Tensor<rank, dim, Number>));
-  AssertCuda(cuda_error);
-
-  // Initialize
-  dim3 block_dim(dim, dim);
-  init_kernel<<<1, block_dim>>>(t_dev);
-  cuda_error = cudaMemcpy(&reference_host,
-                          t_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-
-  // Test multiplication.
-  multiplication_kernel<<<1, 1>>>(t_dev, t1_dev, t2_dev);
-
-  cuda_error = cudaMemcpy(&t_host,
-                          t_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-  cuda_error = cudaMemcpy(&t1_host,
-                          t1_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-  cuda_error = cudaMemcpy(&t2_host,
-                          t2_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-
-  reference_host *= 2;
-  AssertThrow((t_host - reference_host).norm() < tolerance, ExcInternalError());
-  AssertThrow((t1_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-  AssertThrow((t2_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-
-  deallog << "multiplication OK" << std::endl;
-
-  // Test division.
-  division_kernel<<<1, 1>>>(t_dev, t1_dev, t2_dev);
-  cuda_error = cudaMemcpy(&t_host,
-                          t_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-  cuda_error = cudaMemcpy(&t1_host,
-                          t1_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-
-  reference_host /= 2.;
-  AssertThrow((t_host - reference_host).norm() < tolerance, ExcInternalError());
-  AssertThrow((t1_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-
-  deallog << "division OK" << std::endl;
-
-  // Test summation
-  summation_kernel<<<1, 1>>>(t_dev, t1_dev, t2_dev);
-  cuda_error = cudaMemcpy(&t1_host,
-                          t1_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-  cuda_error = cudaMemcpy(&t2_host,
-                          t2_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-
-  reference_host *= 2.;
-  AssertThrow((t1_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-  AssertThrow((t2_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-
-
-  // Test subtraction
-  subtraction_kernel<<<1, 1>>>(t_dev, t1_dev, t2_dev);
-  cuda_error = cudaMemcpy(&t1_host,
-                          t1_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-  AssertCuda(cuda_error);
-  cuda_error = cudaMemcpy(&t2_host,
-                          t2_dev,
-                          sizeof(Tensor<rank, dim, Number>),
-                          cudaMemcpyDeviceToHost);
-
-  reference_host /= 2.;
-  AssertThrow((t1_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-  AssertThrow((t2_host - reference_host).norm() < tolerance,
-              ExcInternalError());
-
-  // Free memory
-  cuda_error = cudaFree(t_dev);
-  AssertCuda(cuda_error);
-  cuda_error = cudaFree(t1_dev);
-  AssertCuda(cuda_error);
-  cuda_error = cudaFree(t2_dev);
-  AssertCuda(cuda_error);
-
-  // Miscellaneous
-  {
-    Number *check_1;
-    Number *check_2;
-    Number *check_3;
-    Number *check_4;
-    Number *check_5;
-
-    cuda_error = cudaMalloc(&check_1, sizeof(Number));
-    AssertCuda(cuda_error);
-    cuda_error = cudaMalloc(&check_2, sizeof(Number));
-    AssertCuda(cuda_error);
-    cuda_error = cudaMalloc(&check_3, sizeof(Number));
-    AssertCuda(cuda_error);
-    cuda_error = cudaMalloc(&check_4, sizeof(Number));
-    AssertCuda(cuda_error);
-    cuda_error = cudaMalloc(&check_5, sizeof(Number));
-    AssertCuda(cuda_error);
-
-    miscellaneous_kernel<rank, dim, Number>
-      <<<1, 1>>>(check_1, check_2, check_3, check_4, check_5);
-
-    Number check_1_host, check_2_host, check_3_host, check_4_host, check_5_host;
-
-    cuda_error = cudaMemcpy(&check_1_host,
-                            check_1,
-                            sizeof(Number),
-                            cudaMemcpyDeviceToHost);
-    AssertCuda(cuda_error);
-    cuda_error = cudaMemcpy(&check_2_host,
-                            check_2,
-                            sizeof(Number),
-                            cudaMemcpyDeviceToHost);
-    AssertCuda(cuda_error);
-    cuda_error = cudaMemcpy(&check_3_host,
-                            check_3,
-                            sizeof(Number),
-                            cudaMemcpyDeviceToHost);
-    AssertCuda(cuda_error);
-    cuda_error = cudaMemcpy(&check_4_host,
-                            check_4,
-                            sizeof(Number),
-                            cudaMemcpyDeviceToHost);
-    AssertCuda(cuda_error);
-    cuda_error = cudaMemcpy(&check_5_host,
-                            check_5,
-                            sizeof(Number),
-                            cudaMemcpyDeviceToHost);
-    AssertCuda(cuda_error);
-
-    AssertThrow(std::abs(check_1_host) < tolerance, ExcInternalError());
-    AssertThrow(std::abs(check_2_host) < tolerance, ExcInternalError());
-    AssertThrow(std::abs(check_3_host) < tolerance, ExcInternalError());
-    AssertThrow(std::abs(check_4_host) < tolerance, ExcInternalError());
-    AssertThrow(std::abs(check_5_host) < tolerance, ExcInternalError());
-
-    cuda_error = cudaFree(check_1);
-    AssertCuda(cuda_error);
-    cuda_error = cudaFree(check_2);
-    AssertCuda(cuda_error);
-    cuda_error = cudaFree(check_3);
-    AssertCuda(cuda_error);
-    cuda_error = cudaFree(check_4);
-    AssertCuda(cuda_error);
-    cuda_error = cudaFree(check_5);
-    AssertCuda(cuda_error);
-  }
-}
-
-int
-main()
-{
-  initlog();
-
-  init_cuda();
-
-  test_gpu<0, 3, double>();
-  test_gpu<1, 3, double>();
-  test_gpu<2, 3, double>();
-  test_gpu<0, 3, float>();
-  test_gpu<1, 3, float>();
-  test_gpu<2, 3, float>();
-}
similarity index 74%
rename from tests/cuda/affine_constraints_set_zero.cc
rename to tests/lac/affine_constraints_set_zero.cc
index 9b497f2fb8963763ad7a9425331b1ebbb6d8e245..e64475d28406695c67c3b58be7927ffac9cd744d 100644 (file)
 #include "../tests.h"
 
 
-__global__ void
-initialize_vector(double *vector, int local_size, int offset)
-{
-  const int index = threadIdx.x + blockIdx.x * blockDim.x;
-  if (index < local_size)
-    vector[index] = 1.0 + index + offset;
-}
-
-
 void
 test()
 {
@@ -56,15 +47,23 @@ test()
   deallog << "CM:" << std::endl;
   cm.print(deallog.get_file_stream());
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> ghosted;
+  using ExecutionSpace = MemorySpace::Default::kokkos_space::execution_space;
+  ExecutionSpace exec;
+
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> ghosted;
   {
     ghosted.reinit(local_active,
                    complete_index_set(2 * numproc),
                    MPI_COMM_WORLD);
-
-    const int n_blocks = 1 + ghosted.size() / CUDAWrappers::block_size;
-    initialize_vector<<<n_blocks, CUDAWrappers::block_size>>>(
-      ghosted.get_values(), numproc, myid * numproc);
+    auto ghosted_values = ghosted.get_values();
+
+    Kokkos::parallel_for(
+      Kokkos::RangePolicy<ExecutionSpace>(exec, 0, numproc),
+      KOKKOS_LAMBDA(int i) {
+        int offset        = myid * numproc;
+        ghosted_values[i] = 1.0 + i + offset;
+      });
+    exec.fence();
     ghosted.compress(VectorOperation::insert);
 
     deallog << "ghosted vector before:" << std::endl;
@@ -76,15 +75,22 @@ test()
     ghosted.print(deallog.get_file_stream());
   }
 
-  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> distributed;
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> distributed;
   {
     distributed.reinit(local_active,
                        complete_index_set(2 * numproc),
                        MPI_COMM_WORLD);
 
-    const int n_blocks = 1 + distributed.size() / CUDAWrappers::block_size;
-    initialize_vector<<<n_blocks, CUDAWrappers::block_size>>>(
-      distributed.get_values(), numproc, myid * numproc);
+    auto distributed_values = distributed.get_values();
+
+    Kokkos::parallel_for(
+      Kokkos::RangePolicy<ExecutionSpace>(exec, 0, numproc),
+      KOKKOS_LAMBDA(int i) {
+        int offset            = myid * numproc;
+        distributed_values[i] = 1.0 + i + offset;
+      });
+    exec.fence();
+
     distributed.compress(VectorOperation::insert);
 
     deallog << "distributed vector before:" << std::endl;
@@ -106,8 +112,6 @@ main(int argc, char *argv[])
   Utilities::MPI::MPI_InitFinalize mpi_initialization(argc, argv, 1);
   MPILogInitAll                    log;
 
-  init_cuda();
-
   test();
   return 0;
 }
similarity index 99%
rename from tests/cuda/vector_reinit_01.cc
rename to tests/lac/vector_reinit_04.cc
index 1bdfed9dc18f49be649badfc6707f53159b299f4..88b3ccdf1df862d7b764faf8e0af75e3ce5e20c4 100644 (file)
@@ -112,7 +112,6 @@ main(int argc, char **argv)
     argc, argv, testing_max_num_threads());
 
   initlog();
-  init_cuda();
 
   do_test<
     LinearAlgebra::distributed::Vector<double, dealii::MemorySpace::CUDA>>();

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.