From 781c4ac6b706b88d2a26d222b9cb37a045de03f2 Mon Sep 17 00:00:00 2001 From: Peter Munch Date: Fri, 27 Sep 2019 07:37:40 +0200 Subject: [PATCH] Update the GPU matrix-free documentation (2) --- .../deal.II/matrix_free/cuda_fe_evaluation.h | 33 +++++++++++++++++-- .../deal.II/matrix_free/cuda_matrix_free.h | 17 +++++++++- 2 files changed, 46 insertions(+), 4 deletions(-) diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h index ba7f3dc0d8..a36434fecc 100644 --- a/include/deal.II/matrix_free/cuda_fe_evaluation.h +++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h @@ -92,13 +92,40 @@ namespace CUDAWrappers class FEEvaluation { public: - using value_type = Number; + /** + * An alias for scalar quantities. + */ + using value_type = Number; + + /** + * An alias for vectorial quantities. + */ using gradient_type = Tensor<1, dim, Number>; - using data_type = typename MatrixFree::Data; - static constexpr unsigned int dimension = dim; + + /** + * An alias to kernel specific information. + */ + using data_type = typename MatrixFree::Data; + + /** + * Dimension. + */ + static constexpr unsigned int dimension = dim; + + /** + * Number of components. + */ static constexpr unsigned int n_components = n_components_; + + /** + * Number of quadrature points per cell. + */ static constexpr unsigned int n_q_points = Utilities::pow(n_q_points_1d, dim); + + /** + * Number of tensor degrees of freedoms per cell. + */ static constexpr unsigned int tensor_dofs_per_cell = Utilities::pow(fe_degree + 1, dim); diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h index 8b741dc90c..e749c854e2 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.h @@ -562,10 +562,16 @@ namespace CUDAWrappers // TODO find a better place to put these things - // Structure to pass the shared memory into a general user function. + + /** + * Structure to pass the shared memory into a general user function. + */ template struct SharedData { + /** + * Constructor. + */ __device__ SharedData(Number *vd, Number *gq[dim]) : values(vd) @@ -574,7 +580,16 @@ namespace CUDAWrappers gradients[d] = gq[d]; } + /** + * Shared memory for dof and quad values. + */ Number *values; + + /** + * Shared memory for computed gradients in reference coordinate system. + * The gradient in each direction is saved in a struct-of-array + * format, i.e. first, all gradients in the x-direction come... + */ Number *gradients[dim]; }; -- 2.39.5