From: Peter Munch Date: Tue, 3 Aug 2021 18:36:18 +0000 (+0200) Subject: Work on GPU X-Git-Tag: v9.4.0-rc1~1102^2~1 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=d4506cbd70ca4c64d821e3ce9ca6ffd9f41519d6;p=dealii.git Work on GPU --- diff --git a/include/deal.II/matrix_free/hanging_nodes_internal.h b/include/deal.II/matrix_free/hanging_nodes_internal.h index 1aed9095fe..bd81181863 100644 --- a/include/deal.II/matrix_free/hanging_nodes_internal.h +++ b/include/deal.II/matrix_free/hanging_nodes_internal.h @@ -913,26 +913,33 @@ namespace CUDAWrappers ((direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_x : dealii::internal::MatrixFreeFunctions::ConstraintKinds:: - unconstrained))); + unconstrained))) != + dealii::internal::MatrixFreeFunctions::ConstraintKinds::unconstrained; // Flag is true if for the given direction, the dof is constrained with // the right type and is on the correct side (left (= 0) or right (= // fe_degree)) const bool constrained_dof = ((direction == 0) && - ((constraint_mask & - dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y) ? + (((constraint_mask & + dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y) != + dealii::internal::MatrixFreeFunctions::ConstraintKinds:: + unconstrained) ? (y_idx == 0) : (y_idx == fe_degree))) || ((direction == 1) && - ((constraint_mask & - dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x) ? + (((constraint_mask & + dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x) != + dealii::internal::MatrixFreeFunctions::ConstraintKinds:: + unconstrained) ? (x_idx == 0) : (x_idx == fe_degree))); if (constrained_face && constrained_dof) { - const bool type = constraint_mask & this_type; + const bool type = (constraint_mask & this_type) != + dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained; if (type) { @@ -994,19 +1001,19 @@ namespace CUDAWrappers const unsigned int y_idx = threadIdx.y; const unsigned int z_idx = threadIdx.z; - const unsigned int this_type = + const auto this_type = (direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x : (direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y : dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_z; - const unsigned int face1_type = + const auto face1_type = (direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_y : (direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_z : dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_x; - const unsigned int face2_type = + const auto face2_type = (direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::type_z : (direction == 1) ? @@ -1015,26 +1022,25 @@ namespace CUDAWrappers // If computing in x-direction, need to match against face_y or // face_z - const unsigned int face1 = + const auto face1 = (direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_y : (direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_z : dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_x; - const unsigned int face2 = + const auto face2 = (direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_z : (direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_x : dealii::internal::MatrixFreeFunctions::ConstraintKinds::face_y; - const unsigned int edge = + const auto edge = (direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::edge_yz : (direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintKinds::edge_zx : dealii::internal::MatrixFreeFunctions::ConstraintKinds::edge_xy; - const unsigned int constrained_face = - constraint_mask & (face1 | face2 | edge); + const auto constrained_face = constraint_mask & (face1 | face2 | edge); const unsigned int interp_idx = (direction == 0) ? x_idx : (direction == 1) ? y_idx : @@ -1047,20 +1053,32 @@ namespace CUDAWrappers y_idx; Number t = 0; - const bool on_face1 = (constraint_mask & face1_type) ? + const bool on_face1 = ((constraint_mask & face1_type) != + dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained) ? (face1_idx == 0) : (face1_idx == fe_degree); - const bool on_face2 = (constraint_mask & face2_type) ? + const bool on_face2 = ((constraint_mask & face2_type) != + dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained) ? (face2_idx == 0) : (face2_idx == fe_degree); const bool constrained_dof = - (((constraint_mask & face1) && on_face1) || - ((constraint_mask & face2) && on_face2) || - ((constraint_mask & edge) && on_face1 && on_face2)); + ((((constraint_mask & face1) != dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained) && + on_face1) || + (((constraint_mask & face2) != dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained) && + on_face2) || + (((constraint_mask & edge) != dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained) && + on_face1 && on_face2)); if (constrained_face && constrained_dof) { - const bool type = constraint_mask & this_type; + const bool type = (constraint_mask & this_type) != + dealii::internal::MatrixFreeFunctions:: + ConstraintKinds::unconstrained; if (type) { for (unsigned int i = 0; i <= fe_degree; ++i)