Make CUDAWrappers::MatrixFree compile with 64-bit indices

diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h
index b097cfc5a7d031a39d4f325a34b387d9c60f7958..9dfb660bffdb85be9ee0fcb6a9b196b6cc4a3128 100644 (file)
--- a/include/deal.II/matrix_free/cuda_fe_evaluation.h
+++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h
@@ -168,9 +168,9 @@ namespace CUDAWrappers
     apply_quad_point_operations(const functor &func);
 
   private:
-    unsigned int *local_to_global;
-    unsigned int  n_cells;
-    unsigned int  padding_length;
+    types::global_dof_index *local_to_global;
+    unsigned int             n_cells;
+    unsigned int             padding_length;
 
     const unsigned int constraint_mask;
 
@@ -225,7 +225,7 @@ namespace CUDAWrappers
       (dim > 1 ? threadIdx.y : 0) * n_q_points_1d +
       (dim > 2 ? threadIdx.z : 0) * n_q_points_1d * n_q_points_1d;
 
-    const unsigned int src_idx = local_to_global[idx];
+    const types::global_dof_index src_idx = local_to_global[idx];
     // Use the read-only data cache.
     values[idx] = __ldg(&src[src_idx]);
 
@@ -257,7 +257,7 @@ namespace CUDAWrappers
       (threadIdx.x % n_q_points_1d) +
       (dim > 1 ? threadIdx.y : 0) * n_q_points_1d +
       (dim > 2 ? threadIdx.z : 0) * n_q_points_1d * n_q_points_1d;
-    const unsigned int destination_idx = local_to_global[idx];
+    const types::global_dof_index destination_idx = local_to_global[idx];
 
     dst[destination_idx] += values[idx];
   }

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 946af3f1c8a626e3374a08e9343b8d4dddd4c241..17e50cab200b15b6a7fc5f831bfac922e2f73d5b 100644 (file)
--- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
+++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
@@ -58,9 +58,9 @@ namespace CUDAWrappers
        */
       template <typename CellIterator>
       void
-      setup_constraints(std::vector<unsigned int> &dof_indices,
-                        const CellIterator &       cell,
-                        unsigned int &             mask) const;
+      setup_constraints(std::vector<types::global_dof_index> &dof_indices,
+                        const CellIterator &                  cell,
+                        unsigned int &                        mask) const;
 
     private:
       /**
@@ -271,9 +271,10 @@ namespace CUDAWrappers
     template <int dim>
     template <typename CellIterator>
     void
-    HangingNodes<dim>::setup_constraints(std::vector<unsigned int> &dof_indices,
-                                         const CellIterator &       cell,
-                                         unsigned int &             mask) const
+    HangingNodes<dim>::setup_constraints(
+      std::vector<types::global_dof_index> &dof_indices,
+      const CellIterator &                  cell,
+      unsigned int &                        mask) const
     {
       mask                         = 0;
       const unsigned int n_dofs_1d = fe_degree + 1;
@@ -282,7 +283,7 @@ namespace CUDAWrappers
 
       std::vector<types::global_dof_index> neighbor_dofs(dofs_per_face);
 
-      std::vector<types::global_dof_index> lex_face_mapping =
+      const auto lex_face_mapping =
         FETools::lexicographic_to_hierarchic_numbering<dim - 1>(
           FE_Q<dim - 1>(fe_degree));
 

diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h
index 3ebe12d846bae66bb5dca13ef4fe0db45bd0b4e3..5d4b5fb79aeb3ee340f483178c916df84ae26dff 100644 (file)
--- a/include/deal.II/matrix_free/cuda_matrix_free.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.h
@@ -131,14 +131,14 @@ namespace CUDAWrappers
      */
     struct Data
     {
-      point_type *  q_points;
-      unsigned int *local_to_global;
-      Number *      inv_jacobian;
-      Number *      JxW;
-      unsigned int  n_cells;
-      unsigned int  padding_length;
-      unsigned int  row_start;
-      unsigned int *constraint_mask;
+      point_type *             q_points;
+      types::global_dof_index *local_to_global;
+      Number *                 inv_jacobian;
+      Number *                 JxW;
+      unsigned int             n_cells;
+      unsigned int             padding_length;
+      unsigned int             row_start;
+      unsigned int *           constraint_mask;
     };
 
     /**
@@ -214,59 +214,73 @@ namespace CUDAWrappers
      * over cells.
      */
     ParallelizationScheme parallelization_scheme;
+
     /**
      * Degree of the finite element used.
      */
     unsigned int fe_degree;
+
     /**
      * Number of degrees of freedom per cell.
      */
     unsigned int dofs_per_cell;
+
     /**
      * Number of constrained degrees of freedom.
      */
     unsigned int n_constrained_dofs;
+
     /**
      * Number of quadrature points per cells.
      */
     unsigned int q_points_per_cell;
+
     /**
      * Number of colors produced by the graph coloring algorithm.
      */
     unsigned int n_colors;
+
     /**
      * Number of cells in each color.
      */
     std::vector<unsigned int> n_cells;
+
     /**
      * Vector of pointers to the quadrature points associated to the cells of
      * each color.
      */
     std::vector<point_type *> q_points;
+
     /**
      * Map the position in the local vector to the position in the global
      * vector.
      */
-    std::vector<unsigned int *> local_to_global;
+    std::vector<types::global_dof_index *> local_to_global;
+
     /**
      * Vector of pointer to the inverse Jacobian associated to the cells of each
      * color.
      */
     std::vector<Number *> inv_jacobian;
+
     /**
      * Vector of pointer to the Jacobian time the weights associated to the
      * cells of each color.
      */
     std::vector<Number *> JxW;
 
-    // Constraints
-    unsigned int *              constrained_dofs;
+    // Pointer to the constrained degrees of freedom.
+    types::global_dof_index *constrained_dofs;
+
+    // Mask deciding where constraints are set on a given cell.
     std::vector<unsigned int *> constraint_mask;
+
     /**
      * Grid dimensions associated to the different colors. The grid dimensions
      * are used to launch the CUDA kernels.
      */
     std::vector<dim3> grid_dim;
+
     /**
      * Block dimensions associated to the different colors. The block dimensions
      * are used to launch the CUDA kernels.

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 d8cc558025d5e4c65f155d45b734c14dfe6e5cdb..8655befcf52192572d3bf30c2fbf948d03021c71 100644 (file)
--- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
@@ -145,23 +145,23 @@ namespace CUDAWrappers
     private:
       MatrixFree<dim, Number> *data;
       // Host data
-      std::vector<unsigned int> local_to_global_host;
-      std::vector<Point<dim>>   q_points_host;
-      std::vector<Number>       JxW_host;
-      std::vector<Number>       inv_jacobian_host;
-      std::vector<unsigned int> constraint_mask_host;
+      std::vector<types::global_dof_index> local_to_global_host;
+      std::vector<Point<dim>>              q_points_host;
+      std::vector<Number>                  JxW_host;
+      std::vector<Number>                  inv_jacobian_host;
+      std::vector<unsigned int>            constraint_mask_host;
       // Local buffer
       std::vector<types::global_dof_index> local_dof_indices;
       FEValues<dim>                        fe_values;
       // Convert the default dof numbering to a lexicographic one
-      const std::vector<unsigned int> &lexicographic_inv;
-      std::vector<unsigned int>        lexicographic_dof_indices;
-      const unsigned int               fe_degree;
-      const unsigned int               dofs_per_cell;
-      const unsigned int               q_points_per_cell;
-      const UpdateFlags &              update_flags;
-      const unsigned int               padding_length;
-      HangingNodes<dim>                hanging_nodes;
+      const std::vector<unsigned int> &    lexicographic_inv;
+      std::vector<types::global_dof_index> lexicographic_dof_indices;
+      const unsigned int                   fe_degree;
+      const unsigned int                   dofs_per_cell;
+      const unsigned int                   q_points_per_cell;
+      const UpdateFlags &                  update_flags;
+      const unsigned int                   padding_length;
+      HangingNodes<dim>                    hanging_nodes;
     };
 
 
@@ -289,7 +289,7 @@ namespace CUDAWrappers
 
       memcpy(&local_to_global_host[cell_id * padding_length],
              lexicographic_dof_indices.data(),
-             dofs_per_cell * sizeof(unsigned int));
+             dofs_per_cell * sizeof(types::global_dof_index));
 
       fe_values.reinit(cell);