From 37eb4e1377a496f840c195129325936a472a8ebf Mon Sep 17 00:00:00 2001
From: john bowen <bowen36@llnl.gov>
Date: Wed, 1 Nov 2023 15:46:38 -0700
Subject: [PATCH] Add comments explaining trial_elements_tuple and refactor for
 clarity.  Clang format

---
 .../functional/boundary_integral_kernels.hpp  | 16 ++--------
 .../functional/detail/hexahedron_H1.inl       |  6 ++--
 .../functional/detail/hexahedron_Hcurl.inl    |  7 ++--
 .../functional/detail/hexahedron_L2.inl       |  6 ++--
 .../functional/detail/metaprogramming.hpp     |  3 +-
 src/serac/numerics/functional/detail/qoi.inl  |  7 ++--
 .../functional/detail/quadrilateral_H1.inl    |  6 ++--
 .../functional/detail/quadrilateral_Hcurl.inl |  7 ++--
 .../functional/detail/quadrilateral_L2.inl    |  6 ++--
 .../numerics/functional/detail/segment_H1.inl | 10 +++---
 .../numerics/functional/detail/segment_L2.inl |  7 ++--
 .../functional/detail/tetrahedron_H1.inl      |  6 ++--
 .../functional/detail/tetrahedron_L2.inl      |  6 ++--
 .../functional/detail/triangle_H1.inl         |  6 ++--
 .../functional/detail/triangle_L2.inl         |  6 ++--
 .../numerics/functional/differentiate_wrt.hpp |  5 +--
 .../functional/domain_integral_kernels.hpp    | 16 ++--------
 .../numerics/functional/finite_element.hpp    |  5 +--
 .../functional/function_signature.hpp         | 32 ++++++++++++++++++-
 src/serac/numerics/functional/functional.hpp  |  5 +--
 .../numerics/functional/quadrature_data.hpp   |  6 ++--
 src/serac/numerics/functional/tensor.hpp      |  6 ++--
 .../numerics/functional/tests/CMakeLists.txt  |  6 +---
 .../functional/tests/hcurl_unit_tests.cpp     |  4 +--
 .../numerics/functional/tests/test_newton.cpp |  3 +-
 src/serac/numerics/functional/tuple.hpp       | 21 ++++++++----
 .../tuple_tensor_dual_functions.hpp           |  3 +-
 src/serac/physics/heat_transfer.hpp           |  3 +-
 28 files changed, 118 insertions(+), 102 deletions(-)

diff --git a/src/serac/numerics/functional/boundary_integral_kernels.hpp b/src/serac/numerics/functional/boundary_integral_kernels.hpp
index 6d2fa1381..3959ac96c 100644
--- a/src/serac/numerics/functional/boundary_integral_kernels.hpp
+++ b/src/serac/numerics/functional/boundary_integral_kernels.hpp
@@ -153,18 +153,6 @@ auto batch_apply_qf(lambda qf, const tensor<double, 3, n>& positions, const tens
   return outputs;
 }
 
-template <mfem::Geometry::Type geom, typename test, typename... trials>
-auto get_trial_elements(FunctionSignature<test(trials...)>)
-{
-  return tuple<finite_element<geom, trials>...>{};
-}
-
-template <mfem::Geometry::Type geom, typename test, typename... trials>
-auto get_test(FunctionSignature<test(trials...)>)
-{
-  return finite_element<geom, test>{};
-}
-
 /// @trial_elements the element type for each trial space
 template <uint32_t differentiation_index, int Q, mfem::Geometry::Type geom, typename test_element,
           typename trial_element_type, typename lambda_type, typename derivative_type, int... indices>
@@ -339,8 +327,8 @@ std::function<void(const std::vector<const double*>&, double*, bool)> evaluation
     signature s, lambda_type qf, const double* positions, const double* jacobians,
     std::shared_ptr<derivative_type> qf_derivatives, uint32_t num_elements)
 {
-  auto trial_elements = get_trial_elements<geom>(s);
-  auto test           = get_test<geom>(s);
+  auto trial_elements = trial_elements_tuple<geom>(s);
+  auto test           = get_test_element<geom>(s);
   return [=](const std::vector<const double*>& inputs, double* outputs, bool /* update state */) {
     evaluation_kernel_impl<wrt, Q, geom>(trial_elements, test, inputs, outputs, positions, jacobians, qf,
                                          qf_derivatives.get(), num_elements, s.index_seq);
diff --git a/src/serac/numerics/functional/detail/hexahedron_H1.inl b/src/serac/numerics/functional/detail/hexahedron_H1.inl
index 09419cabe..2a5ef7479 100644
--- a/src/serac/numerics/functional/detail/hexahedron_H1.inl
+++ b/src/serac/numerics/functional/detail/hexahedron_H1.inl
@@ -167,8 +167,7 @@ struct finite_element<mfem::Geometry::CUBE, H1<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
   {
     // we want to compute the following:
     //
@@ -230,7 +229,8 @@ static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q * q * q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/hexahedron_Hcurl.inl b/src/serac/numerics/functional/detail/hexahedron_Hcurl.inl
index d85fb25ab..b6277937b 100644
--- a/src/serac/numerics/functional/detail/hexahedron_Hcurl.inl
+++ b/src/serac/numerics/functional/detail/hexahedron_Hcurl.inl
@@ -334,8 +334,7 @@ struct finite_element<mfem::Geometry::CUBE, Hcurl<p>> {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& element_values, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& element_values, const TensorProductQuadratureRule<q>&)
   {
     constexpr bool                     apply_weights = false;
     constexpr tensor<double, q, p>     B1            = calculate_B1<apply_weights, q>();
@@ -397,8 +396,8 @@ static auto interpolate(const dof_type& element_values, const TensorProductQuadr
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q * q * q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual,
-                        [[maybe_unused]] int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          [[maybe_unused]] int step = 1)
   {
     constexpr bool                     apply_weights = true;
     constexpr tensor<double, q, p>     B1            = calculate_B1<apply_weights, q>();
diff --git a/src/serac/numerics/functional/detail/hexahedron_L2.inl b/src/serac/numerics/functional/detail/hexahedron_L2.inl
index 03bba1d95..51efeaeec 100644
--- a/src/serac/numerics/functional/detail/hexahedron_L2.inl
+++ b/src/serac/numerics/functional/detail/hexahedron_L2.inl
@@ -171,8 +171,7 @@ struct finite_element<mfem::Geometry::CUBE, L2<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
   {
     // we want to compute the following:
     //
@@ -234,7 +233,8 @@ static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q * q * q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/metaprogramming.hpp b/src/serac/numerics/functional/detail/metaprogramming.hpp
index eec9833b8..216e032e3 100644
--- a/src/serac/numerics/functional/detail/metaprogramming.hpp
+++ b/src/serac/numerics/functional/detail/metaprogramming.hpp
@@ -32,7 +32,8 @@ constexpr auto get(std::integer_sequence<int, n...>)
 namespace detail {
 
 template <typename T>
-struct always_false : std::false_type {};
+struct always_false : std::false_type {
+};
 
 /**
  * @brief unfortunately std::integral_constant doesn't have __host__ __device__ annotations
diff --git a/src/serac/numerics/functional/detail/qoi.inl b/src/serac/numerics/functional/detail/qoi.inl
index 26c8879f4..189b7f205 100644
--- a/src/serac/numerics/functional/detail/qoi.inl
+++ b/src/serac/numerics/functional/detail/qoi.inl
@@ -25,14 +25,14 @@ struct finite_element<g, QOI> {
 
   template <int Q, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<zero, Q>&, const TensorProductQuadratureRule<q>&, dof_type*,
-                        [[maybe_unused]] int step = 1)
+                                          [[maybe_unused]] int step = 1)
   {
     return;  // integrating zeros is a no-op
   }
 
   template <int Q, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<double, Q>& qf_output, const TensorProductQuadratureRule<q>&,
-                        dof_type* element_total, [[maybe_unused]] int step = 1)
+                                          dof_type* element_total, [[maybe_unused]] int step = 1)
   {
     if constexpr (geometry == mfem::Geometry::SEGMENT) {
       static_assert(Q == q);
@@ -78,7 +78,8 @@ struct finite_element<g, QOI> {
   // output to be a tuple with a hardcoded `zero` flux term
   template <typename source_type, int Q, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<serac::tuple<source_type, zero>, Q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_total, [[maybe_unused]] int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_total,
+                                          [[maybe_unused]] int step = 1)
   {
     if constexpr (is_zero<source_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/quadrilateral_H1.inl b/src/serac/numerics/functional/detail/quadrilateral_H1.inl
index edd1c929d..452bd6e17 100644
--- a/src/serac/numerics/functional/detail/quadrilateral_H1.inl
+++ b/src/serac/numerics/functional/detail/quadrilateral_H1.inl
@@ -202,8 +202,7 @@ struct finite_element<mfem::Geometry::SQUARE, H1<p, c> > {
   // A(dy, qx)  := B(qx, dx) * X_e(dy, dx)
   // X_q(qy, qx) := B(qy, dy) * A(dy, qx)
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
   {
     static constexpr bool apply_weights = false;
     static constexpr auto B             = calculate_B<apply_weights, q>();
@@ -247,7 +246,8 @@ static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&
   // tensor<double,dim,dim,dim>}
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q * q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/quadrilateral_Hcurl.inl b/src/serac/numerics/functional/detail/quadrilateral_Hcurl.inl
index 7eed46251..6e332f27e 100644
--- a/src/serac/numerics/functional/detail/quadrilateral_Hcurl.inl
+++ b/src/serac/numerics/functional/detail/quadrilateral_Hcurl.inl
@@ -285,8 +285,7 @@ struct finite_element<mfem::Geometry::SQUARE, Hcurl<p> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& element_values, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& element_values, const TensorProductQuadratureRule<q>&)
   {
     constexpr bool                     apply_weights = false;
     constexpr tensor<double, q, p>     B1            = calculate_B1<apply_weights, q>();
@@ -325,8 +324,8 @@ static auto interpolate(const dof_type& element_values, const TensorProductQuadr
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q * q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual,
-                        [[maybe_unused]] int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          [[maybe_unused]] int step = 1)
   {
     constexpr bool                     apply_weights = true;
     constexpr tensor<double, q, p>     B1            = calculate_B1<apply_weights, q>();
diff --git a/src/serac/numerics/functional/detail/quadrilateral_L2.inl b/src/serac/numerics/functional/detail/quadrilateral_L2.inl
index af2edc93e..09d653b5f 100644
--- a/src/serac/numerics/functional/detail/quadrilateral_L2.inl
+++ b/src/serac/numerics/functional/detail/quadrilateral_L2.inl
@@ -201,8 +201,7 @@ struct finite_element<mfem::Geometry::SQUARE, L2<p, c> > {
   // A(dy, qx)  := B(qx, dx) * X_e(dy, dx)
   // X_q(qy, qx) := B(qy, dy) * A(dy, qx)
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
   {
     static constexpr bool apply_weights = false;
     static constexpr auto B             = calculate_B<apply_weights, q>();
@@ -246,7 +245,8 @@ static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&
   // tensor<double,dim,dim,dim>}
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q * q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/segment_H1.inl b/src/serac/numerics/functional/detail/segment_H1.inl
index 95d8a0924..5651be12d 100644
--- a/src/serac/numerics/functional/detail/segment_H1.inl
+++ b/src/serac/numerics/functional/detail/segment_H1.inl
@@ -94,8 +94,7 @@ struct finite_element<mfem::Geometry::SEGMENT, H1<p, c> > {
   }
 
   template <typename T, int q>
-  SERAC_HOST_DEVICE
-  static auto batch_apply_shape_fn(int jx, tensor<T, q> input, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto batch_apply_shape_fn(int jx, tensor<T, q> input, const TensorProductQuadratureRule<q>&)
   {
     static constexpr bool apply_weights = false;
     static constexpr auto B             = calculate_B<apply_weights, q>();
@@ -122,8 +121,7 @@ struct finite_element<mfem::Geometry::SEGMENT, H1<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
   {
     static constexpr bool apply_weights = false;
     static constexpr auto B             = calculate_B<apply_weights, q>();
@@ -158,8 +156,8 @@ static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual,
-                        [[maybe_unused]] int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          [[maybe_unused]] int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/segment_L2.inl b/src/serac/numerics/functional/detail/segment_L2.inl
index 9b24438c8..214e757dc 100644
--- a/src/serac/numerics/functional/detail/segment_L2.inl
+++ b/src/serac/numerics/functional/detail/segment_L2.inl
@@ -121,8 +121,7 @@ struct finite_element<mfem::Geometry::SEGMENT, L2<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&)
   {
     static constexpr bool apply_weights = false;
     static constexpr auto B             = calculate_B<apply_weights, q>();
@@ -157,8 +156,8 @@ static auto interpolate(const dof_type& X, const TensorProductQuadratureRule<q>&
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q>& qf_output,
-                        const TensorProductQuadratureRule<q>&, dof_type* element_residual,
-                        [[maybe_unused]] int step = 1)
+                                          const TensorProductQuadratureRule<q>&, dof_type* element_residual,
+                                          [[maybe_unused]] int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/tetrahedron_H1.inl b/src/serac/numerics/functional/detail/tetrahedron_H1.inl
index a2540576b..d3ea86061 100644
--- a/src/serac/numerics/functional/detail/tetrahedron_H1.inl
+++ b/src/serac/numerics/functional/detail/tetrahedron_H1.inl
@@ -361,8 +361,7 @@ struct finite_element<mfem::Geometry::TETRAHEDRON, H1<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
   {
     constexpr auto xi = GaussLegendreNodes<q, mfem::Geometry::TETRAHEDRON>();
 
@@ -386,7 +385,8 @@ static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQua
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, nqpts(q)>& qf_output,
-                        const TensorProductQuadratureRule<q>&, tensor<double, c, ndof>* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&,
+                                          tensor<double, c, ndof>* element_residual, int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/tetrahedron_L2.inl b/src/serac/numerics/functional/detail/tetrahedron_L2.inl
index a24c1019a..df30f3409 100644
--- a/src/serac/numerics/functional/detail/tetrahedron_L2.inl
+++ b/src/serac/numerics/functional/detail/tetrahedron_L2.inl
@@ -366,8 +366,7 @@ struct finite_element<mfem::Geometry::TETRAHEDRON, L2<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
   {
     constexpr auto xi = GaussLegendreNodes<q, mfem::Geometry::TETRAHEDRON>();
 
@@ -391,7 +390,8 @@ static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQua
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, nqpts(q)>& qf_output,
-                        const TensorProductQuadratureRule<q>&, tensor<double, c, ndof>* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&,
+                                          tensor<double, c, ndof>* element_residual, int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/triangle_H1.inl b/src/serac/numerics/functional/detail/triangle_H1.inl
index 11425cb77..fe2bc560d 100644
--- a/src/serac/numerics/functional/detail/triangle_H1.inl
+++ b/src/serac/numerics/functional/detail/triangle_H1.inl
@@ -266,8 +266,7 @@ struct finite_element<mfem::Geometry::TRIANGLE, H1<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
   {
     constexpr auto       xi                    = GaussLegendreNodes<q, mfem::Geometry::TRIANGLE>();
     static constexpr int num_quadrature_points = q * (q + 1) / 2;
@@ -292,7 +291,8 @@ static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQua
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q*(q + 1) / 2>& qf_output,
-                        const TensorProductQuadratureRule<q>&, tensor<double, c, ndof>* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&,
+                                          tensor<double, c, ndof>* element_residual, int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/detail/triangle_L2.inl b/src/serac/numerics/functional/detail/triangle_L2.inl
index f55101de4..51f64af5e 100644
--- a/src/serac/numerics/functional/detail/triangle_L2.inl
+++ b/src/serac/numerics/functional/detail/triangle_L2.inl
@@ -275,8 +275,7 @@ struct finite_element<mfem::Geometry::TRIANGLE, L2<p, c> > {
   }
 
   template <int q>
-  SERAC_HOST_DEVICE
-static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
+  SERAC_HOST_DEVICE static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQuadratureRule<q>&)
   {
     constexpr auto       xi                    = GaussLegendreNodes<q, mfem::Geometry::TRIANGLE>();
     static constexpr int num_quadrature_points = q * (q + 1) / 2;
@@ -301,7 +300,8 @@ static auto interpolate(const tensor<double, c, ndof>& X, const TensorProductQua
 
   template <typename source_type, typename flux_type, int q>
   SERAC_HOST_DEVICE static void integrate(const tensor<tuple<source_type, flux_type>, q*(q + 1) / 2>& qf_output,
-                        const TensorProductQuadratureRule<q>&, tensor<double, c, ndof>* element_residual, int step = 1)
+                                          const TensorProductQuadratureRule<q>&,
+                                          tensor<double, c, ndof>* element_residual, int step = 1)
   {
     if constexpr (is_zero<source_type>{} && is_zero<flux_type>{}) {
       return;
diff --git a/src/serac/numerics/functional/differentiate_wrt.hpp b/src/serac/numerics/functional/differentiate_wrt.hpp
index b563b22d6..85262a794 100644
--- a/src/serac/numerics/functional/differentiate_wrt.hpp
+++ b/src/serac/numerics/functional/differentiate_wrt.hpp
@@ -12,7 +12,8 @@ namespace serac {
 static constexpr uint32_t NO_DIFFERENTIATION = uint32_t(1) << 31;
 
 template <uint32_t i>
-struct DifferentiateWRT {};
+struct DifferentiateWRT {
+};
 
 /**
  * @brief this type exists solely as a way to signal to `serac::Functional` that the function
@@ -22,7 +23,7 @@ struct differentiate_wrt_this {
   const mfem::Vector& ref;  ///< the actual data wrapped by this type
 
   /// @brief implicitly convert back to `mfem::Vector` to extract the actual data
-  operator const mfem::Vector&() const { return ref; }
+  operator const mfem::Vector &() const { return ref; }
 };
 
 /**
diff --git a/src/serac/numerics/functional/domain_integral_kernels.hpp b/src/serac/numerics/functional/domain_integral_kernels.hpp
index 7c86fc673..b97ec15c5 100644
--- a/src/serac/numerics/functional/domain_integral_kernels.hpp
+++ b/src/serac/numerics/functional/domain_integral_kernels.hpp
@@ -137,18 +137,6 @@ SERAC_HOST_DEVICE auto batch_apply_qf(lambda qf, const tensor<double, dim, n> x,
   return outputs;
 }
 
-template <mfem::Geometry::Type geom, typename test, typename... trials>
-auto get_trial_elements(FunctionSignature<test(trials...)>)
-{
-  return tuple<finite_element<geom, trials>...>{};
-}
-
-template <mfem::Geometry::Type geom, typename test, typename... trials>
-auto get_test(FunctionSignature<test(trials...)>)
-{
-  return finite_element<geom, test>{};
-}
-
 template <uint32_t differentiation_index, int Q, mfem::Geometry::Type geom, typename test_element,
           typename trial_element_type, typename lambda_type, typename state_type, typename derivative_type,
           int... indices>
@@ -365,8 +353,8 @@ std::function<void(const std::vector<const double*>&, double*, bool)> evaluation
     std::shared_ptr<QuadratureData<state_type> > qf_state, std::shared_ptr<derivative_type> qf_derivatives,
     uint32_t num_elements)
 {
-  auto trial_elements = get_trial_elements<geom>(s);
-  auto test           = get_test<geom>(s);
+  auto trial_elements = trial_elements_tuple<geom>(s);
+  auto test           = get_test_element<geom>(s);
   return [=](const std::vector<const double*>& inputs, double* outputs, bool update_state) {
     domain_integral::evaluation_kernel_impl<wrt, Q, geom>(trial_elements, test, inputs, outputs, positions, jacobians,
                                                           qf, (*qf_state)[geom], qf_derivatives.get(), num_elements,
diff --git a/src/serac/numerics/functional/finite_element.hpp b/src/serac/numerics/functional/finite_element.hpp
index d096a7c6b..28174f703 100644
--- a/src/serac/numerics/functional/finite_element.hpp
+++ b/src/serac/numerics/functional/finite_element.hpp
@@ -41,7 +41,8 @@ struct TensorProductQuadratureRule {
 };
 
 template <auto val>
-struct CompileTimeValue {};
+struct CompileTimeValue {
+};
 
 /**
  * @brief this struct is used to look up mfem's memory layout of
@@ -346,7 +347,7 @@ SERAC_HOST_DEVICE void physical_to_parent(tensor<T, q>& qf_output, const tensor<
  *
  */
 template <mfem::Geometry::Type g, typename family>
-SERAC_HOST_DEVICE struct finite_element;
+struct finite_element;
 
 #include "detail/segment_H1.inl"
 #include "detail/segment_Hcurl.inl"
diff --git a/src/serac/numerics/functional/function_signature.hpp b/src/serac/numerics/functional/function_signature.hpp
index 00bbf12ea..e5ce5d9bb 100644
--- a/src/serac/numerics/functional/function_signature.hpp
+++ b/src/serac/numerics/functional/function_signature.hpp
@@ -1,5 +1,6 @@
-#include <camp/camp.hpp>
 #pragma once
+#include <camp/camp.hpp>
+#include "serac/numerics/functional/finite_element.hpp"
 
 template <typename T>
 struct FunctionSignature;
@@ -20,3 +21,32 @@ struct FunctionSignature<output_type(input_types...)> {
   /// integer sequence used to make iterating over arguments easier
   static constexpr auto index_seq = camp::make_int_seq<int, num_args>{};
 };
+
+/**
+ * @brief This helper function template expands the variadic parameter trials and uses
+ * the pack to construct a tuple of finite elements.
+ *
+ * Expansion of the template parameter trials must occur in a separate template
+ * function from evaluation_kernel_impl.  The __host__ __device__ marker on the
+ * lambda declared in evaluation_kernel_impl prevents use of multiple variadic
+ * template parameters in evaluation_kernel_impl.  See section 14.7.3 of the
+ * CUDA programming guide, item 9.
+ */
+template <mfem::Geometry::Type geom, typename test, typename... trials>
+auto trial_elements_tuple(FunctionSignature<test(trials...)>)
+{
+  return serac::tuple<serac::finite_element<geom, trials>...>{};
+}
+
+/**
+ * @brief This helper function template returns a finite element based on the output
+ * type of FunctionSignature.
+ *
+ * See the comments for trial_elements_tuple to understand why this is needed in
+ * domain_integral_kernels::evaluation_kernel and boundary_integral_kernels::evaluation_kernel.
+ */
+template <mfem::Geometry::Type geom, typename test, typename... trials>
+auto get_test_element(FunctionSignature<test(trials...)>)
+{
+  return serac::finite_element<geom, test>{};
+}
diff --git a/src/serac/numerics/functional/functional.hpp b/src/serac/numerics/functional/functional.hpp
index d9f7c53ab..4b008c295 100644
--- a/src/serac/numerics/functional/functional.hpp
+++ b/src/serac/numerics/functional/functional.hpp
@@ -30,7 +30,8 @@
 namespace serac {
 
 template <int... i>
-struct DependsOn {};
+struct DependsOn {
+};
 
 /**
  * @brief given a list of types, this function returns the index that corresponds to the type `dual_vector`.
@@ -189,7 +190,7 @@ class Functional<test(trials...), exec> {
   static constexpr auto             Q                = std::max({test::order, trials::order...}) + 1;
 
   static constexpr mfem::Geometry::Type elem_geom[4]    = {mfem::Geometry::INVALID, mfem::Geometry::SEGMENT,
-                                                           mfem::Geometry::SQUARE, mfem::Geometry::CUBE};
+                                                        mfem::Geometry::SQUARE, mfem::Geometry::CUBE};
   static constexpr mfem::Geometry::Type simplex_geom[4] = {mfem::Geometry::INVALID, mfem::Geometry::SEGMENT,
                                                            mfem::Geometry::TRIANGLE, mfem::Geometry::TETRAHEDRON};
 
diff --git a/src/serac/numerics/functional/quadrature_data.hpp b/src/serac/numerics/functional/quadrature_data.hpp
index 3627938ec..48f633857 100644
--- a/src/serac/numerics/functional/quadrature_data.hpp
+++ b/src/serac/numerics/functional/quadrature_data.hpp
@@ -35,12 +35,14 @@ namespace serac {
  *        will also store no data, but it will still appear as an argument to the q-function
  *        (to make the material model interface consistent).
  */
-struct Nothing {};
+struct Nothing {
+};
 
 /**
  * @brief see `Nothing` for a complete description of this class and when to use it
  */
-struct Empty {};
+struct Empty {
+};
 
 template <typename T>
 struct QuadratureData;
diff --git a/src/serac/numerics/functional/tensor.hpp b/src/serac/numerics/functional/tensor.hpp
index b6995bb6d..bb89f83d9 100644
--- a/src/serac/numerics/functional/tensor.hpp
+++ b/src/serac/numerics/functional/tensor.hpp
@@ -142,11 +142,13 @@ struct zero {
 
 /** @brief checks if a type is `zero` */
 template <typename T>
-struct is_zero : std::false_type {};
+struct is_zero : std::false_type {
+};
 
 /** @overload */
 template <>
-struct is_zero<zero> : std::true_type {};
+struct is_zero<zero> : std::true_type {
+};
 
 /** @brief the sum of two `zero`s is `zero` */
 SERAC_HOST_DEVICE constexpr auto operator+(zero, zero) { return zero{}; }
diff --git a/src/serac/numerics/functional/tests/CMakeLists.txt b/src/serac/numerics/functional/tests/CMakeLists.txt
index f36ce2e86..4ba9132ff 100644
--- a/src/serac/numerics/functional/tests/CMakeLists.txt
+++ b/src/serac/numerics/functional/tests/CMakeLists.txt
@@ -36,7 +36,7 @@ set(functional_tests_mpi
     functional_comparisons.cpp
     functional_comparison_L2.cpp
     )
-
+    
 serac_add_tests( SOURCES ${functional_tests_mpi}
                  DEPENDS_ON gtest serac_functional serac_state ${functional_depends}
                  NUM_MPI_TASKS 4 )
@@ -58,8 +58,4 @@ if(ENABLE_CUDA)
     serac_add_tests( SOURCES ${functional_tests_cuda}
                      DEPENDS_ON gtest serac_functional serac_state ${functional_depends} cuda)
 
-            
-    #blt_add_executable(NAME functional_cuda
-    #                   SOURCES functional_cuda.cpp
-    #                   DEPENDS_ON gtest serac_functional serac_state ${functional_depends} cuda)
 endif()
diff --git a/src/serac/numerics/functional/tests/hcurl_unit_tests.cpp b/src/serac/numerics/functional/tests/hcurl_unit_tests.cpp
index 9a9a38d17..bfd2a91cc 100644
--- a/src/serac/numerics/functional/tests/hcurl_unit_tests.cpp
+++ b/src/serac/numerics/functional/tests/hcurl_unit_tests.cpp
@@ -18,8 +18,8 @@ static constexpr double curl_tolerance      = 1.0e-10;  // coarser, since compar
 static constexpr int    num_points          = 10;
 static constexpr tensor random_numbers      = {
     {-0.886787, -0.850126, 0.464212,  -0.0733101, -0.397738, 0.302355,   -0.570758, 0.977727,  0.282365,  -0.768947,
-          0.6216,    0.43598,   -0.696321, 0.92545,    0.183003,  0.121761,   -0.877239, 0.0347577, -0.818463, -0.216474,
-          -0.43894,  0.0178874, -0.869944, -0.733499,  0.255124,  -0.0561095, -0.34607,  -0.305958, 0.414472,  -0.744998}};
+     0.6216,    0.43598,   -0.696321, 0.92545,    0.183003,  0.121761,   -0.877239, 0.0347577, -0.818463, -0.216474,
+     -0.43894,  0.0178874, -0.869944, -0.733499,  0.255124,  -0.0561095, -0.34607,  -0.305958, 0.414472,  -0.744998}};
 
 /*
 iterate over the node/direction pairs in the element, and verify that
diff --git a/src/serac/numerics/functional/tests/test_newton.cpp b/src/serac/numerics/functional/tests/test_newton.cpp
index 2d284396c..5ad28b73d 100644
--- a/src/serac/numerics/functional/tests/test_newton.cpp
+++ b/src/serac/numerics/functional/tests/test_newton.cpp
@@ -11,7 +11,8 @@
 
 #include "serac/numerics/functional/tuple_tensor_dual_functions.hpp"
 
-class SlicErrorException : public std::exception {};
+class SlicErrorException : public std::exception {
+};
 
 using namespace serac;
 
diff --git a/src/serac/numerics/functional/tuple.hpp b/src/serac/numerics/functional/tuple.hpp
index 91cab94a5..8c614cf6a 100644
--- a/src/serac/numerics/functional/tuple.hpp
+++ b/src/serac/numerics/functional/tuple.hpp
@@ -25,7 +25,8 @@ namespace serac {
  * see https://en.cppreference.com/w/cpp/utility/tuple for more information about std::tuple
  */
 template <typename... T>
-struct tuple {};
+struct tuple {
+};
 
 /**
  * @brief Type that mimics std::tuple
@@ -182,10 +183,12 @@ SERAC_HOST_DEVICE tuple<T...> make_tuple(const T&... args)
 }
 
 template <class... Types>
-struct tuple_size {};
+struct tuple_size {
+};
 
 template <class... Types>
-struct tuple_size<serac::tuple<Types...>> : std::integral_constant<std::size_t, sizeof...(Types)> {};
+struct tuple_size<serac::tuple<Types...>> : std::integral_constant<std::size_t, sizeof...(Types)> {
+};
 
 /**
  * @tparam i the tuple index to access
@@ -705,7 +708,8 @@ struct tuple_element;
 // recursive case
 /// @overload
 template <size_t I, class Head, class... Tail>
-struct tuple_element<I, tuple<Head, Tail...>> : tuple_element<I - 1, tuple<Tail...>> {};
+struct tuple_element<I, tuple<Head, Tail...>> : tuple_element<I - 1, tuple<Tail...>> {
+};
 
 // base case
 /// @overload
@@ -718,17 +722,20 @@ struct tuple_element<0, tuple<Head, Tail...>> {
  * @brief Trait for checking if a type is a @p serac::tuple
  */
 template <typename T>
-struct is_tuple : std::false_type {};
+struct is_tuple : std::false_type {
+};
 
 /// @overload
 template <typename... T>
-struct is_tuple<serac::tuple<T...>> : std::true_type {};
+struct is_tuple<serac::tuple<T...>> : std::true_type {
+};
 
 /**
  * @brief Trait for checking if a type if a @p serac::tuple containing only @p serac::tuple
  */
 template <typename T>
-struct is_tuple_of_tuples : std::false_type {};
+struct is_tuple_of_tuples : std::false_type {
+};
 
 /**
  * @brief Trait for checking if a type if a @p serac::tuple containing only @p serac::tuple
diff --git a/src/serac/numerics/functional/tuple_tensor_dual_functions.hpp b/src/serac/numerics/functional/tuple_tensor_dual_functions.hpp
index ff07d2e7b..a1ef69f75 100644
--- a/src/serac/numerics/functional/tuple_tensor_dual_functions.hpp
+++ b/src/serac/numerics/functional/tuple_tensor_dual_functions.hpp
@@ -106,7 +106,8 @@ struct one_hot_helper<i, std::integer_sequence<int, I...>, T> {
 };
 
 template <int i, int n, typename T>
-struct one_hot : public one_hot_helper<i, std::make_integer_sequence<int, n>, T> {};
+struct one_hot : public one_hot_helper<i, std::make_integer_sequence<int, n>, T> {
+};
 /// @endcond
 
 /**
diff --git a/src/serac/physics/heat_transfer.hpp b/src/serac/physics/heat_transfer.hpp
index d254ae26d..9d0503f7f 100644
--- a/src/serac/physics/heat_transfer.hpp
+++ b/src/serac/physics/heat_transfer.hpp
@@ -303,7 +303,8 @@ class HeatTransfer<order, dim, Parameters<parameter_space...>, std::integer_sequ
     // to be portable to CUDA platforms, the asserts below prevent serac from compiling
     // if such a lambda is supplied.
   private:
-    class DummyArgumentType {};
+    class DummyArgumentType {
+    };
     static_assert(!std::is_invocable<MaterialType, DummyArgumentType&>::value);
     static_assert(!std::is_invocable<MaterialType, DummyArgumentType*>::value);
     static_assert(!std::is_invocable<MaterialType, DummyArgumentType>::value);