From: Daniel Arndt Date: Wed, 19 Sep 2018 09:18:28 +0000 (+0200) Subject: Make CUDAWrappers::MatrixFree compile with 64-bit indices X-Git-Tag: v9.1.0-rc1~703^2 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=refs%2Fpull%2F7213%2Fhead;p=dealii.git 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 b097cfc5a7..9dfb660bff 100644 --- 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 946af3f1c8..17e50cab20 100644 --- 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 void - setup_constraints(std::vector &dof_indices, - const CellIterator & cell, - unsigned int & mask) const; + setup_constraints(std::vector &dof_indices, + const CellIterator & cell, + unsigned int & mask) const; private: /** @@ -271,9 +271,10 @@ namespace CUDAWrappers template template void - HangingNodes::setup_constraints(std::vector &dof_indices, - const CellIterator & cell, - unsigned int & mask) const + HangingNodes::setup_constraints( + std::vector &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 neighbor_dofs(dofs_per_face); - std::vector lex_face_mapping = + const auto lex_face_mapping = FETools::lexicographic_to_hierarchic_numbering( FE_Q(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 3ebe12d846..5d4b5fb79a 100644 --- 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 n_cells; + /** * Vector of pointers to the quadrature points associated to the cells of * each color. */ std::vector q_points; + /** * Map the position in the local vector to the position in the global * vector. */ - std::vector local_to_global; + std::vector local_to_global; + /** * Vector of pointer to the inverse Jacobian associated to the cells of each * color. */ std::vector inv_jacobian; + /** * Vector of pointer to the Jacobian time the weights associated to the * cells of each color. */ std::vector 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 constraint_mask; + /** * Grid dimensions associated to the different colors. The grid dimensions * are used to launch the CUDA kernels. */ std::vector 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 d8cc558025..8655befcf5 100644 --- 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 *data; // Host data - std::vector local_to_global_host; - std::vector> q_points_host; - std::vector JxW_host; - std::vector inv_jacobian_host; - std::vector constraint_mask_host; + std::vector local_to_global_host; + std::vector> q_points_host; + std::vector JxW_host; + std::vector inv_jacobian_host; + std::vector constraint_mask_host; // Local buffer std::vector local_dof_indices; FEValues fe_values; // Convert the default dof numbering to a lexicographic one - const std::vector &lexicographic_inv; - std::vector 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 hanging_nodes; + const std::vector & lexicographic_inv; + std::vector 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 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);