From: Bruno Turcksin <bruno.turcksin@gmail.com>
Date: Mon, 2 Sep 2019 02:39:37 +0000 (+0000)
Subject: Fix a race condition when computing hanging with the CUDA matrix-free framework
X-Git-Tag: v9.2.0-rc1~1148^2~2
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=611ba85199178f9fce599a513932b2c5ed3ef029;p=dealii.git

Fix a race condition when computing hanging with the CUDA matrix-free framework
---

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 eb7d34b8b9..c30bb4e1ba 100644
--- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
+++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
@@ -177,11 +177,6 @@ namespace CUDAWrappers
       , fe_degree(fe_degree)
       , dof_handler(dof_handler)
     {
-      AssertThrow(
-        (dim == 3) || ((fe_degree % 2) == 1),
-        ExcMessage(
-          "This function is not implemented when dim = 2 and fe_degree is even."));
-
       // Set up line-to-cell mapping for edge constraints (only if dim = 3)
       setup_line_to_cell();
 
@@ -787,6 +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;
@@ -901,6 +897,7 @@ 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;