Skip to content

Design: Adaptive Mesh Refinement (AMR) for 3D Models — Octrees + Face Mortars#

Status: Implemented (Stages 1–6a, plus Stage-5 v2 point-to-point solution migration) — mirrors the implemented 2D design. Stages 1–5 landed with #164; the device-resident solution transfer (Stage 6a's 3D analogue) landed with #165; point-to-point migration, replacing the allgather-based redistribution in both 2D and 3D, landed with #167. What is still outstanding from §7 is noted in that table. Sections written in the future tense below have been updated where the delivered code differs from the plan; where it does not, the plan text is the record of what was built. Scope: 3D hexahedral meshes, octree (2:1) h-refinement built on new 3D 2:1 face-mortar interface support, spectral (Legendre modal) refinement indicator, CPU and GPU backends, MPI-parallel. Pilot model: LinearEuler3D.

Companion documents: Mortar2D-Design.md and AMR2D-Design.md. The 2D stack (mortars → quadtree forest → transfer plan → controller → MPI → GPU) is implemented and validated; this document records only the 3D-specific decisions. Where a 3D module is a mechanical transcription of its 2D counterpart, the 2D document remains the authoritative rationale.

1. Structural mapping from 2D#

2D (implemented) 3D (this plan)
Mesh2D_t%mortarInfo(1:8,:) (1 big + 2 small edges) Mesh3D_t%mortarInfo(1:14,:) (1 big + 4 small faces)
MortarExchange / MortarFluxCollect on 2D mapped data same on MappedScalar3D_t / MappedVector3D_t
1D mortarR/mortarP applied per edge same 1D matrices applied as 2D tensor products per face
flip ∈ {0,1}; sub-edge pairing t vs 3-t flip ∈ 0..7 (dihedral group); quadrant permutation table
flux collect factor −2 (sub-edge Jacobian ½) factor −4 (sub-face Jacobian ¼)
SELF_QuadTreeMesh_2D (4 children) SELF_OctreeMesh_3D (8 children)
childOfSide(1:2,1:4) childOfFace(1:4,1:6)
RefineConnectivity 4 inner face pairs 12 inner face pairs
transfer: 2-pass tensor product, 4 children, ¼ 3-pass tensor product, 8 children, ⅛
indicator: 2D modal transform  = P U Pᵀ 3D modal transform (three passes), same criteria
mortarBuff(N+1,1:4,m,v) mortarBuff(N+1,N+1,1:8,m,v)
dt = dtBase/2**MaxLevel identical

New modules (each mirrors the 2D file of the same name, 2D→3D): SELF_OctreeMesh_3D.f90, SELF_AdaptiveMesh_3D.f90, SELF_MeshRefinement_3D.f90, SELF_RefinementPrimitives_3D.f90, SELF_RefinementIndicator_3D_t.f90 (+ cpu/gpu backends), SELF_SolutionTransfer_3D.f90, SELF_TransferPlan_3D.f90, SELF_AMRController_3D.f90.

2. Conventions (normative)#

2.1 Face trace coordinates#

boundary(i,j,side,elem) face coordinates, from BoundaryInterp:

sides face (i,j)
1 (Bottom), 6 (Top) (ξ₁, ξ₂)
2 (South), 4 (North) (ξ₁, ξ₃)
3 (East), 5 (West) (ξ₂, ξ₃)

2.2 Children#

Children are ordered by CGNS corner-node order: axc = [0,1,1,0,0,1,1,0], ayc = [0,0,1,1,0,0,1,1], azc = [0,0,0,0,1,1,1,1]; child c covers ξ_d ∈ [a-1, a] per direction. Subdivision is orientation-preserving: a child face lying on parent face s is that child's local face s, so neighbor/flip data inherits verbatim (same load-bearing fact as in 2D).

2.3 Face quadrants and childOfFace#

A big face is split into 4 sub-faces ("quadrants") indexed in the big face's trace coordinates: q = kx + 2*(ky-1) with kx, ky ∈ {1,2} the half-interval index along the face's i and j coordinate (1 = [-1,0], 2 = [0,1]). Restriction to quadrant q is the tensor product R_kx ⊗ R_ky of the existing 1D mortarR; projection back is P_kx ⊗ P_ky of mortarP.

childOfFace(q,s) — the child of a refined hex whose face covers quadrant q of parent face s (derived from §2.1 + §2.2):

s face childOfFace(1:4,s)
1 Bottom 1, 2, 4, 3
2 South 1, 2, 5, 6
3 East 2, 3, 6, 7
4 North 4, 3, 8, 7
5 West 1, 4, 5, 8
6 Top 5, 6, 8, 7

Octree helpers: oc_opposite = [6,4,5,2,3,1]; oc_reflect(s,c): s∈{1,6} → [5,6,7,8,1,2,3,4], s∈{2,4} → [4,3,2,1,8,7,6,5], s∈{3,5} → [2,1,4,3,6,5,8,7]; oc_subpos(c,s) is the inverse of childOfFace.

2.4 Flips and quadrant permutation#

3D flips follow the existing 8-state convention (SELF_Mesh_3D_t.f90): (i2,j2) = F_f(i1,j1) maps receiver-face indices to donor-face indices. The same maps applied to half-interval indices give the quadrant permutation mortarQuadPerm(q,f) — the donor-face quadrant that coincides with receiver-face quadrant q:

f perm(1:4) f perm(1:4)
0 1,2,3,4 4 1,3,2,4
1 2,1,4,3 5 2,4,1,3
2 4,3,2,1 6 4,2,3,1
3 3,4,1,2 7 3,1,4,2

On octree-emitted meshes, mortars interior to a root tree always carry flip 0; mortars across a root face inherit the base mesh's face flip. The mortar exchange implements all 8 flips so unstructured (HOHQMesh/HOPr) base meshes are supported.

2.5 mortarInfo(1:14, 1:nMortars)#

Replicated on all ranks; element ids are global. Rows:

 1    big element id            2    big local face id
 3,4  small elem, 10*face+flip  (sub-face covering big-face quadrant 1)
 5,6  same for quadrant 2
 7,8  same for quadrant 3
 9,10 same for quadrant 4
11:14 global side ids of sub-faces 1..4 (MPI message tags)

As in 2D: mortar faces carry sideInfo(3) = 0 and sideInfo(5) = 0 (conforming machinery and BC mapping skip them) and sideInfo(1) = mortar index; the flip stored with each small face maps big-face quadrant coordinates to the small face's own coordinates (the same receiver→donor convention as sideInfo(4)).

2.6 Mortar data plane#

mortarBuff(1:N+1, 1:N+1, 1:8, 1:nMortars, 1:nvar [,1:3]) — slots 1–4 hold the big-face trace (replicated per quadrant; MPI receives land per-quadrant), slots 5–8 hold the small-face traces re-oriented into the big face's coordinates. Exchange algorithm, per mortar:

  1. Stage rank-local big trace into slots 1–4 and (flip-reoriented) small traces into slots 4+q; MPI posts fill remote slots (small ranks receive the big trace into slot q; the big rank receives small traces into 4+q), with received small traces flip-reoriented after the wait.
  2. Small extBoundary(i,j) at face coords F_f(i,j) ← (R_kx ⊗ R_ky) applied to slot q (exact restriction).
  3. Big extBoundary ← Σ_q (P_kx ⊗ P_ky) applied to slots 4+q (L2 projection; mortarP carries ½ per direction ⇒ ¼ total, the correct solution-trace projection).

MortarFluxCollect (on boundaryNormal): big-face integrand := −4 Σ_q (P_kx ⊗ P_ky) g_q. The ¼ from the tensor-product mortarP combines with the ×4 integrand conversion so that the big face's discrete surface integral equals minus the sum of the four small faces' integrals to roundoff — the same conservation argument as 2D with ½/×2.

MPI messages use tags globalSideId + nUniqueSides*(ivar-1) per sub-face, exactly the 2D scheme; the message pattern remains async point-to-point, one round per stage.

2.7 Balance and hanging edges#

2:1 balance is enforced across faces only, the direct analogue of 2D (which balances across edges and tolerates 2-level corner jumps). DG face mortars carry all interface data; hanging edges/corners carry none, so face-2:1 balance is sufficient for well-posed connectivity. EmitMesh guards MaxLevelJump() <= 1 as in 2D.

2.8 Solution transfer#

Operators. Prolongation: U_child = (R_kx ⊗ R_ky ⊗ R_kz) U_parent (exact), (kx,ky,kz) = (axc(c)+1, ayc(c)+1, azc(c)+1). Restriction: U_parent = Σ_c (P_kx ⊗ P_ky ⊗ P_kz) U_child(c) — conservative, and Restrict(Prolong(u)) = u to roundoff via Σ_k P_k R_k = I per direction. Reference-cell conservation: Σ w³ u_parent = (1/8) Σ_c Σ w³ u_child. Each 1D P_k carries the half-interval Jacobian, so the triple product carries 1/8. Implemented in SELF_SolutionTransfer_3D.f90 as three sequential directional contractions; the loop order there is normative, since CLAUDE.md forbids reordering floating-point reductions.

Epoch protocol. The plan (SELF_TransferPlan_3D) is built on the host after the last forest mutation and classifies every new leaf as COPY, PROLONG or RESTRICT, plus a depth and an octant path for any further descent. It is applied around the regrid as a three-step, type-bound protocol on the model, so the backend split that already selects Regrid selects the transfer too:

call model%StageSolutionForTransfer()   ! preserve the field; Regrid may then free it
call model%Regrid(newMesh,newGeom)
call model%ApplyTransferPlan(plan,interp,eFirst,eLast)

StageSolutionForTransfer exists because Regrid reallocates (or resizes) the storage the solution lives in, so the pre-regrid field must be preserved first. The portable implementation stages into DGModel3D_t%transferStage, a host array whose lifetime is exactly stage-to-apply — it is allocated by the first call, consumed and deallocated by the second, and released in Free. Calling ApplyTransferPlan without a preceding stage is a hard error (stop 1), pinned by test/dgmodel3d_guard_transfer_unstaged.f90, because the alternative is reading unallocated storage and failing somewhere less obvious.

ApplyTransferPlan takes an optional uGlobal argument, plus an optional oldFirst: the multi-rank path. On more than one rank the controller hands in host-side old-field data and each rank fills exactly its own new contiguous element range. Two forms exist. By default (Stage-5 v2, §4) uGlobal is the rank's window of the old field — the contiguous run of old elements its new range references — and oldFirst is that window's first global old element index, so the plan's global sourceElem/family indices still resolve correctly. Under SELF_AMR_MIGRATE_GATHER=1 (Stage-5 v1) it is the whole allgathered global old field and oldFirst is absent, which is the same call with oldFirst = 1. Either way migration is a host operation, so that branch stays on the portable host transfer on every backend; serving the device transfer from a migrated window is the named follow-up to #167.

Device path. On a single-rank GPU build both steps are overridden (src/gpu/SELF_DGModel3D.f90): staging is a device-to-device copy into a model-owned buffer, and the plan is applied by TransferSolution_3D_gpu, so an adapting run moves no solution data across the host link at all. Consequence worth knowing: the transferred solution is then left on the device and solution%interior (the host mirror) is stale after Adapt. That matches the rest of the time loop, where the device is authoritative, but it is a behaviour change from the host-only implementation, where the mirror happened to be fresh. See §5 for the kernel, and the Learning page for the measurements.

2.9 Indicator#

3D tensor Legendre modal transform (three passes with the same Pmodal); E_tot, E_clip1, E_clip2 computed over the 3D mode cube with the same top-mode clipping per direction; identical S_e/σ_e definition, amplitude gate, relative energy floor (default 1e-12), hysteresis band, and the two-phase (device/host) split that keeps CPU and GPU flags identical.

2.10 EC / split-form models#

ECDGModel3D raises a runtime error on meshes with nMortars > 0 (entropy-stable mortar operators out of scope), mirroring the 2D guard.

3. What carries over unchanged#

  • Lagrange_t%mortarR/mortarP (+ existing device mirrors) — no new operator math anywhere in the plan.
  • DomainDecomposition_t — contiguous SFC partition of the Morton-ordered leaf list per epoch, elemToRank locality tests, tag conventions.
  • Controller architecture: rank-replicated forest, flag allgather, level cap + halo passes, AdaptFromFlags (coarsen-then-refine, snapshot based), Balance2to1 fixed-point sweeps, no-op detection, BuildTransferPlan walk over stable node ids, EmitMesh regeneration, geometry double-buffering with CopyElements reuse (sourceKind == COPY ⇒ same forest node ⇒ bit-identical coordinates), Regrid via Resize (not Free+Init), Stage-5 v2 point-to-point solution migration (§4), RecommendedTimeStep = dtBase/2**MaxLevel, and the SELF_AMR_GEOM_{FULL,VERIFY,NO_REUSE} and SELF_AMR_MIGRATE_{GATHER,VERIFY} diagnostics.

New 3D-side prerequisites (absent today, required by the controller): Resize on Scalar3D/Vector3D/Tensor3D/mapped variants, SEMHex%Resize/CopyElements/GenerateFromNodeCoords/UploadGeometry + cached scratch + nElem = 0 default, DGModel3D_t%Regrid/ StageSolutionForTransfer/ApplyTransferPlan + transferStage.

Geometry reuse matters more in 3D: CalculateContravariantBasis_SEMHex uses the curl-invariant metric form and is substantially more expensive than its 2D counterpart.

4. MPI strategy#

The 2D document's §7 sketched this; what follows is the record of what was built, for both dimensions, and it supersedes that sketch (see the note at the end).

Ordering and partition. The global leaf order is base-mesh root order with Morton order inside each tree — deterministic and locality preserving. EmitMesh re-runs GenerateDecomposition over the new leaf count every epoch, so each rank owns the contiguous range offsetElem(rank+1)+1 … offsetElem(rank+2). Equal element counts are a valid balance measure because every element carries the same polynomial degree. Repartitioning is therefore implicit in every epoch, and both the old and the new partition are contiguous ranges of the same leaf order — the fact everything below rests on.

Replicated state. The forest is rank-replicated. At Init, InitForestFromDecomposedMesh allgathers the global base-mesh node coordinates, side table and material ids (once, at startup). Each epoch the rank-local indicator flags are allgathered (one int per element) and the indicator's amplitude gate is reduced with MPI_MAX, so every rank applies identical mutations and computes an identical transfer plan. Forest memory is therefore O(global elements) per rank; making the forest distributed is the p4est-style Stage-2 follow-up in #167 and is deliberately not attempted here.

Solution migration (Stage-5 v2, #167). Migration moves only what changes rank:

  • The window. For a rank's new element range [eFirst,eLast], the old elements its plan entries read — sourceElem, or all eight family members for a coarsened family — lie in a contiguous window [wFirst,wLast] of the old element list. PlanWindows computes the min/max hull of those indices for every rank in one pass over the plan. Correctness does not depend on the leaf ordering being monotone: a non-monotone ordering only widens the hull, the worst case being the whole old list, which is exactly what v1 moved. Locality, not correctness, is what the space-filling curve buys.
  • Routing needs no communication. Because the plan and both offsetElem tables are replicated, a rank computes its own window and every peer's window locally. The run it must send to peer r is the intersection of its own old range with peer r's window; the run it must receive from peer r is the intersection of its window with peer r's old range. Both ends of a pair evaluate the same max/min on the same integers, so the schedules match by construction — there is no count exchange, no handshake, and no collective.
  • The exchange. ExchangeOldWindow posts all MPI_Irecvs before the matching MPI_Isends and closes with one MPI_Waitall, the same shape as every other point-to-point exchange in SELF. Nothing is packed: a run of elements at a fixed variable is contiguous in both solution%interior and the window buffer, so each message reads and writes the real storage. That is one message per (peer, variable) — a handful of peers for a balanced repartition, and bounded by the rank count in the worst case — and it removes the pack buffer, the unpack loop and two whole-field copies.
  • Tags. The tag is the variable index. Ambiguity would require two messages between the same ordered rank pair carrying the same variable in one exchange, and each pair exchanges exactly one run per variable. The exchange is drained before Regrid, so migration traffic never coexists with the side, mortar or halo exchanges (which use globalSideId + nUniqueSides*(ivar-1), or tag 0 for the aggregated form) even though all of them share the mesh communicator.
  • Sequencing. The exchange runs before model%Regrid: EmitMesh has already decomposed the new mesh, so both partitions are known, and the sends can read the still-live pre-regrid solution. It also means no traffic is in flight across Regrid, which rebuilds mesh and decomposition state on the same communicator. Overlapping the two is a possible future optimization and would need a tag space of its own.
  • Coarsened families. A family's eight children may be owned by different old ranks. The window hull covers all eight, each child arrives from whichever rank owned it, and RestrictFromChildren runs on the receiving rank in child-octant order — so the projection's operand order, and therefore its result, is independent of the partition.
  • Degenerate cases. A rank owning no new elements gets an empty window (normalized to wFirst=1, wLast=0), posts no receives, skips the apply, and still posts its sends, because peers may need old elements it owns. An empty peer overlap is skipped identically on both sides.
  • Memory. The per-rank migration footprint is the window — local elements plus the overhang past the rank's own old range — instead of a full global field. The window buffer is a controller component, persistent and grow-only, so a settled adapting run performs no allocation in this path.
  • Bit-identity. The two migrations are bit-identical, not merely close: ApplyTransferPlanRange (v1) is ApplyTransferPlanWindow (v2) over the whole old field, so both feed the same operands to the same operators in the same order, and no reduction order changes. SELF_AMR_MIGRATE_VERIFY=1 runs both in one process and compares the window against the allgathered field value by value; SELF_AMR_MIGRATE_GATHER=1 selects v1 outright, keeping the retained fallback exercised.

Exchange-table rebuild. The aggregated halo tables (decomp%halo_*) and each field's persistent halo requests are partition-specific and are torn down on Resize, so Regrid invalidates them and the next exchange rebuilds them. The sub-face global side ids in the regenerated mortarInfo keep the tag convention valid with no new machinery.

Constraint compliance. Every collective and every message lives in the adaptation epoch, between time steps — never in CalculateTendency or an RK stage (CLAUDE.md §5). The communicator is unchanged, no reduction is reordered, and the v1 collective path is retained behind an environment switch rather than deleted (CLAUDE.md §10).

Supersedes: AMR2D-Design.md §7 specified MPI_Alltoallv for redistribution. Replicated routing removes the count exchange that motivates Alltoallv, and the pattern is sparse for a balanced repartition — a rank talks to a couple of neighbours in leaf order rather than to everyone, with the worst case bounded by the rank count and handled — so matched Isend/Irecv pairs are both cheaper and free of an all-ranks synchronization point.

5. GPU strategy#

Same division of labor as 2D: numerics on device, adaptation logic on host. Kernels follow the existing files' pattern: 3D mortar gather/scatter and flux-scatter kernels (SELF_Mortar.cpp), a 3D indicator kernel (SELF_Refinement.cpp, AMR3D_MAXNP bound), and the 3D transfer kernel (SELF_SolutionTransfer.cpp). All are implemented.

The transfer kernel is TransferSolution_3D_gpu. As planned, the 2D kernel's one-block-per-element with Np² threads and one thread per node becomes one-block-per-element with Np² threads where thread (a,b) owns one (a,b) pencil and loops the free index through each of the three directional passes — the same decomposition RefinementIndicator_3D_gpukernel uses. A Np³ thread block is not an option: it is 4096 threads at N=15, past the block limit, and it would push the working buffers into per-thread scratch.

Three Np³ __shared__ buffers do fit, at a bound: AMR3D_MAXNP is 12, giving 3·12³·8 B = 41,472 B, inside both the 48 KB CUDA static shared-memory limit and the 64 KB gfx90a/gfx942 LDS budget. The allocation is static, so it is 41,472 B at every degree, not (N+1)³; a dynamic-shared variant sized to the degree in use measured 1-3% slower on a B300 and was not adopted (see the Learning page's §6.4). The Fortran caller guards interp%N+1 <= 12 so a larger degree fails loudly rather than overrunning.

One deliberate divergence from the host reference, inherited from 2D: the host descent calls ProlongToChildren, which forms all eight children and discards seven, whereas the kernel applies only the operator triple of the child actually on the recorded path — an eighth of the work per descent step. What is dropped is the seven discarded children, not any part of the retained value: each contraction sums the same terms against the same mortar column in the same ascending index order as the host loop, so the reduction order CLAUDE.md protects is preserved. Device and host nonetheless agree only to round-off rather than bitwise, because the device compiler contracts these multiply-accumulates into FMAs. test/solution_transfer_3d_device.f90 pins the kernel's output against the host reference value by value; the AMR regressions assert the invariants (conservation, entropy non-growth), which conservation arguments make exact either way.

The window form (#172). TransferSolution_3D_gpu takes an oldFirst0 argument and nOld means uOld's per-variable element STRIDE, not the global old-leaf count; a plan's global source index is rebased as src - oldFirst0. The whole-field case is exactly oldFirst0 = 0, so the single-rank path is unchanged arithmetic. This is what lets a multi-rank adaptation use the kernel at all: the window of old elements a rank reads is migrated to it point-to-point and, on a GPU build, assembled in device memory — its own run device-to-device, the peers' runs MPI_Irecv'd straight into device memory. That requires a GPU-aware MPI, which is not a new dependency: the per-step aggregated halo exchange has always posted on device allocations with no fallback, and SELF_REQUIRE_GPU_AWARE_MPI is fatal by default. The local run is copied by MigrateWindowLocal_gpu, which synchronizes the device UNCONDITIONALLY before returning, because the caller then posts MPI against solution%interior_gpu; making that contingent on there being something to copy would strand a rank whose local run is empty. HaloPack_3D_gpu synchronizes before its own MPI for the same reason.

A convention this change follows. New index and byte arithmetic belongs in src/, shared with the host path, so that every CI job checks it rather than only the two GPU pipelines (see §6) - and so that a CPU build can fail for a logic error rather than deferring it to hardware. The guard the kernel cannot perform (every source inside the window) likewise lives in the Fortran caller, where it can report; inside the kernel an early return would be neither reportable nor block-uniform, and would strand the __syncthreads() calls.

6. Testing & validation plan#

Serial/CPU (mirroring the 2D suite): - mappedscalarmortarexchange_3d_linear: linear field reproduced exactly across a hand-built 3D mortar mesh (all interior faces conforming or mortar), including a rotated-neighbor variant covering nonzero flips. - mappedvectordgdivergence_3d_mortar: divergence of a linear flux on a mortar mesh matches the conforming result; conservation of the big-face surface integral to roundoff. - lineareuler3d_mortar_soundwave: acoustic pulse crossing a static mortar interface; conservation + entropy checks (a conforming-mesh lineareuler3d_soundwave baseline lands first — LinearEuler3D currently has no dedicated regression test). - octree_balance_3d, mesh3d_uniform_refine, adaptive_mortar_3d, solution_transfer_3d, transfer_plan_3d, refinement_indicator_3d_*: transcriptions of the 2D unit tests (validity invariants, transfer exactness/identity/conservation, indicator smooth/front + guards). - lineareuler3d_amr_soundwave: the 2D soundwave AMR regression in 3D — adapt-to-convergence, conservation per epoch, dt halving, entropy non-increase, forest replication checksum. - Guard tests (WILL_FAIL) for every error stop added, following the 2D guard-test inventory (controller decomp/maxlevel/wrongmesh, unbalanced emit, transfer-plan misuse, indicator argument guards, EC mortar guard).

MPI: mappedscalarmortarexchange_3d_linear_mpi, lineareuler3d_mortar_soundwave_mpi, lineareuler3d_amr_soundwave_mpi (partition-changing adaptation), geometry_3d_reuse_mpi + env-gated verify.

Solution migration (§4, #167) is covered by four additions, chosen so that each one can fail for a different reason:

  • transfer_plan_3d_window (serial): drives PlanWindows and OwnedRun over a table of hand-written (old, new) partition pairs a 2-rank run could never produce - empty new ranges, empty old ranges, a window spanning four peers, a window that is entirely remote - and checks window coverage, window tightness (both ends are actually referenced, so nothing surplus moves), and that the windowed apply is bit-identical to the whole-field apply. It fails if the table stops reaching those configurations.
  • Three WILL_FAIL guards on the windowed apply, one per way a wrong window fails: transfer_plan_3d_guard_window (the window misses a referenced sourceElem), ..._guard_windowfamily (it misses one of the eight children of a coarsened family - the case that arises when a family straddled a rank boundary), and ..._guard_windowbounds (the window is not inside 1..nOld at all). Each must stop 1 rather than read outside the array.
  • amr_migrate_3d_window_mpi (2 and 4 ranks): the real ExchangeOldWindow, on a deliberately ASYMMETRIC adaptation so old elements genuinely change rank. The old field is analytic and encodes each element's global index in its values, so a misrouted element names itself. It asserts that at least one rank received something from a peer - without that, a symmetric refinement can leave every window local, the exchange a no-op, and a passing verify run meaningless.

  • SELF_AMR_MIGRATE_VERIFY=1 / SELF_AMR_MIGRATE_GATHER=1 re-runs of the existing AMR MPI regressions at 2 and 4 ranks: the first asserts v1 and v2 agree value for value in-process on real physics, the second keeps the retained v1 path exercised.

  • amr_migrate_3d_device_apply_mpi (2 and 4 ranks, #172): the same asymmetric epoch, but driven through the MODEL - migrate, regrid, apply - so that on a GPU build it exercises the device window assembly, the receive into device memory and the kernel's oldFirst0, none of which amr_migrate_3d_window_mpi reaches (that one calls ExchangeOldWindow directly, always in host memory). It fails unless some rank received elements from a peer AND some rank's window base is both past 1 and away from that rank's own first old element - the second condition because a window beginning exactly where the rank's old range begins cannot distinguish a correct rebase from a dropped one. Values are checked against ApplyTransferPlanRange over an independently built analytic field: bitwise on a CPU build, to a tolerance on GPU.

  • solution_transfer_3d_device_offset (#172): calls the launcher DIRECTLY with a synthetic window and a nonzero oldFirst0, which separates a kernel indexing error from a migration error. GPU-gated, unavoidably - it is the one gated test in that change, because the kernel is the subject and a CPU build has no second implementation to compare against.
  • SELF_AMR_MIGRATE_VERIFY=1 / SELF_AMR_MIGRATE_GATHER=1 re-runs, plus SELF_AMR_TRANSFER_VERIFY=1 in 2-D. These check DIFFERENT things at different bars and must stay separate: MIGRATE_VERIFY asserts the window arrived intact and is BITWISE, because migration is byte movement; TRANSFER_VERIFY asserts the two applies agree and is tolerance-based, because the device compiler contracts the kernel's multiply-accumulates into FMAs. SELF_AMR_TRANSFER_HOST=1 is deliberately NOT registered as a GitHub Actions re-run: on a CPU build it selects the same code the default takes, so it would cost 33 s across both dimensions for a few lines of branch coverage in a job that already runs within minutes of its 90-minute ceiling. It is exercised where it selects different code - the on-hardware GPU pipelines, and the B300 verification runs, which do an explicit TRANSFER_HOST=1 pass over the whole set.

On patch coverage, and one thing that is easy to get wrong about it. A WILL_FAIL test's stop 1 DOES produce coverage data - gcov's atexit flush runs normally, and the guard lines it reaches show real hit counts. So guard paths are ordinarily coverable and a WILL_FAIL test is the way to cover them; #172 initially assumed otherwise and left twelve reachable lines untested on that mistaken basis.

What genuinely does not get covered is narrower. Assertion branches such as VerifyWindowedApply's mismatch report fire only when the host and device applies disagree beyond tolerance, which needs data corrupted on purpose - leave those, and do not delete an assertion to move a number. Continuation lines of multi-line calls are reported as never executed no matter what, because gcov attributes a statement to its first line. And MPI_IRECV/ISEND/WAITALL inside a multi-rank test are unreliable: the ranks share one .gcda path and clobber each other, so whichever writes last decides which side appears covered.

Codecov's default patch target is the project's own coverage (~93%). Cover what a test can genuinely reach - which is more than it first appears - and say which lines are deliberately left.

Test cost is a real constraint here, but runner speed is the larger term and the two are easy to confuse. Observed gfortran-12 coverage durations on one branch, against a hard timeout-minutes: 90: 85 min, 90 min (cancelled), 82 min, and 43 min - where the 43-minute run carried MORE tests than the 82-minute one. So the job sits near its ceiling and a slow runner, rather than a few added seconds of tests, is what decides whether it gets cancelled. Commit d88910d4 trimmed the 3-D AMR soundwave family for this reason and the constraint clearly recurs, so anything added to the AMR MPI family is still worth costing (Testing/Temporary/CTestCostData.txt) before it is registered. But a trim should not be expected to fix a cancellation on its own, and coverage should not be traded away for seconds a fast runner would have absorbed - #172 did that once and had to put them back.

Two ranks are not enough on their own: with one peer the send and receive runs are forced and no window can straddle more than one peer, which is why the 4-rank entries exist (SELF_MPIEXEC_NUMPROCS_MANY, default 4).

GPU coverage comes from two places, and it is worth being precise about which. Most GitHub Actions jobs (.github/workflows/linux-*) are CPU-only and compile no device kernel, but amd-mi210-gpu-tests.yml is the exception: it runs on the self-hosted galapagos-mi210x4 runner, under rootless podman on the selfish base image, and does an on-hardware AMD MI210 (HIP, gfx90a) coverage build and full ctest. Buildkite (.buildkite/pipeline.yml) supplies the rest: on every non-main branch it uploads an NVIDIA V100 (CUDA, sm_70) on-hardware coverage pipeline and an x86 CPU pipeline. Both GPU pipelines configure SELF_ENABLE_GPU=ON. So a pull request does build and run both device backends, and src/gpu/ coverage is uploaded from both.

Two consequences for the tests above. A GPU-gated test is not dead weight - it runs on both backends per PR - but it is invisible to the much faster and more numerous GitHub Actions jobs, so the default remains NOT to gate: a test that can assert something real on a CPU build should. And keeping new index arithmetic in src/ rather than src/gpu/ still earns its keep, not because the device path is uncompiled, but because arithmetic shared with the host path is checked by every job rather than only by the two GPU ones. Deeper GPU work still wants a dedicated node for measurement (a B300, CUDA sm_103, for #165/#167/#172), since the GPU pipelines are correctness gates and not a benchmarking environment.

Which COMPILERS the estate actually exercises, which is narrower than the support matrix. Every job that runs on a pull request uses gfortran: the GitHub Actions matrices (9/10/11/12, debug/release/coverage), the MI210 GitHub Actions job and the two remaining Buildkite pipelines, which set FC=gfortran and vary only the GPU backend. The linux-amdflang-cmake and linux-nvidia-hpc-cmake workflows exist but are currently disabled_manually in the repository, and there is no ifx job at all. So the compatibility this project asks for in CLAUDE.md - gfortran >= 11, ifx, nvfortran, amdflang - is a standing requirement that CI does not verify for three of the four. Anything resting on a less common language feature (sequence association, non-default integer kinds in unusual positions, type-bound override attribute matching) is checked here only against gfortran, and a green PR should not be read as more than that.

CPU/GPU flag agreement is enforced by construction (shared host FinalizeIndicator); post-adapt solution agreement by tolerance.

7. Work breakdown & phasing#

Stage Content Mirrors
1 3D face mortars: mortarInfo, SimpleMortarMesh3D(+rotated), scalar/vector exchange + flux collect, model integration + EC guard, serial+MPI tests PR #145
2 Octree forest, refinement primitives, uniform refine, EmitMesh, validity tests AMR Stages 2/4a/4b
3 Indicator 3D, transfer plan/solution transfer, Resize/geometry infrastructure, Regrid, controller, unit tests AMR Stages 1/3/4
4 lineareuler3d_amr_soundwave (+ coarsen-wake) integration tests, example #156/#162
5 MPI v1: forest-from-decomposed-mesh, flag allgather, gather-then-slice migration, 2-rank tests Stage 5
5b MPI v2: point-to-point migration - windowed apply, PlanWindows/OwnedRun/ExchangeOldWindow, MIGRATE_GATHER fallback and MIGRATE_VERIFY diagnostic, 2- and 4-rank tests Stage-5 v2 / #167
6 GPU: mortar/indicator/transfer kernels, device staging, geometry upload paths Stage 6

Delivery status: stages 1–3 and 5 landed with #164, together with the lineareuler3d_amr_soundwave serial and 2-rank integration tests of stage 4. Stage 6's mortar and indicator kernels landed with #164; the device-resident solution transfer and its staging landed with #165 (TransferSolution_3D_gpu, StageSolutionForTransfer_DGModel3D, ApplyTransferPlan_DGModel3D), which also delivered stage 4's example, linear_euler3d_amr_spherical_soundwave.

Stage 5b landed with #167, in both dimensions.

Still outstanding: the 3D coarsen-wake regression (the analogue of lineareuler2d_amr_coarsen_wake); the 3D analogues of 2D Stages 6b/6c (amortized high-water-mark storage and geometry reuse) to the extent they are not already inherited from the shared data classes; the device transfer on more than one rank (the migrated window lands in host memory, so a device-side window buffer and an old-element offset in the kernel are what remain); and the distributed forest, which is #167's Stage 2 and is deliberately deferred until measurement justifies it.

Stages are landed as separate reviewable commits on one PR (or stacked PRs at maintainer preference), each leaving the suite green.

8. Compliance with repository constraints (CLAUDE.md)#

Identical posture to the 2D work, which established the precedents: - The mortar interface discretization is the same formulation extension already approved and merged in 2D (#145), extended dimensionally. - Conforming meshes execute byte-identical code paths (nMortars > 0 gates); all existing 3D regression tests must pass bitwise. - No collectives inside time stepping; adaptation runs between steps. - No module renames/moves; additive public API only; Fortran 2008; no external dependencies (octree is home-grown); fprettify formatting. - EC split-form 3D models excluded from nonconforming meshes (guarded).