Remove DEAL_II_COMPILER_CUDA_AWARE

diff --git a/doc/doxygen/options.dox.in b/doc/doxygen/options.dox.in
index 7d1cd2c3b7e81700bc516ea9782fc66e4e651881..6af5f28efdbca7724877ea20b570d006e2f02192 100644 (file)
--- a/doc/doxygen/options.dox.in
+++ b/doc/doxygen/options.dox.in
@@ -200,7 +200,6 @@ PREDEFINED             = DOXYGEN=1 \
                          DEAL_II_WITH_TASKFLOW=1 \
                          DEAL_II_WITH_COMPLEX_VALUES=1 \
                          DEAL_II_WITH_CUDA=1 \
-                         DEAL_II_COMPILER_CUDA_AWARE=1 \
                          DEAL_II_WITH_GINKGO=1 \
                          DEAL_II_WITH_GMSH=1 \
                          DEAL_II_GMSH_WITH_API=1 \

diff --git a/include/deal.II/base/config.h.in b/include/deal.II/base/config.h.in
index 0a36b957ac4bae5370674c6b6d4d58b6e78a1711..3745b038b5bd32de44a9fb0e4be0b104fab8c173 100644 (file)
--- a/include/deal.II/base/config.h.in
+++ b/include/deal.II/base/config.h.in
@@ -103,14 +103,6 @@
 #cmakedefine DEAL_II_RESTRICT @DEAL_II_RESTRICT@
 #cmakedefine DEAL_II_COMPILER_HAS_DIAGNOSTIC_PRAGMA
 
-/*
- * A variable to tell if the compiler used in the current compilation process
- * understands CUDA code.
- */
-#if defined(DEAL_II_WITH_CUDA) && defined(__CUDACC__)
-#  define DEAL_II_COMPILER_CUDA_AWARE
-#endif
-
 /***********************************************************************
  * CPU features:
  *

diff --git a/include/deal.II/base/cuda.h b/include/deal.II/base/cuda.h
index f6cf36917d22f4e17602a8e8ba21ae2c174e4730..d882804dd4b102c29902b7305b3c0be4d9b84f05 100644 (file)
--- a/include/deal.II/base/cuda.h
+++ b/include/deal.II/base/cuda.h
@@ -21,7 +21,7 @@
 #include <deal.II/base/array_view.h>
 #include <deal.II/base/exceptions.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 #  include <cusolverDn.h>
 #  include <cusolverSp.h>
 #  include <cusparse.h>

diff --git a/include/deal.II/base/numbers.h b/include/deal.II/base/numbers.h
index 8c8b7382ff267131f7e95bf5a1f04d4675f8a438..0e13bb2a415200fd7e665f2ff45c90f012044f1e 100644 (file)
--- a/include/deal.II/base/numbers.h
+++ b/include/deal.II/base/numbers.h
@@ -21,7 +21,7 @@
 
 #include <deal.II/base/types.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 #  include <cuComplex.h>
 #endif
 
@@ -809,7 +809,7 @@ namespace internal
     }
   };
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
   template <>
   struct NumberType<cuComplex>
   {

diff --git a/include/deal.II/base/partitioner.templates.h b/include/deal.II/base/partitioner.templates.h
index 78ab02140f1e861275138cb83eadc7c8f9ad0ffd..c1d7a6856ea7cfb0a0802a48555ff4af085856bd 100644 (file)
--- a/include/deal.II/base/partitioner.templates.h
+++ b/include/deal.II/base/partitioner.templates.h
@@ -107,8 +107,7 @@ namespace Utilities
         }
 
       Number *temp_array_ptr = temporary_storage.data();
-#    if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
-      defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+#    if defined(DEAL_II_WITH_CUDA) && defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
       // When using CUDAs-aware MPI, the set of local indices that are ghosts
       // indices on other processors is expanded in arrays. This is for
       // performance reasons as this can significantly decrease the number of
@@ -121,8 +120,7 @@ namespace Utilities
 
       for (unsigned int i = 0; i < n_import_targets; ++i)
         {
-#    if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
-      defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+#    if defined(DEAL_II_WITH_CUDA) && defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
           if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
             {
               const auto chunk_size = import_indices_plain_dev[i].second;
@@ -228,7 +226,7 @@ namespace Utilities
                     }
                   else
                     {
-#    ifdef DEAL_II_COMPILER_CUDA_AWARE
+#    ifdef DEAL_II_WITH_CUDA
                       cudaError_t cuda_error =
                         cudaMemcpy(ghost_array.data() + ghost_range.first,
                                    ghost_array.data() + offset,
@@ -378,7 +376,7 @@ namespace Utilities
                         }
                       else
                         {
-#    ifdef DEAL_II_COMPILER_CUDA_AWARE
+#    ifdef DEAL_II_WITH_CUDA
                           cudaError_t cuda_error =
                             cudaMemcpy(ghost_array_ptr + offset,
                                        ghost_array.data() + my_ghosts->first,
@@ -414,8 +412,7 @@ namespace Utilities
             ExcMessage("Index overflow: Maximum message size in MPI is 2GB. "
                        "The number of ghost entries times the size of 'Number' "
                        "exceeds this value. This is not supported."));
-#    if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
-      defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+#    if defined(DEAL_II_WITH_CUDA) && defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
           if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
             cudaDeviceSynchronize();
 #    endif
@@ -526,7 +523,7 @@ namespace Utilities
                    "import_from_ghosted_array_start as is passed "
                    "to import_from_ghosted_array_finish."));
 
-#      ifdef DEAL_II_COMPILER_CUDA_AWARE
+#      ifdef DEAL_II_WITH_CUDA
           if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
             {
               cudaMemset(ghost_array.data(),
@@ -560,8 +557,7 @@ namespace Utilities
       const unsigned int n_import_targets = import_targets_data.size();
       const unsigned int n_ghost_targets  = ghost_targets_data.size();
 
-#    if (defined(DEAL_II_COMPILER_CUDA_AWARE) && \
-         defined(DEAL_II_MPI_WITH_CUDA_SUPPORT))
+#    if (defined(DEAL_II_WITH_CUDA) && defined(DEAL_II_MPI_WITH_CUDA_SUPPORT))
       // When using CUDAs-aware MPI, the set of local indices that are ghosts
       // indices on other processors is expanded in arrays. This is for
       // performance reasons as this can significantly decrease the number of
@@ -583,8 +579,7 @@ namespace Utilities
           AssertThrowMPI(ierr);
 
           const Number *read_position = temporary_storage.data();
-#    if !(defined(DEAL_II_COMPILER_CUDA_AWARE) && \
-          defined(DEAL_II_MPI_WITH_CUDA_SUPPORT))
+#    if !(defined(DEAL_II_WITH_CUDA) && defined(DEAL_II_MPI_WITH_CUDA_SUPPORT))
           // If the operation is no insertion, add the imported data to the
           // local values. For insert, nothing is done here (but in debug mode
           // we assert that the specified value is either zero or matches with
@@ -730,8 +725,7 @@ namespace Utilities
         {
           Assert(ghost_array.begin() != nullptr, ExcInternalError());
 
-#    if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
-      defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+#    if defined(DEAL_II_WITH_CUDA) && defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
           if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
             {
               Assert(std::is_trivial<Number>::value, ExcNotImplemented());

diff --git a/include/deal.II/base/tensor.h b/include/deal.II/base/tensor.h
index 551731a4028168cf6deca249bc351281a15cb51b..55db04f176d385d8fff5e29561000aaff4e5904a 100644 (file)
--- a/include/deal.II/base/tensor.h
+++ b/include/deal.II/base/tensor.h
@@ -1432,8 +1432,12 @@ constexpr DEAL_II_ALWAYS_INLINE
   DEAL_II_HOST_DEVICE const typename Tensor<rank_, dim, Number>::value_type &
   Tensor<rank_, dim, Number>::operator[](const unsigned int i) const
 {
-#  ifndef DEAL_II_COMPILER_CUDA_AWARE
+#  if KOKKOS_VERSION < 30700
+#    ifdef KOKKOS_ACTIVE_MEMORY_SPACE_HOST
   AssertIndexRange(i, dim);
+#    endif
+#  else
+  KOKKOS_IF_ON_HOST((AssertIndexRange(i, dim);))
 #  endif
 
   return values[i];
@@ -1444,9 +1448,16 @@ template <int rank_, int dim, typename Number>
 constexpr inline DEAL_II_ALWAYS_INLINE const Number &
 Tensor<rank_, dim, Number>::operator[](const TableIndices<rank_> &indices) const
 {
-#  ifndef DEAL_II_COMPILER_CUDA_AWARE
+#  if KOKKOS_VERSION < 30700
+#    ifdef KOKKOS_ACTIVE_MEMORY_SPACE_HOST
   Assert(dim != 0,
          ExcMessage("Cannot access an object of type Tensor<rank_,0,Number>"));
+#    endif
+#  else
+  KOKKOS_IF_ON_HOST(
+    (Assert(dim != 0,
+            ExcMessage(
+              "Cannot access an object of type Tensor<rank_,0,Number>"));))
 #  endif
 
   return TensorAccessors::extract<rank_>(*this, indices);
@@ -1458,9 +1469,16 @@ template <int rank_, int dim, typename Number>
 constexpr inline DEAL_II_ALWAYS_INLINE Number &
 Tensor<rank_, dim, Number>::operator[](const TableIndices<rank_> &indices)
 {
-#  ifndef DEAL_II_COMPILER_CUDA_AWARE
+#  if KOKKOS_VERSION < 30700
+#    ifdef KOKKOS_ACTIVE_MEMORY_SPACE_HOST
   Assert(dim != 0,
          ExcMessage("Cannot access an object of type Tensor<rank_,0,Number>"));
+#    endif
+#  else
+  KOKKOS_IF_ON_HOST(
+    (Assert(dim != 0,
+            ExcMessage(
+              "Cannot access an object of type Tensor<rank_,0,Number>"));))
 #  endif
 
   return TensorAccessors::extract<rank_>(*this, indices);

diff --git a/include/deal.II/lac/affine_constraints.templates.h b/include/deal.II/lac/affine_constraints.templates.h
index 044e35ade92ba0ede0de335f5e196a95ffbf0be2..9b2c26cf730bcbee2a5655e3ad84a2a9b84acf2d 100644 (file)
--- a/include/deal.II/lac/affine_constraints.templates.h
+++ b/include/deal.II/lac/affine_constraints.templates.h
@@ -2256,7 +2256,7 @@ namespace internal
       vec.zero_out_ghost_values();
     }
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
     template <typename Number>
     __global__ void
     set_zero_kernel(const size_type *  constrained_dofs,

diff --git a/include/deal.II/lac/cuda_atomic.h b/include/deal.II/lac/cuda_atomic.h
index 9edcbcde4631ad4e6667b56f4a18a5d9a6a789dc..82c451c1c256947a64ae3dab87a37427b67d4c06 100644 (file)
--- a/include/deal.II/lac/cuda_atomic.h
+++ b/include/deal.II/lac/cuda_atomic.h
@@ -18,7 +18,7 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 DEAL_II_NAMESPACE_OPEN
 

diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h
index 564a461f800fc14a90129dec0d6810922bb052f4..178ad998d7c0727d5c657b31b5b39068f6a4fa1c 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.h
+++ b/include/deal.II/lac/cuda_kernels.h
@@ -18,7 +18,7 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 
 #  include <deal.II/base/cuda_size.h>

diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h
index b5bcdf69fd6e76499acd9087241142b94fd42d76..86136320b57ee6525865572a893553e9eb359815 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.templates.h
+++ b/include/deal.II/lac/cuda_kernels.templates.h
@@ -20,7 +20,7 @@
 
 #include <deal.II/lac/cuda_kernels.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 DEAL_II_NAMESPACE_OPEN
 

diff --git a/include/deal.II/lac/cuda_precondition.h b/include/deal.II/lac/cuda_precondition.h
index 401bb2e31d79230753a5545afef3fc43a472ef50..bba3c077d859b5060727c953f419fc8eb8dac18f 100644 (file)
--- a/include/deal.II/lac/cuda_precondition.h
+++ b/include/deal.II/lac/cuda_precondition.h
@@ -23,7 +23,7 @@
 
 #include <memory>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 DEAL_II_NAMESPACE_OPEN
 

diff --git a/include/deal.II/lac/cuda_solver_direct.h b/include/deal.II/lac/cuda_solver_direct.h
index 643397eee659d2288cb11c6e87c86d9dff412d0f..dff21c90af4282b1165d96575c95512f536f1f2b 100644 (file)
--- a/include/deal.II/lac/cuda_solver_direct.h
+++ b/include/deal.II/lac/cuda_solver_direct.h
@@ -18,7 +18,7 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 #  include <deal.II/base/cuda.h>
 
 #  include <deal.II/lac/cuda_sparse_matrix.h>

diff --git a/include/deal.II/lac/cuda_sparse_matrix.h b/include/deal.II/lac/cuda_sparse_matrix.h
index 98e7ae00b5023d47822a3cd49e2848471409d4e5..166c11ffe72fcda9a5a435d25a499f42fb7e35a4 100644 (file)
--- a/include/deal.II/lac/cuda_sparse_matrix.h
+++ b/include/deal.II/lac/cuda_sparse_matrix.h
@@ -22,7 +22,7 @@
 
 #include <iomanip>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 #  include <deal.II/base/cuda.h>
 
 #  include <deal.II/lac/cuda_vector.h>

diff --git a/include/deal.II/lac/precondition.h b/include/deal.II/lac/precondition.h
index 6ff12b63e1426c4880e367338c6da5d1c947ed70..c41d9698935456f1209f51bf208ad3165854fb0f 100644 (file)
--- a/include/deal.II/lac/precondition.h
+++ b/include/deal.II/lac/precondition.h
@@ -3386,7 +3386,7 @@ namespace internal
     }
 
 
-#  ifdef DEAL_II_COMPILER_CUDA_AWARE
+#  ifdef DEAL_II_WITH_CUDA
     template <typename Number>
     __global__ void
     set_initial_guess_kernel(const types::global_dof_index offset,
@@ -3425,7 +3425,7 @@ namespace internal
       const Number mean_value = vector.mean_value();
       vector.add(-mean_value);
     }
-#  endif // DEAL_II_COMPILER_CUDA_AWARE
+#  endif // DEAL_II_WITH_CUDA
 
     struct EigenvalueTracker
     {

diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h
index f887b1e66404cd84cca0121d448ff1152708fe83..ed694816b61e6c1658efe05e6a8a35e80d064140 100644 (file)
--- a/include/deal.II/matrix_free/cuda_fe_evaluation.h
+++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h
@@ -18,7 +18,7 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 #  include <deal.II/base/tensor.h>
 #  include <deal.II/base/utilities.h>

diff --git a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
index 842840951808ccae3da51be10451e1854934dfe4..e15070c511d9e68dd5e7e9baa10c4a89c2d8876a 100644 (file)
--- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
+++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
@@ -18,7 +18,7 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 #  include <deal.II/base/cuda_size.h>
 

diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h
index dd72ee4eafacd5de85b36e65565aa45572554bcb..d5bee6a61faa774b2cfff39cc1c62de1ae32e6a9 100644 (file)
--- a/include/deal.II/matrix_free/cuda_matrix_free.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.h
@@ -19,7 +19,7 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 #  include <deal.II/base/cuda_size.h>
 #  include <deal.II/base/mpi_stub.h>

diff --git a/include/deal.II/matrix_free/cuda_matrix_free.templates.h b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
index 297a0e468a6f575d6776375a465e370f53f9209f..bc5efd162e5d5efbd2a9ea14b3270829b7cf0a76 100644 (file)
--- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
@@ -21,7 +21,7 @@
 
 #include <deal.II/matrix_free/cuda_matrix_free.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 #  include <deal.II/base/cuda.h>
 #  include <deal.II/base/cuda_size.h>

diff --git a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h
index 2d55db16e90e81fd0d8032f969f0a89d7b962ebb..c423b48dd608e7fa45ae69356eea2d27d14281e0 100644 (file)
--- a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h
+++ b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h
@@ -23,7 +23,7 @@
 
 #include <deal.II/matrix_free/cuda_matrix_free.templates.h>
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 
 DEAL_II_NAMESPACE_OPEN
 

diff --git a/tests/tests.h b/tests/tests.h
index 6778eeeceddfdf15699368dbd1d06e4c50aaacdb..16185387d93096a08af533e4f20117ec858a2982 100644 (file)
--- a/tests/tests.h
+++ b/tests/tests.h
@@ -618,7 +618,7 @@ struct MPILogInitAll
 };
 
 
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
+#ifdef DEAL_II_WITH_CUDA
 // By default, all the ranks will try to access the device 0.
 // If we are running with MPI support it is better to address different graphic
 // cards for different processes even if only one node is used. The choice below