From: Peter Munch Date: Fri, 27 Sep 2019 05:37:40 +0000 (+0200) Subject: Update the GPU matrix-free documentation (2) X-Git-Tag: v9.2.0-rc1~1008^2 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=refs%2Fpull%2F8872%2Fhead;p=dealii.git Update the GPU matrix-free documentation (2) --- 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]; };