From: Peter Munch <peterrmuench@gmail.com>
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=781c4ac6b706b88d2a26d222b9cb37a045de03f2;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<dim, Number>::Data;
-    static constexpr unsigned int dimension    = dim;
+
+    /**
+     * An alias to kernel specific information.
+     */
+    using data_type = typename MatrixFree<dim, Number>::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 <int dim, typename Number>
   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];
   };