MeshFields
GPU accelerated mesh-based fields
MeshField.hpp
1 #ifndef MESHFIELD_MESHFIELD_HPP
2 #define MESHFIELD_MESHFIELD_HPP
3 
4 #include "KokkosController.hpp"
5 #include "MeshField_Element.hpp"
6 #include "MeshField_Fail.hpp"
7 #include "MeshField_For.hpp"
8 #include "MeshField_ShapeField.hpp"
9 #include "Omega_h_file.hpp" //move
10 #include "Omega_h_mesh.hpp" //move
11 #include "Omega_h_simplex.hpp" //move
12 
13 namespace {
14 
15 MeshField::MeshInfo getMeshInfo(Omega_h::Mesh &mesh) {
16  MeshField::MeshInfo meshInfo;
17  meshInfo.dim = mesh.dim();
18  meshInfo.numVtx = mesh.nverts();
19  if (mesh.dim() > 1)
20  meshInfo.numEdge = mesh.nedges();
21  if (mesh.family() == OMEGA_H_SIMPLEX) {
22  if (mesh.dim() > 1)
23  meshInfo.numTri = mesh.nfaces();
24  if (mesh.dim() == 3)
25  meshInfo.numTet = mesh.nregions();
26  } else { // hypercube
27  if (mesh.dim() > 1)
28  meshInfo.numQuad = mesh.nfaces();
29  if (mesh.dim() == 3)
30  meshInfo.numHex = mesh.nregions();
31  }
32  return meshInfo;
33 }
34 // we use functors because cannot use auto return type with KOKKOS_LAMBDAS
35 template <int dim, typename Field, typename Coords, typename CoordinateMap>
36 struct SetCoordFieldFunctor
37 {
38  SetCoordFieldFunctor(Field& field, Coords& coords, int meshDim,
39  const CoordinateMap& map) :
40  field_(field), coords_(coords), meshDim_(meshDim), map_(map) {}
41  KOKKOS_INLINE_FUNCTION
42  void operator()(const size_t vtx) const {
43  Kokkos::Array<MeshField::Real, dim> x{};
44  for (int d = 0; d < meshDim_; ++d)
45  x[d] = coords_[vtx * meshDim_ + d];
46  const auto mapped = map_(x);
47  for (int d = 0; d < meshDim_; ++d)
48  field_(vtx, 0, d, MeshField::Vertex) = mapped[d];
49  }
50  Field field_;
51  Coords coords_;
52  int meshDim_;
53  const CoordinateMap map_;
54 };
55 
56 
57 template <int dim, typename Element, typename CoordinateMap>
58 struct InterpolateElementFunctor
59 {
60  InterpolateElementFunctor(const Element& element, const CoordinateMap& map)
61  : element_(element), map_(map) {}
62 
63  KOKKOS_INLINE_FUNCTION
64  auto operator()(const int ent, Kokkos::Array<MeshField::Real, dim> const& xi) const {
65  return map_(element_.getValue(ent, xi));
66  }
67 
68  const Element element_;
69  const CoordinateMap map_;
70 };
71 
72 } // anonymous namespace
73 
74 namespace MeshField {
75 
76 namespace Omegah {
77 struct LinearTriangleToVertexField {
78  Omega_h::LOs triVerts;
79  LinearTriangleToVertexField(Omega_h::Mesh &mesh)
80  : triVerts(mesh.ask_elem_verts()) {
81  if (mesh.dim() != 2 && mesh.family() != OMEGA_H_SIMPLEX) {
82  MeshField::fail(
83  "The mesh passed to %s must be 2D and simplex (triangles)\n",
84  __func__);
85  }
86  }
87 
88  static constexpr KOKKOS_FUNCTION Kokkos::Array<MeshField::Mesh_Topology, 1>
89  getTopology() {
90  return {MeshField::Triangle};
91  }
92 
93  KOKKOS_FUNCTION MeshField::ElementToDofHolderMap
94  operator()(MeshField::LO triNodeIdx, MeshField::LO triCompIdx,
95  MeshField::LO tri, MeshField::Mesh_Topology topo) const {
96  assert(topo == MeshField::Triangle);
97  const auto triDim = 2;
98  const auto vtxDim = 0;
99  const auto ignored = -1;
100  const auto localVtxIdx =
101  Omega_h::simplex_down_template(triDim, vtxDim, triNodeIdx, ignored);
102  const auto triToVtxDegree = Omega_h::simplex_degree(triDim, vtxDim);
103  const MeshField::LO vtx = triVerts[(tri * triToVtxDegree) + localVtxIdx];
104  return {0, triCompIdx, vtx, MeshField::Vertex}; // node, comp, ent, topo
105  }
106 };
107 struct LinearTetrahedronToVertexField {
108  Omega_h::LOs tetVerts;
109  LinearTetrahedronToVertexField(Omega_h::Mesh &mesh)
110  : tetVerts(mesh.ask_elem_verts()) {
111  if (mesh.dim() != 3 && mesh.family() != OMEGA_H_SIMPLEX) {
112  MeshField::fail(
113  "The mesh passed to %s must be 3D and simplex (tetrahedron)\n",
114  __func__);
115  }
116  }
117  static constexpr KOKKOS_FUNCTION Kokkos::Array<MeshField::Mesh_Topology, 1>
118  getTopology() {
119  return {MeshField::Tetrahedron};
120  }
121 
122  KOKKOS_FUNCTION MeshField::ElementToDofHolderMap
123  operator()(MeshField::LO tetNodeIdx, MeshField::LO tetCompIdx,
124  MeshField::LO tet, MeshField::Mesh_Topology topo) const {
125  assert(topo == MeshField::Tetrahedron);
126  const auto tetDim = 3;
127  const auto vtxDim = 0;
128  const auto ignored = -1;
129  const auto localVtxIdx = Omega_h::simplex_down_template(tetDim, vtxDim, tetNodeIdx, ignored);
130  const auto tetToVtxDegree = Omega_h::simplex_degree(tetDim, vtxDim);
131  const MeshField::LO vtx = tetVerts[(tet * tetToVtxDegree) + localVtxIdx];
132  return {0, tetCompIdx, vtx, MeshField::Vertex}; // node, comp, ent, topo
133  }
134 };
137  Omega_h::LOs triVerts;
138  Omega_h::LOs triEdges;
139  QuadraticTriangleToField(Omega_h::Mesh &mesh)
140  : triVerts(mesh.ask_elem_verts()),
141  triEdges(mesh.ask_down(mesh.dim(), 1).ab2b) {
142  if (mesh.dim() != 2 && mesh.family() != OMEGA_H_SIMPLEX) {
143  MeshField::fail(
144  "The mesh passed to %s must be 2D and simplex (triangles)\n",
145  __func__);
146  }
147  }
148 
149  static constexpr KOKKOS_FUNCTION Kokkos::Array<MeshField::Mesh_Topology, 1>
150  getTopology() {
151  return {MeshField::Triangle};
152  }
153 
154  KOKKOS_FUNCTION MeshField::ElementToDofHolderMap
155  operator()(MeshField::LO triNodeIdx, MeshField::LO triCompIdx,
156  MeshField::LO tri, MeshField::Mesh_Topology topo) const {
157  assert(topo == MeshField::Triangle);
158  // Omega_h has no concept of nodes so we can define the map from
159  // triNodeIdx to the dof holder index
160  const MeshField::LO triNode2DofHolder[6] = {
161  /*vertices*/ 0, 1, 2,
162  /*edges*/ 0, 1, 2};
163  const MeshField::Mesh_Topology triNode2DofHolderTopo[6] = {
164  /*vertices*/
165  MeshField::Vertex, MeshField::Vertex, MeshField::Vertex,
166  /*edges*/
167  MeshField::Edge, MeshField::Edge, MeshField::Edge};
168  const auto dofHolderIdx = triNode2DofHolder[triNodeIdx];
169  const auto dofHolderTopo = triNode2DofHolderTopo[triNodeIdx];
170  // Given the topo index and type find the Omega_h vertex or edge index that
171  // bounds the triangle
172  Omega_h::LO osh_ent;
173  if (dofHolderTopo == MeshField::Vertex) {
174  const auto triDim = 2;
175  const auto vtxDim = 0;
176  const auto ignored = -1;
177  const auto localVtxIdx = Omega_h::simplex_down_template(
178  triDim, vtxDim, dofHolderIdx, ignored);
179  const auto triToVtxDegree = Omega_h::simplex_degree(triDim, vtxDim);
180  osh_ent = triVerts[(tri * triToVtxDegree) + localVtxIdx];
181  } else if (dofHolderTopo == MeshField::Edge) {
182  const auto triDim = 2;
183  const auto edgeDim = 1;
184  const auto triToEdgeDegree = Omega_h::simplex_degree(triDim, edgeDim);
185  // passing dofHolderIdx as Omega_h_simplex.hpp does not provide
186  // a function that maps a triangle and edge index to a 'canonical' edge
187  // index. This may need to be revisited...
188  osh_ent = triEdges[(tri * triToEdgeDegree) + dofHolderIdx];
189  } else {
190  assert(false);
191  }
192  return {0, triCompIdx, osh_ent, dofHolderTopo};
193  }
194 };
196 
198  Omega_h::LOs tetVerts;
199  Omega_h::LOs tetEdges;
200  QuadraticTetrahedronToField(Omega_h::Mesh &mesh)
201  : tetVerts(mesh.ask_elem_verts()),
202  tetEdges(mesh.ask_down(mesh.dim(), 1).ab2b) {
203  if (mesh.dim() != 3 && mesh.family() != OMEGA_H_SIMPLEX) {
204  MeshField::fail(
205  "The mesh passed to %s must be 3D and simplex (tetrahedron)",
206  __func__);
207  }
208  }
209 
210  static constexpr KOKKOS_FUNCTION Kokkos::Array<MeshField::Mesh_Topology, 1>
211  getTopology() {
212  return {MeshField::Tetrahedron};
213  }
214 
215  KOKKOS_FUNCTION MeshField::ElementToDofHolderMap
216  operator()(MeshField::LO tetNodeIdx, MeshField::LO tetCompIdx,
217  MeshField::LO tet, MeshField::Mesh_Topology topo) const {
218  assert(topo == MeshField::Tetrahedron);
219  const MeshField::LO tetNode2DofHolder[10] = {0, 1, 2, 3, 0, 2, 3, 1, 5, 4};
220  const MeshField::Mesh_Topology tetNode2DofHolderTopo[10] = {
221  MeshField::Vertex, MeshField::Vertex, MeshField::Vertex,
222  MeshField::Vertex, MeshField::Edge, MeshField::Edge,
223  MeshField::Edge, MeshField::Edge, MeshField::Edge,
224  MeshField::Edge};
225  const auto dofHolderIdx = tetNode2DofHolder[tetNodeIdx];
226  const auto dofHolderTopo = tetNode2DofHolderTopo[tetNodeIdx];
227  Omega_h::LO osh_ent;
228  if (dofHolderTopo == MeshField::Vertex) {
229  const auto tetDim = 3;
230  const auto vtxDim = 0;
231  const auto ignored = -1;
232  const auto localVtxIdx = Omega_h::simplex_down_template(
233  tetDim, vtxDim, dofHolderIdx, ignored);
234  const auto tetToVtxDegree = Omega_h::simplex_degree(tetDim, vtxDim);
235  osh_ent = tetVerts[(tet * tetToVtxDegree) + localVtxIdx];
236  } else if (dofHolderTopo == MeshField::Edge) {
237  const auto tetDim = 3;
238  const auto edgeDim = 1;
239  const auto tetToEdgeDegree = Omega_h::simplex_degree(tetDim, edgeDim);
240  osh_ent = tetEdges[(tet * tetToEdgeDegree) + dofHolderIdx];
241  } else {
242  assert(false);
243  }
244  return {0, tetCompIdx, osh_ent, dofHolderTopo};
245  }
246 };
247 
249 template <int ShapeOrder> auto getTriangleElement(Omega_h::Mesh &mesh) {
250  static_assert(ShapeOrder == 1 || ShapeOrder == 2);
251  if constexpr (ShapeOrder == 1) {
252  struct result {
254  LinearTriangleToVertexField map;
255  };
256  return result{MeshField::LinearTriangleShape(),
257  LinearTriangleToVertexField(mesh)};
258  } else if constexpr (ShapeOrder == 2) {
259  struct result {
261  QuadraticTriangleToField map;
262  };
263  return result{MeshField::QuadraticTriangleShape(),
264  QuadraticTriangleToField(mesh)};
265  }
266 }
268 template <int ShapeOrder> auto getTetrahedronElement(Omega_h::Mesh &mesh) {
269  static_assert(ShapeOrder == 1 || ShapeOrder == 2);
270  if constexpr (ShapeOrder == 1) {
271  struct result {
273  LinearTetrahedronToVertexField map;
274  };
275  return result{MeshField::LinearTetrahedronShape(),
276  LinearTetrahedronToVertexField(mesh)};
277  } else if constexpr (ShapeOrder == 2) {
278  struct result {
280  QuadraticTetrahedronToField map;
281  };
282  return result{MeshField::QuadraticTetrahedronShape(),
283  QuadraticTetrahedronToField(mesh)};
284  }
285 }
286 
287 } // namespace Omegah
288 
293  template <typename Point>
294  KOKKOS_INLINE_FUNCTION Point operator()(Point const &x) const {
295  return x;
296  }
297 };
298 
299 template <typename ExecutionSpace, size_t dim,
300  template <typename...> typename Controller =
301  MeshField::KokkosController>
302 class OmegahMeshField {
303 private:
304  Omega_h::Mesh &mesh;
305  const MeshField::MeshInfo meshInfo;
306  using CoordField =
307  decltype(MeshField::CreateCoordinateField<ExecutionSpace, Controller,
308  dim>(MeshField::MeshInfo()));
309  CoordField coordField;
310 
311 public:
312  OmegahMeshField(Omega_h::Mesh &mesh_in)
313  : mesh(mesh_in), meshInfo(getMeshInfo(mesh)),
314  coordField(CreateLagrangeCoordinateField<1>()) {
315  static_assert(dim == 1 || dim == 2 || dim == 3);
316  }
317 
318  template <typename DataType, size_t order, size_t numComp>
319  // Ordering of field indexing changed to 'entity, node, component'
320  auto CreateLagrangeField() const {
321  return MeshField::CreateLagrangeField<ExecutionSpace, Controller, DataType,
322  order, dim, numComp>(meshInfo);
323  }
324 
340  template <size_t order, typename CoordinateMap = IdentityCoordinateMap>
341  auto CreateLagrangeCoordinateField(
342  const CoordinateMap &map = CoordinateMap()) const {
343  auto fieldWithCtrlr = CreateLagrangeField<Real, order, dim>();
344  auto field = fieldWithCtrlr.field;
345  // note order 1 is called in the constructor to initialize the coordinate field
346  // it must be called before higher order coordinate fields are constructed
347  if constexpr (order == 1) {
348  const auto meshDim = meshInfo.dim;
349  const auto coords = mesh.coords();
350  MeshField::parallel_for<decltype(fieldWithCtrlr.ctrlr)>(
351  {0}, {meshInfo.numVtx},
352  SetCoordFieldFunctor<dim, decltype(field), decltype(coords),
353  CoordinateMap>(field, coords, meshDim, map),
354  "setCoordField");
355  } else {
356  const auto linear = CreateLagrangeElement(coordField.field);
357  auto target = CreateLagrangeElement(field);
358  MeshField::setDofByInterpolation(
359  target,
360  InterpolateElementFunctor<dim, decltype(linear), CoordinateMap>(
361  linear, map));
362  }
363  return fieldWithCtrlr;
364  }
365 
366  auto getCoordField() { return coordField; }
367 
372  template <typename Field>
373  auto CreateLagrangeElement(const Field &field) const {
374  static_assert(dim == 2 || dim == 3,
375  "CreateLagrangeElement supports triangles and tetrahedra");
376  constexpr int order = Field::Order;
377  if constexpr (dim == 2) {
378  const auto [shp, map] = Omegah::getTriangleElement<order>(mesh);
379  return MeshField::FieldElement(meshInfo.numTri, field, shp, map);
380  } else {
381  const auto [shp, map] = Omegah::getTetrahedronElement<order>(mesh);
382  return MeshField::FieldElement(meshInfo.numTet, field, shp, map);
383  }
384  }
385 
386  // FIXME support 2d and 3d and fields with order>1
387  template <typename Field> void writeVtk(Field &field) const {
388  using FieldDataType = typename decltype(field.vtxField)::BaseType;
389  // HACK assumes there is a vertex field.. in the Field Mixin object
390  auto field_view = field.vtxField.serialize();
391  Omega_h::Write<FieldDataType> field_write(field_view);
392  mesh.add_tag(0, "field", 1, Omega_h::read(field_write), false,
393  Omega_h::ArrayType::VectorND);
394  Omega_h::vtk::write_parallel("foo.vtk", &mesh, mesh.dim());
395  }
396 
397  template <typename ViewType = Kokkos::View<MeshField::LO *>>
398  ViewType createOffsets(size_t numTri, size_t numPtsPerElem) const {
399  ViewType offsets("offsets", numTri + 1);
400  Kokkos::parallel_for(
401  "setOffsets", numTri,
402  KOKKOS_LAMBDA(int i) { offsets(i) = i * numPtsPerElem; });
403  Kokkos::deep_copy(Kokkos::subview(offsets, offsets.size() - 1),
404  numTri * numPtsPerElem);
405  return offsets;
406  }
407 
408  // evaluate a field at the specified local coordinate for each triangle
409  template <typename ViewType, typename ShapeField>
410  auto triangleLocalPointEval(const ViewType &localCoords, size_t NumPtsPerElem,
411  const ShapeField &field) const {
412  auto offsets = createOffsets(meshInfo.numTri, NumPtsPerElem);
413  auto eval = triangleLocalPointEval<ViewType, ShapeField>(localCoords,
414  offsets, field);
415  return eval;
416  }
417 
418  // evaluate a field at the specified local coordinates for each triangle
419  template <typename ViewType, typename ShapeField>
420  auto triangleLocalPointEval(const ViewType &localCoords,
421  Kokkos::View<LO *> offsets,
422  const ShapeField &field) const {
423  const auto MeshDim = 2;
424  if (mesh.dim() != MeshDim) {
425  MeshField::fail("input mesh must be 2d\n");
426  }
427  const auto ShapeOrder = ShapeField::Order;
428  if (ShapeOrder != 1 && ShapeOrder != 2) {
429  MeshField::fail("input field order must be 1 or 2\n");
430  }
431 
432  const auto [shp, map] = Omegah::getTriangleElement<ShapeOrder>(mesh);
433 
434  MeshField::FieldElement<ShapeField, decltype(shp), decltype(map)> f(
435  meshInfo.numTri, field, shp, map);
436  auto eval = MeshField::evaluate(f, localCoords, offsets);
437  return eval;
438  }
439 
440  template <typename ViewType, typename ShapeField>
441  auto tetrahedronLocalPointEval(const ViewType &localCoords,
442  size_t NumPtsPerElem,
443  const ShapeField &field) const {
444  auto offsets = createOffsets(meshInfo.numTet, NumPtsPerElem);
445  auto eval = tetrahedronLocalPointEval(localCoords, offsets, field);
446  return eval;
447  }
448 
449  template <typename ViewType, typename ShapeField>
450  auto tetrahedronLocalPointEval(const ViewType &localCoords,
451  Kokkos::View<LO *> offsets,
452  const ShapeField &field) const {
453  const auto MeshDim = 3;
454  if (mesh.dim() != MeshDim) {
455  MeshField::fail("input mesh must be 3d\n");
456  }
457  const auto ShapeOrder = ShapeField::Order;
458  if (ShapeOrder != 1 && ShapeOrder != 2) {
459  MeshField::fail("input field order must be 1 or 2\n");
460  }
461  const auto [shp, map] = Omegah::getTetrahedronElement<ShapeOrder>(mesh);
462  MeshField::FieldElement<ShapeField, decltype(shp), decltype(map)> f(
463  meshInfo.numTet, field, shp, map);
464  auto eval = MeshField::evaluate(f, localCoords, offsets);
465  return eval;
466  }
467 };
468 
469 } // namespace MeshField
470 
471 #endif
Supports mapping between mesh (i.e., Omega_h) ordering and MeshFields ordering.
Supports the evaluation of a field, and other per-element operations, given the definition of the ele...
the CoordinateMap that leaves the mesh geometry unchanged
Definition: MeshField.hpp:292
Linear (P1) shape functions for 2D triangular elements.
On-process mesh metadata.
Quadratic (P2) shape functions for 3D tetrahedral elements.
Quadratic (P2) shape functions for 2D triangular elements.