]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Fix a race condition when computing hanging with the CUDA matrix-free framework
authorBruno Turcksin <bruno.turcksin@gmail.com>
Mon, 2 Sep 2019 02:39:37 +0000 (02:39 +0000)
committerBruno Turcksin <bruno.turcksin@gmail.com>
Mon, 2 Sep 2019 02:39:37 +0000 (02:39 +0000)
include/deal.II/matrix_free/cuda_hanging_nodes_internal.h

index eb7d34b8b9c9ddbf29c8c790392c851a13d25275..c30bb4e1ba7bde3378b2dacfadb1590be4dd0a2e 100644 (file)
@@ -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;

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.