]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Update the GPU matrix-free documentation (2) 8872/head
authorPeter Munch <peterrmuench@gmail.com>
Fri, 27 Sep 2019 05:37:40 +0000 (07:37 +0200)
committerPeter Munch <peterrmuench@gmail.com>
Fri, 27 Sep 2019 05:37:40 +0000 (07:37 +0200)
include/deal.II/matrix_free/cuda_fe_evaluation.h
include/deal.II/matrix_free/cuda_matrix_free.h

index ba7f3dc0d8af2544de5d2d821f63e9ec635dc175..a36434fecc591fb32ff21eed68fd9a7f4aaf16a5 100644 (file)
@@ -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);
 
index 8b741dc90cc7b8cd882e71a308691edb198a02c0..e749c854e2e79a6d52d77068ca3eabd98cb66571 100644 (file)
@@ -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];
   };
 

In the beginning the Universe was created. This has made a lot of people very angry and has been widely regarded as a bad move.

Douglas Adams


Typeset in Trocchi and Trocchi Bold Sans Serif.