From: Daniel Arndt <arndtd@ornl.gov>
Date: Tue, 3 Sep 2019 19:35:43 +0000 (+0000)
Subject: Reorganize __syncthreads for hanging nodes
X-Git-Tag: v9.2.0-rc1~1147^2
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=19447a80d244b5479f7e4795b3c18bc9455be595;p=dealii.git

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 2bbc299a5c..b2bb935ccf 100644
--- 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 c30bb4e1ba..1560cbc764 100644
--- 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