From: Daniel Arndt 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=refs%2Fpull%2F8693%2Fhead;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(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(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(x_idx, y_idx, z_idx)] = t; + + __syncthreads(); } } // namespace internal