Reorganize __syncthreads for hanging nodes

diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h
index 2bbc299a5cbec693cb06e5868888f25fed8c54cc..b2bb935ccfa457352bee0c4447edd6e00d16a1e7 100644 (file)
--- a/include/deal.II/matrix_free/cuda_fe_evaluation.h
+++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h
@@ -258,10 +258,10 @@ namespace CUDAWrappers
     // Use the read-only data cache.
     values[idx] = __ldg(&src[src_idx]);
 
+    __syncthreads();
+
     internal::resolve_hanging_nodes<dim, fe_degree, false>(constraint_mask,
                                                            values);
-
-    __syncthreads();
   }
 
 

diff --git a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
index c30bb4e1ba7bde3378b2dacfadb1590be4dd0a2e..1560cbc7640b385e415b8bfd013b7b70c4dcf9ae 100644 (file)
--- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
+++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
@@ -782,7 +782,7 @@ namespace CUDAWrappers
           ((direction == 1) && ((constraint_mask & internal::constr_type_x) ?
                                   (x_idx == 0) :
                                   (x_idx == fe_degree)));
-        __syncthreads();
+
         if (constrained_face && constrained_dof)
           {
             const bool type = constraint_mask & this_type;
@@ -829,9 +829,10 @@ namespace CUDAWrappers
         // The synchronization is done for all the threads in one block with
         // each block being assigned to one element.
         __syncthreads();
-
         if (constrained_face && constrained_dof)
           values[index2<fe_degree + 1>(x_idx, y_idx)] = t;
+
+        __syncthreads();
       }
 
 
@@ -897,7 +898,6 @@ namespace CUDAWrappers
            ((constraint_mask & face2) && on_face2) ||
            ((constraint_mask & edge) && on_face1 && on_face2));
 
-        __syncthreads();
         if (constrained_face && constrained_dof)
           {
             const bool type = constraint_mask & this_type;
@@ -946,10 +946,14 @@ namespace CUDAWrappers
               }
           }
 
+        // The synchronization is done for all the threads in one block with
+        // each block being assigned to one element.
         __syncthreads();
 
         if (constrained_face && constrained_dof)
           values[index3<fe_degree + 1>(x_idx, y_idx, z_idx)] = t;
+
+        __syncthreads();
       }
     } // namespace internal