ECT: atomic test-and-subtract for the face-extension area accounting

[clarkse-he]

Summary

Split out of #7104 (1/3), per review request there.

In the ECT face-extension passes (ComputeOneWayExtensions,
ComputeEightWaysExtensions), several enlarged faces can attempt to borrow area from
the same intruded face concurrently on GPU. The plain read-test-write lets an intruded
face give the same area away more than once (#2257; the fix is equivalent to the one
proposed in #2298): the borrowing outcome then differs from run to run, and the
double-lent area de-stabilizes the ECT update.

What's changed

• The area is taken with an atomic test-and-subtract (amrex::Gpu::Atomic::If with
  amrex::Minus): the subtraction commits only if the remaining area stays positive,
  so each unit of area is lent exactly once. The commit predicate acts on the would-be
  combined value, i.e. remainder > 0 is exactly the old area > S_ext test.
• The borrowing metadata (size, inds_pointer) is re-recorded from the commits that
  actually happened, so a face whose lender was drained concurrently stays flagged for
  the next extension pass (or the BCK correction) instead of carrying stale borrowing
  records into the conformal B push.

Behavior

• CPU: bit-identical. The passes execute serially per fab; the host implementation
  of Atomic::If is the same read-modify-write as before, and a - b > 0 ⟺ a > b
  exactly in IEEE double. Verified on the 2D rotated-square ECT cavity (68 steps):
  single box, 16 boxes on 1 and 2 ranks — max relative difference 0.0 in all six field
  components vs development.
• GPU: deterministic borrowing. Together with the owner-unique borrowing follow-up
  (ECT: owner-unique cross-box borrowing and seam-synchronized conformal B push #7104), run-to-run max relative differences on the 4-box rotated-square cavity drop
  from 4.5e-1 to 2.5e-14 (see the measurement in ECT: owner-unique cross-box borrowing and seam-synchronized conformal B push #7104).

Test coverage

The race has no lanes to run in on the CPU CI runners, so no new test can catch it
there; the existing ECT tests (embedded_boundary_rotated_cube,
embedded_boundary_cube, ...) cover CPU bit-identity. The GPU determinism measurement
is in #7104.

🤖 Generated with Claude Code

https://claude.ai/code/session_01FyFeNNpr5jir3ygSakbZCm

[RemiLehe]

Not need to mention the issue and PR here.

Once merged, the comment would stay in the code but would not be relevant.

[clarkse-he]

Removed, thanks.

[RemiLehe]

Could you add a comment that explains the syntax of amrex::Gpu::Atomic::If?
Right now, it is difficult to parse what this code is doing.

[RemiLehe]

The syntax is here: https://amrex-codes.github.io/amrex/doxygen/namespaceamrex_1_1Gpu_1_1Atomic.html#a54472d6cb90e23b6479dc667f7e02669

and is not necessarily very intuitive (the condition is given at the end)

[clarkse-he]

Added — each argument is now documented inline at the call site, with a link to the AMReX docs. You're right that the trailing condition is the confusing part; I also noted that the flag_ext_face test inside the predicate is what stops the loop after the first successful borrow.

[RemiLehe]

No need to add issue number and PR number

[clarkse-he]

Removed.

[RemiLehe]

Can you check whether this is still needed?

[clarkse-he]

Yes, but it needed narrowing — and it is now smaller than what you were looking at.

EvolveBCartesianECT iterates borrowing_size(i, j, k) entries starting at *borrowing_inds_pointer(i, j, k) with no flag guard, and the counting pass sizes that list from a non-atomic read of S_mod. So a face whose atomic borrow fails would leave the solver reading entries that were never filled in.

Since flag_ext_face is cleared exactly when a borrow succeeds, it already carries that information, so I dropped the n_borrowed counter entirely and keyed the reset off the flag instead.

[RemiLehe]

Could you add a comment explaining why this is needed

[clarkse-he]

Good catch — this came in from #7104 along with the rest of the partial-extension bookkeeping. There, a face's lent area is tracked in a dedicated lent_area MultiFab and reduced across boxes by sync_lent_areas(), which is what makes a partially satisfied borrow representable in the first place.

Without that ledger it was simply wrong here: a face that got only some of its patches left the intruded neighbours' S_mod short with nothing to compensate, which breaks the S_mod + (sum of lent areas) == S identity that the intruded-face branch of EvolveBCartesianECT relies on.

So I've made the extension all-or-nothing instead: if any patch cannot be taken, the successful ones are given back and the face is left flagged, so it gets stabilized by the BCK correction. all_borrowed now only drives that rollback, and says so in a comment. Partial and cross-box borrowing stay with #7104, which will want a rebase on top of whatever lands here.

[RemiLehe]

Could you explain why this is needed

[clarkse-he]

Same as above — this is what triggers the rollback now, and it is commented.

While re-checking it I also found a hole in the count < nborrow cap: the borrowing pass can find more available neighbours than the counting pass reserved slots for, and silently skipping them could unflag a face that never actually reached S_stab. That path now sets all_borrowed = false as well, so the face rolls back and goes to BCK.

[RemiLehe]

Can you check whether this is still needed?

[clarkse-he]

Yes, for the same reason as the one-way case: the solver trusts borrowing_size when it walks the borrowing list, and count could come out below the nborrow slots the counting pass reserved. With the all-or-nothing change it is now either 0 or the full count.

[clarkse-he]

Thanks for the review @RemiLehe — pushed in 0986772.

The two "is this still needed?" threads turned out to be pointing at real cruft. The partial-extension bookkeeping had been imported from #7104, where a face's lent area is tracked in a dedicated lent_area field and reduced across boxes; without that ledger it left intruded neighbours' S_mod short with nothing to compensate. The eight-ways extension is now all-or-nothing — if any patch cannot be taken, the successful ones are given back and the face is left to the BCK correction — so a face is either fully extended or has borrowed nothing, and no partial-state bookkeeping is needed here. Cross-box/partial borrowing stays with #7104.

That also let ComputeOneWayExtensions get smaller: flag_ext_face is cleared exactly when a borrow succeeds, so the n_borrowed counter was redundant and is gone.

One extra fix: the count < nborrow cap could skip a still-available neighbour without recording it, which could unflag a face that never reached S_stab.

Local regression run (2D+3D, EB, OMP):

test	solver	result
test_2d_embedded_boundary_rotated_cube	ECT	run + analysis + checksum pass
test_3d_embedded_boundary_rotated_cube	ECT	run + analysis + checksum pass
test_2d_embedded_boundary_cube	Yee	pass

Both ECT tests pass with OMP_NUM_THREADS unset (112 threads here), so the atomic path was exercised under real concurrency and still reproduces the committed benchmarks. The 3D embedded_boundary_cube checksums differ locally on near-zero components only (Ey/Ez abs ~1e-05 against Ex ~4e6), and those decks do not set algo.maxwell_solver = ect so they never reach this code — I have not touched their benchmarks.

[aarontran]

Hi Eric, I'm working with Remi to study this & subsequent PRs.

Do you have an example situation where borrowing pass gets count > nborrow and thus triggers this check?

If the first lambda function (counting pass) in amrex::Scan::PrefixSum<int> is guaranteed to complete before any calls to the second lambda function (borrowing pass), then it seems count must be strictly <= nborrow because during the first pass, the neighbor faces still have their "maximal" available area in S_mod as no atomic updates have occurred yet. If no such guarantee exists, then I could see the check triggering...

So the question might be translated to: what guarantees, if any, do we have on the execution order of the lambdas in PrefixSum (and where can humans best find information about any such guarantees or lack thereof?) ( ... deferring the question of cross-box or cross-multifab contention til your PR 7104 ... )

[clarkse-he]

Hi Aaron — thanks for the careful read, and for the measurements above.

Short answer: no, there is no such guarantee on CUDA or HIP — and separately, this check isn't really about fin/fout ordering, it's about the fout calls racing with each other, so it would survive even if the ordering guarantee existed.

1. What amrex::Scan::PrefixSum actually guarantees

I couldn't find this documented anywhere, so I went through the implementation at the AMReX commit WarpX 26.08 pins (0579402, Src/Base/AMReX_Scan.H). It's backend-dependent:

backend	structure	all fin before any fout?
CUDA (L810)	one kernel; CUB BlockLoad + BlockScan with cub::TilePrefixCallbackOp decoupled look-back	No. Each block calls fin only over its own tile (L876), and block 0 goes straight to fout (L935) right after its BlockScan; block k only waits on look-back over tiles < k. Nothing ever waits on a later block's fin.
HIP (L642)	same shape, rocprim::detail::lookback_scan_prefix_op	No (fin L732, fout L786).
SYCL, nblocks > 1 (L196, PrefixSum_mp)	three separate kernel launches: all fin (L247), scan the block sums (L299), all fout (L377)	Yes, strictly separated — unless AMREX_SYCL_NO_MULTIPASS_SCAN.
SYCL, single block (L400)	hand-rolled Merrill & Garland	n/a
no GPU (L1117)	plain serial loop, x = fin(i); …; fout(i, y)	interleaved per cell, single-threaded

So on CUDA/HIP the two lambdas overlap in wall-clock time across blocks. As far as I can tell, reading AMReX_Scan.H is the only way to know this; a one-line note on PrefixSum saying fin and fout are not separated by a barrier on the look-back backends would be a worthwhile small AMReX PR if you agree it's contract-level information.

One correction while I'm here, since it bears on how you read my earlier summary: on host, amrex::Gpu::Atomic::If is a plain non-atomic read-test-write (AMReX_GpuAtomic.H L256, the AMREX_IF_ON_HOST branch). Combined with the serial non-GPU scan and the fact that neither ComputeOneWayExtensions nor ComputeEightWaysExtensions has an OpenMP pragma over its MFIter, a CPU build never exercises the atomic path at all within a box. I overstated it above when I said the OMP_NUM_THREADS-unset runs exercised the atomic path under real concurrency — they didn't; those were only a regression check that the restructuring still reproduces the committed benchmarks.

2. Why the check survives even given a strict fin/fout fence

count > nborrow requires the borrowing pass to find more live entries in local_avail than the counting pass did. Working through the inputs:

• S is never written in these kernels.
• flag_info_face is only ever written available → intruded (L969), and local_avail accepts both, so flag updates can't change membership.
• That leaves S_mod. A neighbour is dropped in the while loop when S_mod(neigh) - patch <= 0, so the set can only grow if some neighbour's S_mod is larger at borrowing time than it was at counting time.

Nearly every write to S_mod decreases it, and those land in the !borrowed / denom < S_ext branches instead (count < nborrow). The exception is the rollback this PR introduced, at L984:

amrex::Gpu::Atomic::AddNoRet(
    ::GetNeighPtr(S_mod, i, j, k, vec(0), vec(1), idim),
    borrowing_area[ps + n]);

Face B failing all-or-nothing hands area back to a shared neighbour N, concurrently with face A's borrowing loop. A's counting pass had excluded N as drained; A's borrowing pass recomputes local_avail after the refund and includes it — one slot more than was reserved. A and B are both in fout, so a fence between the passes wouldn't help.

(The other writes that raise S_mod — S_mod(i,j,k) = S(i,j,k), += patch, and = S(i,j,k) + S_ext in the one-way pass — are all on the extending face's own cell, and are harmless here: a face awaiting extension keeps flag_info_face == FaceInfo::extended, which is neither available nor intruded, so it's never a member of anybody's local_avail.)

3. Do I have a concrete trace of it firing?

No — honestly it's a guard, not a fix for something I caught in the act. I added it because the failure mode without it isn't a wrong number: it's borrowing_inds[ps + count] running past this face's reserved slots into the next face's region of the scan-sized DeviceVector. I'd rather pay a compare than rely on the argument that two independent re-derivations of local_avail stay in lockstep across a race.

If you'd like to settle it empirically, your sphere case is a much better instrument than anything I have (5,552 unstable faces per dimension, and you're on GPU). A device counter in that branch —

if (count == nborrow) {
    amrex::Gpu::Atomic::AddNoRet(p_ncapped, 1);
    all_borrowed = false;
    continue;
}

— over a few hundred steps would say whether the path is live in practice. By the analysis in §2 it can only fire on a step where some face also rolled back, so I'd expect it to be rare if it fires at all, and on a CPU build it can't fire by construction. I'd genuinely like to know either way: if it turns out to be unreachable, I'd rather it become an AMREX_ASSERT than stay as silent flow control.

Happy to fold any of this into #7104, which has to rebase on this anyway.

[aarontran]

Thanks, that clarifies & i think no followup needed for now,

1. Agree no guarantees of 1st/2nd function order across threadblocks (based only on quick skim of your reply & AMReX source)
2. Agree on count > nborrow possible due to rollback (after sketching out one possible execution order of two adjoining threadblocks that can trigger this code path on paper; omitted for brevity)
3. Agree on rationale for pre-emptive catch

[aarontran]

I have measured non-determinism arising from ECT face extensions using four versions of the source code:

• WarpX 26.08, one-way pass
• WarpX 26.08, eight-way pass
• This PR, one-way pass
• This PR, eight-way pass
  To force WarpX to use eight-way pass, the param eb_ect_balanced_borrow was backwards-ported from PR# 7105.

TL;DR this PR successfully reduces non-determinism, especially for the eight-way extension code branch (with the edge case of contention across AMReX box boundaries deferred to PR# 7104). Seems OK with me to merge.

The test setup is 3D spherical cavity's TM_{011} eigenmode from Xiao & Liu (2008). Each dimension has 5,552 unstable faces to be stabilized within WarpXFaceExtensions.cpp. WarpX input parameters in this file:

inputs_test_3d_embedded_boundary_sphere_ppw64.txt

More details about the test setup at this link:

https://docs.google.com/presentation/d/1twPuWJrKBFMwqh3q1L-SJCZd19iPKtS3Mz80TlFsbTw/edit?usp=sharing

Main results below. In WarpX 26.08 the eight-way pass leads to noticeably stronger non-determinism at the embedded boundary, compared to the one-way pass. This PR reduces non-deterministic effects in both one- and eight-way passes; note the colorbar magnitudes.