]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Work on GPU
authorPeter Munch <peterrmuench@gmail.com>
Tue, 3 Aug 2021 18:36:18 +0000 (20:36 +0200)
committerPeter Munch <peterrmuench@gmail.com>
Tue, 3 Aug 2021 18:36:18 +0000 (20:36 +0200)
include/deal.II/matrix_free/hanging_nodes_internal.h

index 1aed9095fe0d10541010a3c7ddcc30c94d617123..bd8118186331dcdb58cc6db452cb9bf52a3b592d 100644 (file)
@@ -913,26 +913,33 @@ namespace CUDAWrappers
           ((direction == 1) ?
              dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_x :
              dealii::internal::MatrixFreeFunctions::ConstraintKinds::
-               unconstrained)));
+               unconstrained))) !=
+        dealii::internal::MatrixFreeFunctions::ConstraintKinds::unconstrained;
 
       // Flag is true if for the given direction, the dof is constrained with
       // the right type and is on the correct side (left (= 0) or right (=
       // fe_degree))
       const bool constrained_dof =
         ((direction == 0) &&
-         ((constraint_mask &
-           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y) ?
+         (((constraint_mask &
+            dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y) !=
+           dealii::internal::MatrixFreeFunctions::ConstraintKinds::
+             unconstrained) ?
             (y_idx == 0) :
             (y_idx == fe_degree))) ||
         ((direction == 1) &&
-         ((constraint_mask &
-           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x) ?
+         (((constraint_mask &
+            dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x) !=
+           dealii::internal::MatrixFreeFunctions::ConstraintKinds::
+             unconstrained) ?
             (x_idx == 0) :
             (x_idx == fe_degree)));
 
       if (constrained_face && constrained_dof)
         {
-          const bool type = constraint_mask & this_type;
+          const bool type = (constraint_mask & this_type) !=
+                            dealii::internal::MatrixFreeFunctions::
+                              ConstraintKinds::unconstrained;
 
           if (type)
             {
@@ -994,19 +1001,19 @@ namespace CUDAWrappers
       const unsigned int y_idx = threadIdx.y;
       const unsigned int z_idx = threadIdx.z;
 
-      const unsigned int this_type =
+      const auto this_type =
         (direction == 0) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x :
         (direction == 1) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y :
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_z;
-      const unsigned int face1_type =
+      const auto face1_type =
         (direction == 0) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y :
         (direction == 1) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_z :
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x;
-      const unsigned int face2_type =
+      const auto face2_type =
         (direction == 0) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_z :
         (direction == 1) ?
@@ -1015,26 +1022,25 @@ namespace CUDAWrappers
 
       // If computing in x-direction, need to match against face_y or
       // face_z
-      const unsigned int face1 =
+      const auto face1 =
         (direction == 0) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_y :
         (direction == 1) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_z :
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_x;
-      const unsigned int face2 =
+      const auto face2 =
         (direction == 0) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_z :
         (direction == 1) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_x :
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_y;
-      const unsigned int edge =
+      const auto edge =
         (direction == 0) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::edge_yz :
         (direction == 1) ?
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::edge_zx :
           dealii::internal::MatrixFreeFunctions::ConstraintKinds::edge_xy;
-      const unsigned int constrained_face =
-        constraint_mask & (face1 | face2 | edge);
+      const auto constrained_face = constraint_mask & (face1 | face2 | edge);
 
       const unsigned int interp_idx = (direction == 0) ? x_idx :
                                       (direction == 1) ? y_idx :
@@ -1047,20 +1053,32 @@ namespace CUDAWrappers
                                                          y_idx;
 
       Number     t        = 0;
-      const bool on_face1 = (constraint_mask & face1_type) ?
+      const bool on_face1 = ((constraint_mask & face1_type) !=
+                             dealii::internal::MatrixFreeFunctions::
+                               ConstraintKinds::unconstrained) ?
                               (face1_idx == 0) :
                               (face1_idx == fe_degree);
-      const bool on_face2 = (constraint_mask & face2_type) ?
+      const bool on_face2 = ((constraint_mask & face2_type) !=
+                             dealii::internal::MatrixFreeFunctions::
+                               ConstraintKinds::unconstrained) ?
                               (face2_idx == 0) :
                               (face2_idx == fe_degree);
       const bool constrained_dof =
-        (((constraint_mask & face1) && on_face1) ||
-         ((constraint_mask & face2) && on_face2) ||
-         ((constraint_mask & edge) && on_face1 && on_face2));
+        ((((constraint_mask & face1) != dealii::internal::MatrixFreeFunctions::
+                                          ConstraintKinds::unconstrained) &&
+          on_face1) ||
+         (((constraint_mask & face2) != dealii::internal::MatrixFreeFunctions::
+                                          ConstraintKinds::unconstrained) &&
+          on_face2) ||
+         (((constraint_mask & edge) != dealii::internal::MatrixFreeFunctions::
+                                         ConstraintKinds::unconstrained) &&
+          on_face1 && on_face2));
 
       if (constrained_face && constrained_dof)
         {
-          const bool type = constraint_mask & this_type;
+          const bool type = (constraint_mask & this_type) !=
+                            dealii::internal::MatrixFreeFunctions::
+                              ConstraintKinds::unconstrained;
           if (type)
             {
               for (unsigned int i = 0; i <= fe_degree; ++i)

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.