DCEL: rewrite MeshT to drop its owned Pool pointer for device-mirror safety

[rmrsk]

Summary

Background

This is PR 4 of the GPU-port DCEL migration, stacked on top of #139 (Pool-backed PODVector
storage). #139 gave MeshT a Pool* cached at construction and dereferenced from every query --
correct on the host, but fundamentally unsafe for the eventual device-mirror goal: a pointer
cached in one address space is meaningless once the same bytes are mirrored into another (and
Pool::mirror()'s destination Pool object is itself still host-resident even when its block
is device memory), so MeshT could never actually be copied to a device and remain usable.

Solution

• MeshT no longer stores a Pool*. It is now default-constructible and stores only a resolved
  void* m_base, set once via the new bind(const Pool&) (asserts the Pool is frozen).
• Every method that resolves a vertex/edge/face index now has two overloads: an explicit-base
  form (e.g. getVertex(void* a_base, uint32_t), valid at any point including mid-build, before a
  Pool is frozen) and a no-argument, bound form (e.g. getVertex(uint32_t)) that resolves
  against m_base for convenient querying once building is finished. This mirrors the
  PODVector::at()-vs-bind() convention already established in the memory foundation.
  Build-phase mutators (reserveVertices/Edges/Faces, addVertex/Edge/Face) take Pool&
  explicitly, matching PODVector::reserveFrom/push_back.
• Internally, reconcile()/sanityCheck()/the direct signed-distance queries delegate to
  FaceT/EdgeT/VertexT methods that take a const Mesh& and always resolve via the mesh's
  bound accessors -- unusable pre-bind. Rather than adding base parameters to those classes'
  public interfaces (which would leak memory-semantics concerns into types that shouldn't need
  them), MeshT gained a small boundView(const void* a_base) helper: a disposable, stack-only,
  trivially-copyable view of the same mesh data bound to an explicit base, safe to pass wherever a
  bound const Mesh& is required before the real mesh can be bound.
• Soup::soupToDCEL/reconcilePairEdgesDCEL, the STL/PLY/OBJ/VTK convertToDCEL entry
  points, and Parser::readIntoDCEL build purely against an explicit Pool& and never freeze it
  -- the Pool may still be shared with more files/meshes.
• FlatMeshSDF/MeshSDF's constructors are the actual freeze+bind boundary: retaining a mesh for
  long-term querying is the point at which building must be considered finished, so their
  constructors now take Pool& and freeze (idempotent) + bind it. TriMeshSDF's mesh-based
  constructor deliberately does neither -- it only reads the mesh once to extract flat triangles
  and does not retain it, so forcing a freeze would needlessly block any other mesh still being
  built into a shared Pool.
• Parser's multi-file readIntoMesh/readIntoPackedBVH now build every file's mesh before
  wrapping any of them (rather than looping the single-file overload), since wrapping the first
  file would otherwise freeze the shared Pool out from under the rest.
• Adds a new Sphinx "Memory model" page documenting Pool/PODVector/MemoryResource in
  general, and a matching "Memory model" subsection on ImplemDCEL.rst for how MeshT is built
  on top of it specifically, including the sharing/freezing pitfalls this change introduces.
  Updates Parsers.rst for the new soupToDCEL signature and the freeze/pool-sharing caveats.

Side-effects

• Source-breaking API change for anyone calling MeshT's constructor, Soup::soupToDCEL, or
  FlatMeshSDF/MeshSDF/TriMeshSDF's mesh-based constructors directly: MeshT is now
  default-constructed (no Pool at construction); soupToDCEL takes an explicit Pool&;
  FlatMeshSDF/MeshSDF/TriMeshSDF's mesh-based constructors take an additional Pool&.
  Parser::readInto*'s own public signatures ((filename, pool, ...)) are unchanged.
• A Pool shared across multiple FlatMeshSDF/MeshSDF builds must have all of its meshes
  built before the first one is wrapped (documented as a new warning in the Memory model docs);
  Examples/MeshSDF previously built three independent representations from one shared Pool and
  needed a genuine fix (three separate Pools) rather than a mechanical signature update.
• No behavior change to signed-distance results, mesh topology, or any other observable output:
  validated by the full unit test suite (double and float), all Examples via ctest, and the
  debug-san (ASan/UBSan) preset on every DCEL-related test binary -- all clean.

Alternative solutions

• Considered keeping MeshT::reconcile()/sanityCheck() working by inlining FaceT/EdgeT/
  VertexT's topology-walking and geometry-computation logic directly into MeshT, avoiding the
  boundView() helper entirely. Rejected: it would duplicate real logic that already lives
  correctly on those classes, for no benefit over a small, clearly-scoped, disposable stack-local
  view.
• Considered freezing a Pool automatically inside Parser::readIntoDCEL itself (so a bare parsed
  mesh is always immediately queryable). Rejected: readIntoDCEL's own multi-file overload
  depends on the ability to keep building into a still-open Pool across several files, which an
  automatic freeze on the first file would break; freeze+bind is deferred to the SDF-wrapper layer
  instead, where "done building" is unambiguous.

Reviewer checklist (to be completed by a human)

• The test suite compiles and runs to completion without warnings or errors.
• All relevant new features are documented in the user documentation (Sphinx).
• This contribution does not break existing sections in the user documentation.
• All relevant APIs are documented in the doxygen documentation.
• Appropriate labels have been assigned to this PR.
• New or revised proper licensing and copyright information is in place.
• A PR review has been run using @claude review.
• The continuous integration and testing hooks at GitHub run to completion.

---

Generated by Claude Code

[rmrsk]

@claude review

---

Generated by Claude Code

[github-actions]

Claude finished @rmrsk's task in 3m 57s —— View job

---

Claude finished reviewing PR #140

Tasks

• Gather context (PR body, changed files list)
• Fetch base branch and diff
• Read changed source files (EBGeometry_DCEL_Mesh.hpp, MeshImplem.hpp, MeshDistanceFunctions*, Soup*, parsers, format converters)
• Read changed test files (TestDCEL.cpp, TestBVH.cpp, TestPolygon2D.cpp, TestOBJ/PLY/STL/VTK.cpp)
• Read new/changed docs (MemoryModel.rst, ImplemDCEL.rst, Parsers.rst, index.rst, Implementation.rst)
• Post review feedback

(Note: could not run a local build/test — cmake/g++ invocations required approval that wasn't available in this non-interactive review context, so this is a careful read-through rather than a compiled/executed verification.)

Summary

This is a clean, well-executed rewrite. The explicit-base vs. bound-accessor split is applied consistently across MeshT, Soup, the format converters, Parser, and FlatMeshSDF/MeshSDF/TriMeshSDF; the boundView() bridge is a reasonable way to avoid leaking memory-semantics parameters into VertexT/EdgeT/FaceT. Tests were updated correctly throughout — freeze()/bind() calls are consistently inserted right before the first no-argument accessor use, and pools that must stay separate (because a wrapper's constructor now freezes them) were correctly split into multiple Pool instances. The new MemoryModel.rst page and the ImplemDCEL.rst/Parsers.rst updates accurately describe the new contract, and I found no stale references to the old MeshT(Pool&) constructor or the old soupToDCEL(mesh, vertices, facets, id) signature anywhere in the docs.

Finding: Integrations/AMReX/PaintEB/main.cpp is broken by this API change

This file is not CI-tested (per CLAUDE.md), so nothing caught it automatically, but it's a real regression introduced by this PR's source-breaking changes:

• Source/EBGeometry_MeshDistanceFunctions.hpp:200: MeshSDF's constructor now takes (mesh, pool, build). Integrations/AMReX/PaintEB/main.cpp:49 still calls the old two-argument form:

  m_sdf = std::make_shared<EBGeometry::MeshSDF<T, Meta, K>>(mesh, EBGeometry::BVH::Build::SAH);

  This will fail to compile.

• Integrations/AMReX/PaintEB/main.cpp:46: mesh->getFace(i).getMetaData() = 1.0 * i; uses the no-argument (bound) accessor on a mesh straight out of readIntoDCEL. Per the new contract (correctly documented in ImplemDCEL.rst's new "Memory model" section), readIntoDCEL deliberately never freezes/binds the pool, so mesh->m_base is still nullptr here — even after fixing the constructor call above, this line would be undefined behavior (a null-pointer-offset dereference), not a compile error.

Both would need fixing together, e.g.:

auto mesh = EBGeometry::Parser::readIntoDCEL<T, Meta>(a_filename, *m_pool);

m_pool->freeze();
mesh->bind(*m_pool);

for (uint32_t i = 0; i < mesh->numFaces(); i++) {
  mesh->getFace(i).getMetaData() = 1.0 * i;
}

m_sdf = std::make_shared<EBGeometry::MeshSDF<T, Meta, K>>(mesh, *m_pool, EBGeometry::BVH::Build::SAH);

(freezing before the metadata-mutation loop is fine here since MeshSDF's own constructor freeze is idempotent, and nothing else needs to build into *m_pool afterward).

The sibling files Integrations/AMReX/MeshSDF/main.cpp and Integrations/Chombo/MeshSDF/main.cpp only call the public Parser::readIntoMesh/readIntoTriangleBVH entry points, whose signatures are explicitly unchanged by this PR, so those two are fine.

Fix this →

Minor / non-blocking observations

• Source/EBGeometry_DCEL_MeshImplem.hpp:676-686 (DirectSignedDistance2): closestIndex is initialized to 0 and the loop below also starts at i = 0, so the first iteration redundantly re-computes face 0's distance against itself (harmless, pre-existing pattern, not introduced by this PR — just noting it since the surrounding code was touched).
• Docs/Sphinx/source/MemoryModel.rst and the Sec:DCELMemoryModel subsection are thorough and clearly cross-referenced; no literalinclude usage introduced, consistent with the repo-wide ban.
• Tests/InstantiateAll.cpp was correctly left untouched — MeshT isn't a new class, and the reused-meshPool-across-multiple-readInto*-calls pattern there is harmless because that file is compile-only (never executed), so the runtime freeze/reserve assertion this would otherwise trip never fires.

[rmrsk]

Good catch — fixed in 5655310. Integrations/ isn't CI-built (no AMReX dependency here), so this was a read-through fix mirroring the same freeze+bind pattern applied everywhere else in this PR, not a compiled/executed verification. The minor DirectSignedDistance2 observation is pre-existing and out of scope for this PR.

---

Generated by Claude Code

[rmrsk]

@claude review

---

Generated by Claude Code

[github-actions]

Claude finished @rmrsk's task in 5m 40s —— View job

---

Claude's review of PR #140

Tasks

• Gather context (PR body, prior review thread, changed files list)
• Identify what's new since the last @claude review pass (commits 6746443, e9c3635 — GPU host/device annotations on VertexT/EdgeT/FaceT/EdgeIteratorT/MeshT; commit 5655310, the PaintEB fix, was already reviewed/confirmed)
• Read the new diff (Source/EBGeometry_DCEL_{Mesh,Edge,Face,Vertex,Iterator}{,Implem}.hpp, Tests/TestDCEL.cpp, Tests/CMakeLists.txt)
• Cross-check every EBGEOMETRY_HOST/EBGEOMETRY_HOST_DEVICE declaration against its out-of-line definition, and against what it actually touches (Pool, std::vector, std::cerr)
• Post review feedback

Summary

This second pass only reviews what changed since my last comment: the previously-flagged Integrations/AMReX/PaintEB breakage is fixed in 5655310 and confirmed by @rmrsk as a read-through fix (untested, since Integrations/ has no AMReX dependency in CI) — no new issue there. The new material is 6746443/e9c3635, which annotate the whole DCEL query surface (VertexT, EdgeT, FaceT, EdgeIteratorT, MeshT) with the EBGEOMETRY_HOST/EBGEOMETRY_HOST_DEVICE macros from EBGeometry_GPU.hpp, plus a new [gpu]-tagged TestDCEL.cpp case that launches a real kernel exercising MeshT::boundView()/getFace()/getVertex()/unsignedDistance2()/EdgeIteratorT and compares it against the host result.

I went through this annotation work function-by-function rather than spot-checking:

• Every .hpp declaration's macro matches its *Implem.hpp out-of-line definition — I diffed occurrence counts per class (Mesh, Edge, Face, Vertex, Iterator) and manually reconciled every count mismatch back to a class-level doc comment mentioning the macro by name (not a missing/extra annotation on an actual declaration).
• Every function left as EBGEOMETRY_HOST genuinely needs to be: it either takes/touches a Pool& directly (bind, deepCopy, reserve*, add*), returns/builds a std::vector (getAllVertexCoordinates, gatherVertexIndices, gatherEdgeIndices, getSmallestCoordinate/getHighestCoordinate, computeCentroid/computeNormal/computeArea and thus Face::reconcile), or can reach std::cerr on a diagnostic path (sanityCheck, incrementWarning/printWarnings, computeVertexNormalAngleWeighted, the enum-dispatching signedDistance overloads, and therefore Mesh::reconcile/reconcileFaces/reconcileVertices, which call into those). None of the EBGEOMETRY_HOST_DEVICE-annotated functions reach any of these paths.
• FaceT::computeWindingNumber's local isLeft lambda was inlined into the cross-product expression (EBGeometry_DCEL_FaceImplem.hpp:433) specifically so the function stays callable from device code without --extended-lambda. I checked the algebra: (P2.x-P1.x)*(a_point.y-P1.y) - (a_point.x-P1.x)*(P2.y-P1.y) is exactly isLeft(P1, P2, a_point) under the old definition — a pure refactor, no behavior change. I also grepped the rest of the newly-annotated *Implem.hpp files for other lambdas/std::vector/std::cerr inside EBGEOMETRY_HOST_DEVICE bodies and found none.
• Defaulted special member functions (copy/move ctors, destructor, copy/move assignment on VertexT/EdgeT/FaceT/EdgeIteratorT) are left unannotated — consistent with the pre-existing convention on every other GPU-ported value type in this codebase (Vec, PODVector, PointSoA/AoSoA, TriangleSoA/AoSoA, BoundingVolumes), so not a gap introduced here.
• New device test (Tests/TestDCEL.cpp, gated #if defined(EBGEOMETRY_CUDA) || defined(EBGEOMETRY_HIP), skips cleanly via SKIP("no GPU device available") otherwise): reuses the existing buildTetrahedron() helper correctly (it already freezes+binds the source Pool), mirrors it with the pre-existing Pool::mirror(), and passes mesh->boundView(devicePool.base()) into the kernel by value — consistent with MeshT being trivially copyable. Tests/CMakeLists.txt's EBGEOMETRY_GPU_TESTS list is updated to include TestDCEL alongside the other GPU-tested files. I did not have cmake/nvcc/hipcc available in this review environment to actually compile the device path, so this is a careful read-through, not a compiled verification (same caveat as my first pass on this PR).
• No Sphinx doc changes were needed for this increment: no Sphinx page in this repo currently documents the EBGEOMETRY_HOST/EBGEOMETRY_HOST_DEVICE macros or their per-class distribution (not even for the classes annotated in earlier PRs in this stack), so this is consistent with existing practice, not a new gap.

No issues found in this increment.
· branch dcel-mesh-explicit-base

[rmrsk]

@claude review

---

Generated by Claude Code

[github-actions]

Claude finished @rmrsk's task in 2m 26s —— View job

---

Claude's review of PR #140 (third pass)

Tasks

• Gather context (PR body, prior review threads, changed files list)
• Identify what's new since the last @claude review pass (last reviewed commit: e9c3635)
• Read the new diff (e9c3635..HEAD): CI trigger widening (dc066ec), HIP std::clamp fix (0292ded), and the "adversarial device-mirror-safety" commit (2689747)
• Cross-check the five claimed fixes in 2689747 against the actual code
• Post review feedback

(No network/build access in this run either — this is a careful read-through, not a compiled/executed verification, same caveat as both prior passes on this PR.)

Summary

This pass only covers what changed since my last comment. The previously-flagged Integrations/AMReX/PaintEB issue is already fixed and confirmed. The new commits are:

1. dc066ec — widens CI's pull_request trigger to [main, dev] (explaining why this whole PR stack showed 0 check runs before) and wires ctest -L gpu-device into the GPU-CUDA/GPU-HIP jobs via a second catch_discover_tests(... TEST_SPEC "[gpu]" TEST_PREFIX "gpu:") registration in Tests/CMakeLists.txt. The "gpu:" prefix keeps these from colliding with the same file's "unit"-labeled registration under ebgeometry_add_test(). Correct.
2. 0292ded — replaces std::clamp in EdgeT::unsignedDistance2 with a hand-rolled ternary clamp, matching Vec3T::clamp()'s existing rationale (libstdc++ hardened-mode std::clamp pulls in a __host__-only assert helper that HIP's device compiler rejects from __host__ __device__ code). I checked every other std::clamp call site in the codebase (EBGeometry_CSGImplem.hpp:820, EBGeometry_AnalyticDistanceFunctions.hpp:980,1212-1213, EBGeometry_DCEL_VertexImplem.hpp:258) — none of them sit inside an EBGEOMETRY_HOST_DEVICE-annotated function (the CSG/AnalyticDistanceFunctions files have no GPU annotations at all yet, and the VertexImplem.hpp one is inside computeVertexNormalAngleWeighted, which stays EBGEOMETRY_HOST because it also calls std::cerr). The commit's "only reachable one" claim holds.
3. 2689747 — addresses 5 self-identified adversarial findings; I verified each against the diff rather than trusting the commit message:
   • MeshT::signedDistance()'s 4 public overloads are promoted EBGEOMETRY_HOST → EBGEOMETRY_HOST_DEVICE, and the algorithm-dispatch default: branch's std::cerr call is replaced with EBGEOMETRY_EXPECT(false) alone. Confirmed minDist still has a sane default (std::numeric_limits<T>::max(), initialized before the switch) on that corrupted-enum path, and EBGEOMETRY_EXPECT already had a device-safe (assert()-based) expansion under EBGEOMETRY_DEVICE_COMPILE from EBGeometry_Macros.hpp — nothing new needed there.
   • static_assert(std::is_trivially_copyable_v<Meta>, ...) added to VertexT/EdgeT/FaceT, mirroring the existing PointAoSoA/TriangleAoSoA convention. Straightforward.
   • EdgeIteratorT's copy constructor changed from = default to a user-provided memberwise copy, specifically to make std::is_trivially_copyable_v false (per-standard, one non-trivial copy/move ctor or assignment is sufficient to disqualify the whole type, regardless of the move ctor/copy-assignment/move-assignment still being = default) so PODVector<EdgeIteratorT<...>> fails to compile instead of silently mirroring a dangling host pointer. The two new static_assert(!std::is_trivially_copyable_v<...>) lines confirm the intended effect. Correct.
   • boundView()'s return type changed Mesh → const Mesh, which does block a chained mutating call on the temporary (e.g. mesh.boundView(base).flip() no longer compiles); the doc comment correctly warns that copy-initializing into a non-const local (Mesh view = mesh.boundView(base);) still produces an independent, fully mutable copy — this is a real, inherent limit of the approach (not a gap in this PR), and it's called out rather than left implicit.
   • The new [gpu] kernel/host test additions in Tests/TestDCEL.cpp (flip() self-consistency, WindingNumber/SubtendedAngle branches, the newly-public device signedDistance()) mirror each other correctly between the kernel and host-side expectation, including the comment explaining why the device pool is mirrored before the host-side mutations (setInsideOutsideAlgorithm, flip()) that both sides then replay independently from the same starting snapshot.

No issues found in this increment — the fixes are correctly targeted and match what they claim to do.
· branch dcel-mesh-explicit-base