From feb1ff2d09ce2efd63e6410458e540ec254bae58 Mon Sep 17 00:00:00 2001 From: Peter Munch Date: Sat, 31 Jul 2021 11:28:26 +0200 Subject: [PATCH] Swith from unsigned int to ConstraintTypes --- .../deal.II/matrix_free/cuda_fe_evaluation.h | 3 +- .../deal.II/matrix_free/cuda_matrix_free.h | 10 +- .../matrix_free/cuda_matrix_free.templates.h | 16 +- include/deal.II/matrix_free/dof_info.h | 3 +- .../deal.II/matrix_free/dof_info.templates.h | 2 +- .../deal.II/matrix_free/evaluation_kernels.h | 33 ++-- .../matrix_free/evaluation_template_factory.h | 19 ++- .../evaluation_template_factory.templates.h | 16 +- include/deal.II/matrix_free/fe_evaluation.h | 11 +- .../matrix_free/hanging_nodes_internal.h | 157 ++++++++++++------ include/deal.II/matrix_free/tools.h | 4 +- source/matrix_free/dof_info.cc | 5 +- tests/matrix_free/hanging_node_kernels_01.cc | 92 +++++----- 13 files changed, 232 insertions(+), 139 deletions(-) diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h index 01928dc8c7..06efd8e5ba 100644 --- a/include/deal.II/matrix_free/cuda_fe_evaluation.h +++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h @@ -238,7 +238,8 @@ namespace CUDAWrappers unsigned int padding_length; const unsigned int mf_object_id; - const unsigned int constraint_mask; + const dealii::internal::MatrixFreeFunctions::ConstraintTypes + constraint_mask; const bool use_coloring; diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h index 3181bb014b..3557fd4218 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.h @@ -38,6 +38,8 @@ # include # include +# include + DEAL_II_NAMESPACE_OPEN @@ -210,7 +212,7 @@ namespace CUDAWrappers /** * Mask deciding where constraints are set on a given cell. */ - unsigned int *constraint_mask; + dealii::internal::MatrixFreeFunctions::ConstraintTypes *constraint_mask; /** * If true, use graph coloring has been used and we can simply add into @@ -589,7 +591,8 @@ namespace CUDAWrappers /** * Mask deciding where constraints are set on a given cell. */ - std::vector constraint_mask; + std::vector + constraint_mask; /** * Grid dimensions associated to the different colors. The grid dimensions @@ -815,7 +818,8 @@ namespace CUDAWrappers /** * Mask deciding where constraints are set on a given cell. */ - std::vector constraint_mask; + std::vector + constraint_mask; /** * If true, use graph coloring has been used and we can simply add into diff --git a/include/deal.II/matrix_free/cuda_matrix_free.templates.h b/include/deal.II/matrix_free/cuda_matrix_free.templates.h index 3252ccabd6..1f6750010c 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h @@ -31,7 +31,6 @@ # include # include -# include # include # include @@ -230,7 +229,8 @@ namespace CUDAWrappers std::vector> q_points_host; std::vector JxW_host; std::vector inv_jacobian_host; - std::vector constraint_mask_host; + std::vector + constraint_mask_host; // Local buffer std::vector local_dof_indices; FEValues fe_values; @@ -378,7 +378,8 @@ namespace CUDAWrappers for (unsigned int i = 0; i < dofs_per_cell; ++i) lexicographic_dof_indices[i] = local_dof_indices[lexicographic_inv[i]]; - const ArrayView cell_id_view(constraint_mask_host[cell_id]); + const ArrayView + cell_id_view(constraint_mask_host[cell_id]); hanging_nodes.setup_constraints(cell, partitioner, @@ -493,10 +494,11 @@ namespace CUDAWrappers n_cells * dim * dim * padding_length); } - alloc_and_copy(&data->constraint_mask[color], - ArrayView(constraint_mask_host.data(), - constraint_mask_host.size()), - n_cells); + alloc_and_copy( + &data->constraint_mask[color], + ArrayView( + constraint_mask_host.data(), constraint_mask_host.size()), + n_cells); } diff --git a/include/deal.II/matrix_free/dof_info.h b/include/deal.II/matrix_free/dof_info.h index b4b2aaceaf..a465e8537b 100644 --- a/include/deal.II/matrix_free/dof_info.h +++ b/include/deal.II/matrix_free/dof_info.h @@ -30,6 +30,7 @@ #include #include +#include #include #include #include @@ -527,7 +528,7 @@ namespace internal * Masks indicating for each cell and component if the optimized * hanging-node constraint is applicable and if yes which type. */ - std::vector hanging_node_constraint_masks; + std::vector hanging_node_constraint_masks; /** * This variable describes the position of constraints in terms of the diff --git a/include/deal.II/matrix_free/dof_info.templates.h b/include/deal.II/matrix_free/dof_info.templates.h index 1ea3d41022..6e019507db 100644 --- a/include/deal.II/matrix_free/dof_info.templates.h +++ b/include/deal.II/matrix_free/dof_info.templates.h @@ -283,7 +283,7 @@ namespace internal const TriaIterator> &cell, std::vector & dof_indices) { - const ArrayView mask_view( + const ArrayView mask_view( hanging_node_constraint_masks.data() + cell_number * cell->get_fe().n_components(), cell->get_fe().n_components()); diff --git a/include/deal.II/matrix_free/evaluation_kernels.h b/include/deal.II/matrix_free/evaluation_kernels.h index 2d87ded745..cbbabd1bf6 100644 --- a/include/deal.II/matrix_free/evaluation_kernels.h +++ b/include/deal.II/matrix_free/evaluation_kernels.h @@ -4428,14 +4428,15 @@ namespace internal { template static bool - run(const unsigned int n_desired_components, + run(const unsigned int n_desired_components, const FEEvaluationBaseData & fe_eval, - const bool transpose, - const std::array &c_mask, - Number * values) + Number> &fe_eval, + const bool transpose, + const std::array + & c_mask, + Number *values) { if (transpose) run_internal(n_desired_components, @@ -4649,14 +4650,14 @@ namespace internal template static void - run_internal( - const unsigned int n_desired_components, - const FEEvaluationBaseData & fe_eval, - const std::array &constraint_mask, - Number * values) + run_internal(const unsigned int n_desired_components, + const FEEvaluationBaseData &fe_eval, + const std::array & constraint_mask, + Number * values) { const Number *weights = fe_eval.get_shape_info() .data.front() @@ -4680,11 +4681,11 @@ namespace internal { for (unsigned int v = 0; v < Number::size(); ++v) { - if (constraint_mask[v] == 0) - continue; - const auto mask = constraint_mask[v]; + if (mask == MatrixFreeFunctions::ConstraintTypes::unconstrained) + continue; + if (dim == 2) // 2D: only faces { // direction 0: diff --git a/include/deal.II/matrix_free/evaluation_template_factory.h b/include/deal.II/matrix_free/evaluation_template_factory.h index e344762058..ca38ec24c5 100644 --- a/include/deal.II/matrix_free/evaluation_template_factory.h +++ b/include/deal.II/matrix_free/evaluation_template_factory.h @@ -22,6 +22,7 @@ #include #include +#include #include @@ -199,10 +200,11 @@ namespace internal apply(const unsigned int n_components, const unsigned int fe_degree, const FEEvaluationBaseData - & fe_eval, - const bool transpose, - const std::array &c_mask, - VectorizedArrayType * values); + & fe_eval, + const bool transpose, + const std::array &c_mask, + VectorizedArrayType * values); /** * For faces. @@ -213,10 +215,11 @@ namespace internal apply(const unsigned int n_components, const unsigned int fe_degree, const FEEvaluationBaseData - & fe_eval, - const bool transpose, - const std::array &c_mask, - VectorizedArrayType * values); + & fe_eval, + const bool transpose, + const std::array &c_mask, + VectorizedArrayType * values); }; } // end of namespace internal diff --git a/include/deal.II/matrix_free/evaluation_template_factory.templates.h b/include/deal.II/matrix_free/evaluation_template_factory.templates.h index 73ab1e7fa0..ce74b34fb8 100644 --- a/include/deal.II/matrix_free/evaluation_template_factory.templates.h +++ b/include/deal.II/matrix_free/evaluation_template_factory.templates.h @@ -410,10 +410,11 @@ namespace internal const unsigned int n_components, const unsigned int fe_degree, const FEEvaluationBaseData - & fe_eval, - const bool transpose, - const std::array &c_mask, - VectorizedArrayType * values) + & fe_eval, + const bool transpose, + const std::array &c_mask, + VectorizedArrayType * values) { instantiation_helper_run< 1, @@ -435,9 +436,10 @@ namespace internal const unsigned int n_components, const unsigned int fe_degree, const FEEvaluationBaseData &fe_eval, - const bool transpose, - const std::array &c_mask, - VectorizedArrayType * values) + const bool transpose, + const std::array &c_mask, + VectorizedArrayType * values) { instantiation_helper_run< 1, diff --git a/include/deal.II/matrix_free/fe_evaluation.h b/include/deal.II/matrix_free/fe_evaluation.h index 93aeaaef9f..871640c8e1 100644 --- a/include/deal.II/matrix_free/fe_evaluation.h +++ b/include/deal.II/matrix_free/fe_evaluation.h @@ -5312,8 +5312,9 @@ FEEvaluationBase:: this->dof_info ->n_vectorization_lanes_filled[this->dof_access_index][this->cell]; - constexpr unsigned int n_lanes = VectorizedArrayType::size(); - std::array constraint_mask; + constexpr unsigned int n_lanes = VectorizedArrayType::size(); + std::array + constraint_mask; bool hn_available = false; @@ -5346,14 +5347,16 @@ FEEvaluationBase:: this->dof_info->hanging_node_constraint_masks[cell_dof_index]; constraint_mask[v] = mask; - hn_available |= (mask != 0); + hn_available |= + (mask != internal::MatrixFreeFunctions::ConstraintTypes::unconstrained); } if (hn_available == false) return; // no hanging node on cell batch -> nothing to do for (unsigned int v = n_vectorization_actual; v < n_lanes; ++v) - constraint_mask[v] = 0; + constraint_mask[v] = + internal::MatrixFreeFunctions::ConstraintTypes::unconstrained; internal::FEEvaluationHangingNodesFactory:: apply(n_components, diff --git a/include/deal.II/matrix_free/hanging_nodes_internal.h b/include/deal.II/matrix_free/hanging_nodes_internal.h index 97421e7f3a..d5c1a3aa47 100644 --- a/include/deal.II/matrix_free/hanging_nodes_internal.h +++ b/include/deal.II/matrix_free/hanging_nodes_internal.h @@ -33,6 +33,93 @@ namespace internal { namespace MatrixFreeFunctions { + /** + * Here is the system for how we store constraint types in a binary mask. + * This is not a complete contradiction-free system, i.e., there are + * invalid states that we just assume that we never get. + * + * If the mask is zero, there are no constraints. Then, there are three + * different fields with one bit per dimension. The first field determines + * the type, or the position of an element along each direction. The + * second field determines if there is a constrained face with that + * direction as normal. The last field determines if there is a + * constrained edge of a given pair of coordinate planes, but where + * neither of the corresponding faces are constrained (only valid in 3D). + * + * The element is placed in the 'first position' along *-axis. These also + * determine which face is constrained. For example, in 2D, if + * face_x and type are set, then x = 0 is constrained. + */ + enum ConstraintTypes : unsigned short + { + unconstrained = 0, + + type_x = 1 << 0, + type_y = 1 << 1, + type_z = 1 << 2, + + // Element has as a constraint at * = 0 or * = fe_degree face + face_x = 1 << 3, + face_y = 1 << 4, + face_z = 1 << 5, + + // Element has as a constraint at * = 0 or * = fe_degree edge + edge_xy = 1 << 6, + edge_yz = 1 << 7, + edge_zx = 1 << 8 + }; + + + + /** + * Global operator which returns an object in which all bits are set which + * are either set in the first or the second argument. This operator exists + * since if it did not then the result of the bit-or operator | + * would be an integer which would in turn trigger a compiler warning when + * we tried to assign it to an object of type UpdateFlags. + */ + inline ConstraintTypes + operator|(const ConstraintTypes f1, const ConstraintTypes f2) + { + return static_cast(static_cast(f1) | + static_cast(f2)); + } + + + + /** + * Global operator which sets the bits from the second argument also in the + * first one. + */ + inline ConstraintTypes & + operator|=(ConstraintTypes &f1, const ConstraintTypes f2) + { + f1 = f1 | f2; + return f1; + } + + + +#ifdef DEAL_II_COMPILER_CUDA_AWARE + __device__ inline ConstraintTypes + operator|(const ConstraintTypes f1, const ConstraintTypes f2) + { + return static_cast(static_cast(f1) | + static_cast(f2)); + } + + + + __device__ inline ConstraintTypes & + operator|=(ConstraintTypes &f1, const ConstraintTypes f2) + { + f1 = f1 | f2; + return f1; + } +#endif + + + /** * This class creates the mask used in the treatment of hanging nodes in * CUDAWrappers::MatrixFree. @@ -60,7 +147,7 @@ namespace internal const std::shared_ptr &partitioner, const std::vector & lexicographic_mapping, std::vector &dof_indices, - const ArrayView & mask) const; + const ArrayView & mask) const; private: /** @@ -94,41 +181,6 @@ namespace internal line_to_cells; }; - /** - * Here is the system for how we store constraint types in a binary mask. - * This is not a complete contradiction-free system, i.e., there are - * invalid states that we just assume that we never get. - * - * If the mask is zero, there are no constraints. Then, there are three - * different fields with one bit per dimension. The first field determines - * the type, or the position of an element along each direction. The - * second field determines if there is a constrained face with that - * direction as normal. The last field determines if there is a - * constrained edge of a given pair of coordinate planes, but where - * neither of the corresponding faces are constrained (only valid in 3D). - * - * The element is placed in the 'first position' along *-axis. These also - * determine which face is constrained. For example, in 2D, if - * face_x and type are set, then x = 0 is constrained. - */ - enum ConstraintTypes : unsigned int - { - unconstrained = 0, - - type_x = 1 << 0, - type_y = 1 << 1, - type_z = 1 << 2, - - // Element has as a constraint at * = 0 or * = fe_degree face - face_x = 1 << 3, - face_y = 1 << 4, - face_z = 1 << 5, - - // Element has as a constraint at * = 0 or * = fe_degree edge - edge_xy = 1 << 6, - edge_yz = 1 << 7, - edge_zx = 1 << 8 - }; template @@ -238,7 +290,7 @@ namespace internal const std::shared_ptr &partitioner, const std::vector & lexicographic_mapping, std::vector &dof_indices, - const ArrayView & masks) const + const ArrayView & masks) const { bool cell_has_hanging_node_constraints = false; @@ -277,7 +329,7 @@ namespace internal ++c, ++comp) { auto &mask = masks[comp]; - mask = 0; + mask = ConstraintTypes::unconstrained; const auto &fe = cell->get_fe().base_element(base_element_index); @@ -478,7 +530,7 @@ namespace internal // For each line on cell, which faces does it belong to, what is // the edge mask, what is the types of the faces it belong to, // and what is the type along the edge. - const unsigned int line_to_edge[12][4] = { + const ConstraintTypes line_to_edge[12][4] = { {ConstraintTypes::face_x | ConstraintTypes::face_z, ConstraintTypes::edge_zx, ConstraintTypes::type_x | ConstraintTypes::type_z, @@ -501,7 +553,7 @@ namespace internal ConstraintTypes::type_y}, {ConstraintTypes::face_x | ConstraintTypes::face_z, ConstraintTypes::edge_zx, - 0, + ConstraintTypes::unconstrained, ConstraintTypes::type_y}, {ConstraintTypes::face_y | ConstraintTypes::face_z, ConstraintTypes::edge_yz, @@ -509,7 +561,7 @@ namespace internal ConstraintTypes::type_x}, {ConstraintTypes::face_y | ConstraintTypes::face_z, ConstraintTypes::edge_yz, - 0, + ConstraintTypes::unconstrained, ConstraintTypes::type_x}, {ConstraintTypes::face_x | ConstraintTypes::face_y, ConstraintTypes::edge_xy, @@ -525,7 +577,7 @@ namespace internal ConstraintTypes::type_z}, {ConstraintTypes::face_x | ConstraintTypes::face_y, ConstraintTypes::edge_xy, - 0, + ConstraintTypes::unconstrained, ConstraintTypes::type_z}}; for (unsigned int local_line = 0; @@ -817,7 +869,10 @@ namespace CUDAWrappers bool transpose, typename Number> __device__ inline void - interpolate_boundary_2d(const unsigned int constraint_mask, Number *values) + interpolate_boundary_2d( + const dealii::internal::MatrixFreeFunctions::ConstraintTypes + constraint_mask, + Number *values) { const unsigned int x_idx = threadIdx.x % (fe_degree + 1); const unsigned int y_idx = threadIdx.y; @@ -836,10 +891,12 @@ namespace CUDAWrappers (constraint_mask & (((direction == 0) ? dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y : - 0) | + dealii::internal::MatrixFreeFunctions::ConstraintTypes:: + unconstrained) | ((direction == 1) ? dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x : - 0))); + dealii::internal::MatrixFreeFunctions::ConstraintTypes:: + 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 (= @@ -911,7 +968,10 @@ namespace CUDAWrappers bool transpose, typename Number> __device__ inline void - interpolate_boundary_3d(const unsigned int constraint_mask, Number *values) + interpolate_boundary_3d( + const dealii::internal::MatrixFreeFunctions::ConstraintTypes + constraint_mask, + Number *values) { const unsigned int x_idx = threadIdx.x % (fe_degree + 1); const unsigned int y_idx = threadIdx.y; @@ -1043,7 +1103,10 @@ namespace CUDAWrappers */ template __device__ void - resolve_hanging_nodes(const unsigned int constraint_mask, Number *values) + resolve_hanging_nodes( + const dealii::internal::MatrixFreeFunctions::ConstraintTypes + constraint_mask, + Number *values) { if (dim == 2) { diff --git a/include/deal.II/matrix_free/tools.h b/include/deal.II/matrix_free/tools.h index a0fda74183..71a3859131 100644 --- a/include/deal.II/matrix_free/tools.h +++ b/include/deal.II/matrix_free/tools.h @@ -465,7 +465,9 @@ namespace MatrixFreeTools AlignedVector values_dofs( dofs_per_component); - std::array + std::array constraint_mask; constraint_mask[0] = mask; diff --git a/source/matrix_free/dof_info.cc b/source/matrix_free/dof_info.cc index 0d78a5ad5f..9dbbb3fe54 100644 --- a/source/matrix_free/dof_info.cc +++ b/source/matrix_free/dof_info.cc @@ -336,7 +336,7 @@ namespace internal new_dof_indices.reserve(dof_indices.size()); new_constraint_indicator.reserve(constraint_indicator.size()); - std::vector new_hanging_node_constraint_masks; + std::vector new_hanging_node_constraint_masks; new_hanging_node_constraint_masks.reserve( new_hanging_node_constraint_masks.size()); @@ -425,7 +425,8 @@ namespace internal .second = new_constraint_indicator.size(); if (hanging_node_constraint_masks.size() > 0) - new_hanging_node_constraint_masks.push_back(0); + new_hanging_node_constraint_masks.push_back( + ConstraintTypes::unconstrained); } position_cell += n_vect; } diff --git a/tests/matrix_free/hanging_node_kernels_01.cc b/tests/matrix_free/hanging_node_kernels_01.cc index dec80c3d68..fb56496180 100644 --- a/tests/matrix_free/hanging_node_kernels_01.cc +++ b/tests/matrix_free/hanging_node_kernels_01.cc @@ -45,13 +45,15 @@ namespace dealii { template static bool - run(const FEEvaluationBaseData & fe_eval, - const bool transpose, - const std::array &c_mask, - Number * values) + run( + const FEEvaluationBaseData &fe_eval, + const bool transpose, + const std::array & c_mask, + Number * values) { Assert(is_face == false, ExcInternalError()); @@ -102,12 +104,14 @@ namespace dealii template static void - run_2D(const FEEvaluationBaseData & fe_eval, - const std::array &constraint_mask, - Number * values) + run_2D( + const FEEvaluationBaseData &fe_eval, + const std::array & constraint_mask, + Number * values) { const auto &constraint_weights = fe_eval.get_shape_info().data.front().subface_interpolation_matrix; @@ -232,12 +236,14 @@ namespace dealii template static void - run_3D(const FEEvaluationBaseData & fe_eval, - const std::array &constraint_mask, - Number * values) + run_3D( + const FEEvaluationBaseData &fe_eval, + const std::array & constraint_mask, + Number * values) { const auto &constraint_weights = fe_eval.get_shape_info().data.front().subface_interpolation_matrix; @@ -391,7 +397,8 @@ namespace dealii template void -test(const unsigned int degree, const unsigned int mask_value) +test(const unsigned int degree, + const dealii::internal::MatrixFreeFunctions::ConstraintTypes mask_value) { Triangulation tria; GridGenerator::subdivided_hyper_cube(tria, 2); @@ -420,8 +427,13 @@ test(const unsigned int degree, const unsigned int mask_value) FEEvaluation eval(matrix_free); eval.reinit(0); - std::array::size()> cmask; - std::fill(cmask.begin(), cmask.end(), 0); + std::array::size()> + cmask; + std::fill( + cmask.begin(), + cmask.end(), + dealii::internal::MatrixFreeFunctions::ConstraintTypes::unconstrained); cmask[0] = mask_value; for (unsigned int b = 0; b < 2; ++b) @@ -473,49 +485,47 @@ main(int argc, char **argv) for (unsigned int degree = 1; degree <= 3; ++degree) { - test<2>(degree, 0); + test<2>( + degree, + dealii::internal::MatrixFreeFunctions::ConstraintTypes::unconstrained); deallog << std::endl; test<2>(degree, - 0 | - dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x | + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x | dealii::internal::MatrixFreeFunctions::ConstraintTypes::type_x | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: type_y); // face 0/0 test<2>(degree, - 0 | - dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x | + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: type_x); // face 0/1 test<2>(degree, - 0 | - dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x | + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: type_y); // face 1/0 - test<2>(degree, - 0 | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: - face_x); // face 1/1 + test<2>( + degree, + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_x); // face + // 1/1 deallog << std::endl; test<2>(degree, - 0 | - dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y | + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y | dealii::internal::MatrixFreeFunctions::ConstraintTypes::type_y | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: type_x); // face 2/0 test<2>(degree, - 0 | - dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y | + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: type_y); // face 2/1 test<2>(degree, - 0 | - dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y | + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: type_x); // face 3/0 - test<2>(degree, - 0 | dealii::internal::MatrixFreeFunctions::ConstraintTypes:: - face_y); // face 3/1 + test<2>( + degree, + dealii::internal::MatrixFreeFunctions::ConstraintTypes::face_y); // face + // 3/1 deallog << std::endl; } -- 2.39.5