From 19447a80d244b5479f7e4795b3c18bc9455be595 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Tue, 3 Sep 2019 19:35:43 +0000 Subject: [PATCH] Reorganize __syncthreads for hanging nodes --- include/deal.II/matrix_free/cuda_fe_evaluation.h | 4 ++-- .../deal.II/matrix_free/cuda_hanging_nodes_internal.h | 10 +++++++--- 2 files changed, 9 insertions(+), 5 deletions(-) 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 -- 2.39.5