Repository navigation
update cross_proj - #1
Merged
Merged
Conversation
…4773) ## Summary We already use features of `FindCUDAToolkit` for nvtx3 that are CMake 3.25+. Modernize and remove pre CMake 3.20 CUDA logic. Move CUDA architecture handling to the standard CMake interface and drop the deprecated CMP0104 OLD path and the FindCUDA/select_compute_arch mangling (convert_cuda_archs / set_cuda_architectures / set_nvcc_arch_flags), which broke on Hopper-next/Blackwell (sm_100/sm_120) and CUDA 12.9. Closes AMReX-Codes#3948. - Default to `native`: CMake resolves it at **configure** time from the GPU(s) visible on the machine running `cmake`, and the resolved architecture is what gets built, reported and exported. If no GPU is visible and no architecture was given, configuration stops with an error pointing at `-DCMAKE_CUDA_ARCHITECTURES=...`. - `AMReX_CUDA_LTO=ON` now performs real CUDA device LTO via `INTERPROCEDURAL_OPTIMIZATION` (requires RDC; CMake >= 3.25). It covers the AMReX libraries *and* the targets set up with `setup_target_for_cuda_compilation`, so dependent code takes part in the LTO instead of getting nvlink's *partial* LTO. Host code generation is deliberately left untouched. - Minimum CMake is now 3.25 — for building AMReX and for consuming an installed AMReX (the installed config used to demand the exact CMake version AMReX was built with). User transition: - Select architectures with the standard `-DCMAKE_CUDA_ARCHITECTURES=80` (or the `CUDAARCHS` env var). Integers only (`80`, not `8.0`); `native`, `all`, `all-major`, `-real`/`-virtual` and `90a` forms are supported, as is `OFF` for packagers who pass the flags themselves. - `-DAMReX_CUDA_ARCH` and the `AMREX_CUDA_ARCH` env var still work but warn. They map to `CMAKE_CUDA_ARCHITECTURES` / `CUDAARCHS` with unchanged precedence, and their historical spellings (`Auto`, `Common`, `Volta`, `8.0`, `7.0+PTX`, whitespace-separated lists) are translated. - On GPU-less machines (login nodes, CI, container builds), pass an explicit architecture instead of relying on the `native` default. This includes configuring on a login node and building on a GPU node. - Architectures below compute capability 6.0 are dropped with a message (also when they are part of what `all`/`all-major` stands for); configuration fails if nothing is left. - `CUDA_FORWARD_UNKNOWN_FLAGS_HOST` removed (CMake >= 3.17 forwards natively). Also removes now-dead CMake < 3.20 guards. ## Additional background [CMake 3.25 was released in November 2022](https://github.com/Kitware/CMake/releases/tag/v3.25.0) and is widely available and very easy to install as a user, if missing. Also a bit of prep for AMReX-Codes#4572 Close AMReX-Codes#3948 Close AMReX-Codes#4572 Close AMReX-Codes#5524 ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [x] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [x] include documentation in the code and/or rst files, if appropriate --------- Co-authored-by: Sam Mish <samuelpmish@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: Weiqun Zhang <weiqunzhang@lbl.gov>
## Summary `amrex::Initialize` and `amrex::Finalize` drained their hook stacks with `top()()` followed by `pop()`. A hook that registers another one with `ExecOnInitialize`/`ExecOnFinalize` therefore left the *calling* hook on top of the stack: it ran a second time, and the following `pop()` silently discarded the newly registered hook without ever running it. Move the hook off the stack before calling it, so that a hook which registers another one runs exactly once and the newly registered hook runs as well. ## Additional background Found while making the ABLASTR warn manager state per AMReX cycle in BLAST-WarpX/warpx#7186: a warning recorded from a finalize hook resurrects the warn manager singleton and loses the teardown hook that would have released it again. ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate
Avoids the `MPI_Testall` for empty request vectors. This was observed to improve performance in an MLMG case where the bottom solver only used one MPI rank (the whole simulation used four ranks). In this case, the bottom solver did not emit MPI requests and it is not needed to call `MPI_Testall`. Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
## Summary
`clang-tidy`'s `modernize-pass-by-value` fires on
`detail::gpu_tuple_element`'s constructor and asks for the argument to
be taken by value and moved. Following that advice would be a
pessimization here. This adds a `NOLINTNEXTLINE` with the reasoning
written down, because it is not obvious from the code.
## Why by value is wrong here
The forwarding constructor immediately below already takes everything
that can be deduced — lvalues and rvalues of `T` alike. The
const-reference overload is reached only by arguments it *cannot*
deduce, a braced initializer list above all:
```cpp
struct P { int a, b; };
P p{1,2};
E<P> a(p); // forwarding constructor
E<P> b(P{3,4}); // forwarding constructor
E<P> c({5,6}); // this one -- U cannot be deduced from a braced-init-list
```
(verified by instrumenting the two constructors and running the three
cases)
So taking it by value would add a move for exactly the callers this
overload exists to serve, and save a copy for nobody, since rvalues
never reach it. Removing the overload instead is not an option either —
the braced-init case would stop compiling.
## Why it is showing up now
Nothing changed in `AMReX_Tuple.H`. The SIMD CI job runs `clang-tidy`
over ccache *misses*, so which headers get linted depends on what was
recompiled. A change to a widely included header —
AMReX-Codes#5644 touches `AMReX_Math.H` —
pulls `AMReX_Tuple.H` into the linted set and the finding appears. It is
pre-existing on `development` and unrelated to that work, which is why
it is split out here: this is a two-line comment change that can be
reviewed on its own, and AMReX-Codes#5644 can rebase on it afterwards.
## Testing
```
clang-tidy-21 --config-file=.clang-tidy --header-filter='.*AMReX.*'
```
over `Tests/SIMD/main.cpp` and
`Tests/Particles/ParticleReduceSIMD/main.cpp`: `modernize-pass-by-value`
findings across all AMReX headers go from 1 to 0, and this was the only
one. `AMReX_Tuple.H` still compiles and behaves the same
(`amrex::makeTuple` with scalar and aggregate members).
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…X-Codes#5630) ## Summary This PR adds a volume-fraction-based fallback pathway in the `itracker` construction for EB state redistribution when the original normal-based neighbor selection fails. ## Problem - State redistribution merges cutcells with neighbors to achieve target volume. Currently, neighbors are selected based on the EB normal direction, and the selected local indices are stored in `itracker` array. This is a geometrically reasonable approach when the normal vector accurately represents local EB topology. - However, when the EB is poorly resolved or cells have limited connectivity (e.g., cells near the tip of a narrow canyon), the normal-based selection can point to a covered cell (vfrac=0). The current code aborts in this situation with `Abort: Trying to merge with covered cell`. - This stops simulations even when valid merge targets exist in other directions. ## Solution - This PR adds a fallback selection strategy that activates only when the normal-based method fails. - The additional selection method chooses the primary and secondary merge target with the largest volume fraction. - The implementation has no impact on existing simulations, because it preserves the original normal-based logic when it succeeds. ## Validation - Tested on an ERF application with complex urban geometry. This enables redistribution in previously-aborting cases. The proposed changes: - [ ] fix a bug or incorrect behavior in AMReX - [x] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate
…des#5637) A face-centered box holds one more staggered cell than it holds cell-centered cells, and nodal is how that is expressed. areafrac, facecent and the edge centroids across the stagger were cell-typed, so their last staggered cell fell in the ghost region, out of reach of the ParallelCopy in fill*FC, and the covered-region mask added in AMReX-Codes#5636 cleared that plane with nothing able to write it back. The edge centroid along face_dir had the converse defect: coarsen_from_fine shrinks its box back to cell, so its top plane was never written and kept the 1.0 of a fully open edge, which then overwrote the mask inside the body. One rule now -- face_dir nodal, the other directions as in the cell-centered path -- plus the face_dir edge box grown back to the staggered range it is indexed by. The value at a given index does not move; the last staggered cell of each box becomes valid data the copy can deliver. The same geometry built with one EB box and with many now agrees in every field, where the edge centroid along face_dir used to differ at 1633-3844 cells. Addresses issue 1 in AMReX-Codes#5635, on top of AMReX-Codes#5636.
Fixes AMReX-Codes#5638. **What was achieved.** `amrex::Random()` is documented as `[0,1)` but could return exactly `1.0` — on CUDA/HIP because `1 - curand_uniform()` rounds back up for a tiny draw (~3e-8 per call in single precision), and on the host because `std::uniform_real_distribution` is permitted to reach its upper bound (LWG 2524). Both endpoints are now obtained by **relocating a single endpoint of the underlying generator rather than by arithmetic**, which is exact by construction and cheaper. Adds **`amrex::RandomPositive()`**, a uniform real in `(0,1]`, for samplers singular at zero (`-log(u)`, `pow(u,-a)`, Box-Muller) — user code previously wrote `1 - amrex::Random()`, which is *not* a valid guard since it evaluates to exactly `0.0` in single precision. `RandomGamma` now uses it; both are documented in a new "Random Numbers" section of the user's guide. **What was tested.** New `Tests/Base/Random` in CTest: a deterministic maximal-URBG case (the endpoints are far too rare to sample), the endpoint behaviour of the interval conversions, interval assertions through the `RandomEngine` overloads, and `RandomGamma` finiteness — validated by breaking a guard on purpose to confirm it fails. Passes on CPU (9/9, Small) and single-precision CUDA. On an A2000 the original bug reproduces at **109 violations in 3.9e9 draws** (predicted ~118) and goes to **0**; chi-square over 1024 bins is clean for both generators (z = +0.28 and −0.00). Both precision branches compile clean under IEEE and `--use_fast_math -ftz=true`. A third commit pins the **MSVC minimum to 19.44 (VS 2022 17.14)**, up from an untested 19.34 — see the details block. **Does it break user interfaces?** The API is purely additive — one new function, no signature or default changed. **`Random()` does return different values on GPU for a given engine state**, because the interval conversion is no longer `1 - c`, so GPU results and any checksums taken from them will shift. This is not a distribution change: `P(== 0)` is unchanged and chi-square is indistinguishable. CPU streams are unaffected. <details> <summary>Details, corner cases, and rejected ideas</summary> ### The clamp constant Used on the host and SYCL paths, where the native interval is already `[0,1)` and only the vendor defect needs guarding. It is `1 - eps/2`, exactly `std::nextafter(Real(1), Real(0))` in both precisions (verified for `float` and `double`, with and without `-ffast-math`). Being `constexpr`, it works in device code where `std::nextafter` is unavailable — so neither workaround discussed in AMReX-Codes#5638 is needed: not `1 - epsilon` (2 ULPs below one, needlessly discarding a representable value), and not a host-computed value stashed in `RandomEngine`. Incidentally, libstdc++ already clamps a maximal URBG to exactly this value. ### Safety under all math modes Both guarantees hold however AMReX is built, because no path performs arithmetic on the draw: every interval is established by a **comparison**, so there is nothing to reassociate, contract or flush. `almost_one` is exactly representable, so constant folding is stable, and the smallest value `RandomPositive()` can return is a **normal** number in both precisions. Verified: all GPU results are byte-identical under `-O3` and `--use_fast_math -ftz=true`. ### Does any of this perturb uniformity? No. In every case only a single endpoint value is affected and every other value passes through bit-for-bit unchanged, so the distribution restricted to the documented interval is untouched. - **The GPU remap** relocates `2^-25` of mass from `1.0`, which was never in contract, to `0.0` — which is exactly where the old `1 - c` form put it too, so `P(Random() == 0)` is unchanged at ~`2^-25` (measured: 3.75e-08 over 4e8 draws). - **The host/SYCL clamp** moves `2^-25` of mass by one ULP, bounding any moment shift by `2^-49` (~2e-15). - **`RandomPositive`** performs no relocation at all on CUDA/HIP, and on host/SYCL moves `P(u == 0)` onto `1.0` (measured `P(RandomPositive() == 1)` = 2.5e-08). Measured chi-square over 1024 bins at 4e8 draws in single precision on GPU: `Random()` z = +0.28, `RandomPositive()` z = −0.00. The engine stream is untouched in all cases — no extra draws are consumed and engine state is not altered. One corner case on the record: under the GPU remap the atom at exactly `0.0` sits next to the generator's fine `~2^-32` grid, so zero is locally overweighted relative to its immediate neighbours, where under `1 - c` the grid near zero was a uniform `2^-24`. `P(== 0)` is identical either way and this is invisible at any binning coarser than ULP scale, but it is the one respect in which the remap's distribution is less locally regular. ### Why the interval conversion is an endpoint remap, not `1 - c` `(c == 1) ? 0 : c` converts `(0,1]` to `[0,1)` exactly, because the two intervals differ in precisely one endpoint. Subtracting from one also converts the interval, but not exactly, and costs more: - for a draw below half an ULP of one, `1 - draw` rounds back up to `1.0` — the bug; - `1 - draw` is exact by Sterbenz above `0.5`, which collapses every returned value below `0.5` onto a `2^-24` grid. So the old form could not return anything between `0` and `2^-24`, discarding roughly nine bits for small values. The remap keeps the generator's own `~2^-33` floor (measured: smallest positive `5.8e-10` vs `6.0e-08`). - it is one comparison instead of a subtract plus a guard. Measured over 3.93e9 draws, the two are statistically indistinguishable (chi-square 1089.6 vs 1089.2) and have identical `P(== 0)`, so the resolution and cost are gained for nothing. The cost is that every GPU draw changes value, which is called out above. An endpoint-only variant that keeps the flip — `v = 1 - c; return (v == 1) ? 0 : v;` — would have preserved GPU streams (differing in only 111 draws in 3.9e9), but it **doubles** `P(Random() == 0)` (5.59e-08 vs 2.77e-08), and zero is the endpoint that makes `std::log(u)` infinite. That is also why the host and SYCL paths, which need only a defect guard rather than an interval conversion, *clamp* instead of relocating: clamping does not raise the probability of returning zero. Full analysis with measurements in [this comment](AMReX-Codes#5643 (comment)). ### Why not just `amrex::RandomNormal` for the Box-Muller case? For a plain normal deviate, yes — the guide's distribution table lists it directly below, and the doxygen carries a `\see`. Not a coincidence that it is unaffected: `curand_uniform` excludes zero *precisely because* curand's own normal generator needs `log(u)`; the `1 - c` flip discarded exactly the property the vendor put there on purpose. Hand-written Box-Muller still earns its place when the **radial variable itself** is needed — a radially truncated Gaussian (rejection on `RandomNormal` would be an unbounded loop with divergent warps, variable draw counts, and a *square* rather than *circular* cut), or an isotropic direction from a normalized vector of deviates. ### Pinning the MSVC minimum (third commit) The declared floor was MSVC 19.34 (VS 2022 17.4, Nov 2022), but nothing exercises it. Since AMReX requires C++20, the Windows toolchains it is actually built with are much newer: - **CI**: the `windows-2022` and `windows-latest` images currently provide **MSVC 19.44.35228** and **19.51.36252**. So 19.44 is the oldest MSVC any test has ever run against. - **conda-forge**: the `amrex` feedstock, and the downstream `warpx` and `impactx` feedstocks, all build Windows with the `vs2022` compiler package (`c_stdlib: vs`, no lower pin), so distributed binaries are not built with anything older either. Raised to 19.44 in `CMakeLists.txt` and `BuildingAMReX_Chapter.rst` so the claim matches what is verified. Note this is a **policy change**, not a consequence of the RNG work — happy to split it out or pick a different number. The natural alternative is **19.42 (VS 17.12)**, the first release containing the `generate_canonical` fix — less aggressive, though still untested. Note the tested floor is a moving target: `windows-2022` is on its final servicing baseline and `windows-latest` has already moved to VS 2026, so this number will need revisiting when that image retires. This also corrected the framing of the host clamp. An earlier revision justified it by "MSVC 19.34 is defective and is supported"; with the floor at 19.44 every supported standard library clamps already. The clamp stays, but for the accurate reason: **the C++20 wording itself does not guarantee the bound** (LWG 2524; MSVC only conformed once it implemented P0952R2 in VS 2022 17.12 = MSVC 19.42, so the whole of the old 19.34–19.41 window is affected and none of the new one is). The interval should be a property of AMReX, not of whichever standard library is in use. ### Known limitations - The deterministic test pins the *guard* rather than `Random()` itself, because there is no URBG injection point; it catches a regression in the guard or its constant, but the call-site coverage remains statistical. A seam for injecting an engine would fix that — out of scope here, happy to add. - `Random()`'s interval is a guarantee about the returned value, not about expressions computed from it: `problo + Random()*dx` can still round onto the upper cell face (verified in both precisions). Noted in the guide. - `FillRandom` remains backend-dependent. Its CPU path is now clamped (free — already looping), but the GPU paths forward to vendor *bulk* generators where a guarantee needs a full extra pass over the array. Documented accurately rather than papered over. - Naming is bikesheddable: `RandomPositive` fits the `Random*` prefix, `RandomNonZero` reads more literally. - A `(0,1)` variant excluding *both* endpoints would additionally serve logistic/Cauchy inversion, at the cost of a clamp on every backend. Left out. ### Downstream ImpactX BLAST-ImpactX/impactx#1627 can drop its `amrex::max(u1, numeric_limits::min())` clamps in favour of `RandomPositive`; WarpX has four sites that want it (`SampleGaussianFluxDistribution.H:49`, `InjectorMomentum.H:320,322`, `DefaultInitialization.H:71`). </details> ## To Do - [x] draft 🤖 Generated with [Claude Code](https://claude.com/claude-code) - [x] self-review - [x] LLM-based review - [x] final self-review --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
`OpenBCSolver` rebuilds its spectral Green's function whenever `setGreensFunction()` is called, and there is no way to hand it one that was built earlier. Add `greensFunctionFFT()` accessors so a caller can keep its own set and move them into and out of the solver with `std::swap` rather than recomputing. `solve()` only reads the spectral Green's function, and `prepare_openbc()` concerns the R2C plans rather than the Green's function values, so one moved in this way needs no further preparation. The one caveat, documented at the accessor: in 2D, `m_G_fft` is an alias of the Green's function R2C's spectral data, which a later `setGreensFunction()` writes into. Callers may write through the reference there, but must not swap or otherwise re-seat it. Additive only, so nothing changes for existing users. Exercised downstream in BLAST-WarpX/warpx#7178 & BLAST-ImpactX/impactx#1621, where reusing Green's functions this way reproduces rebuilding them bit for bit, and where it takes an accelerating beam simulation from 100 Green's function builds to 10. - [x] 🤖 Generated with [Claude Code](https://claude.com/claude-code) - [x] tested with ABLASTR/WarpX and ImpactX - [x] manually self-reviewed --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…5649) These were stubs that aborted, so the hypre bottom solver could not be used with MLEBNodeFDLaplacian. The assembled row reproduces Fapply with homogeneous boundary data, which is what MLMG's bottom solve needs since it works on the residual correction: * Cut cells use the same Shortley-Weller distances and row scaling as Fapply. A covered neighbor drops out of the row and only contributes to the diagonal, matching the zero EB value that Fsmooth also assumes. * sigma given as a MultiFab is averaged onto the faces, volume weighted where there are cut cells. * Nodes left out of the linear system (Dirichlet, covered, or owned by another box) are recognized by an invalid gid and their columns are dropped. * A node outside a Neumann boundary is folded onto its mirror image, which is what MLNodeLinOp::applyBC puts into those ghost cells. Unlike MLNodeLaplacian this is a finite-difference operator, so fillRHS needs no half-volume factor at Neumann nodes. RZ still aborts. Tests/LinearSolvers/NodeEBFDLap solves the same problem with the native and the hypre bottom solver and compares. With max_coarsening_level = 0 the bottom solve does all of the work, so the two solutions can only agree if the assembled matrix is right; they agree to 1e-16 relative in 2D and 3D, with and without cut cells, with constant and variable sigma, and for periodic, Dirichlet and Neumann domain boundaries. --------- Signed-off-by: Roelof Groenewald <rgroenewald@realtafusion.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…5663) ## Summary When used in pyAMReX bindings, `amrex::ParmParse` leaks state between sequential simulations in the same process. This fixes it. ## Additional background Seen in ImpactX with a simulation that turns a non-default option on, like CSR modeling, followed by a simulation that relies on the default (off).
## Summary Rewrites `mlebabeclap_adotx` (2D and 3D) from a `Box`-based kernel that loops internally (`amrex::Loop(box, ncomp, ...)`) into a per-cell `(i,j,k,n)` kernel, and switches its EB geometry arguments from ten separate `Array4`s (`flag`, `vfrc`, `apx/apy/apz`, `fcx/fcy/fcz`, `ba`, `bc`) to the existing `EBData` view, matching what `mlebabeclap_gsrb` already does. `MLEBABecLap::Fapply` now calls it through `AMREX_HOST_DEVICE_PARALLEL_FOR_4D` instead of `AMREX_LAUNCH_HOST_DEVICE_LAMBDA`, and gets the geometry via `factory->getEBData(mfi)`. The arithmetic inside the kernel is untouched: the loop wrapper is removed and the body is only re-indented, and the EB arrays are re-bound to identically named local references (`auto const& apx = ebdata.get<EBData_t::apx>();` etc.), so the expressions are character for character the same. Results are bitwise unchanged. ## Additional background This is the first of several small PRs that split up the stale WIP AMReX-Codes#4922 ("GPU specific kernels for MLEBABecLap"), as requested there. It supersedes the `mlebabeclap_adotx` part of AMReX-Codes#4922. The templating on `T` that AMReX-Codes#4922 also introduced has been dropped, per the review comment on that PR. The `(i,j,k,n)` form is the prerequisite for a later PR that adds a fused (`MultiFab`-wide `ParallelFor`) GPU path for `Fapply`; that follow-up is not included here. Testing (all commands run from the repo root unless noted): * `cmake -S . -B build-split -DAMReX_SPACEDIM=3 -DAMReX_EB=ON -DAMReX_LINEAR_SOLVERS_EM=OFF -DAMReX_ENABLE_TESTS=ON -DAMReX_TEST_TYPE=Small -DAMReX_MPI=ON -DCMAKE_BUILD_TYPE=Release` then `cmake --build build-split -j8` and `ctest --test-dir build-split --output-on-failure` → builds clean, 9/9 tests pass. * `Tests/LinearSolvers/CellEB`, `make -j8 COMP=llvm USE_MPI=FALSE DIM=3` and `DIM=2`, run over seven configurations (`sphere`, `sphere` + `eb_is_dirichlet=1`, `rotated_box`, `two_spheres`, `flower`, two-level `sphere`, periodic `sphere`; `n_cell=64`, `verbose=2`). The MLMG/BiCGStab residual histories and the initial/final max, 1- and 2-norm residuals are **bitwise identical** to `development` in both 2D and 3D. * CUDA build on an NVIDIA RTX A5000 (CUDA 13.2, gcc 11.4): `Tests/LinearSolvers/CellEB`, `make -j8 COMP=gnu USE_MPI=FALSE USE_CUDA=TRUE CUDA_ARCH=86 DIM=3`, same seven configurations. MLMG iteration counts are identical to `development` (9, 12, 11, 3, 11, 28, 11); the residual values differ only at the level of the run-to-run nondeterminism of the *unmodified* binary (GPU reductions are not bit-reproducible), and all solves converge to `resid/resid0 ~ 1e-13`. * The same CellEB matrix was re-run on Linux/gcc 11.4 (`make -j8 COMP=gnu USE_MPI=FALSE`, `DIM=3` and `DIM=2`) against a `development` build in a sibling worktree: residual histories again **bitwise identical** in both dimensions. ### Performance Neither timing changed measurably; this PR is a refactor, not an optimisation. CPU, `Tests/LinearSolvers/CellEB` `main3d.gnu.TEST.ex` (`make -j8 COMP=gnu USE_MPI=FALSE DIM=3`, gcc 11.4, single rank), `inputs n_cell=128 eb2.geom_type=sphere eb_is_dirichlet=1 verbose=1`. All binaries were built first and then timed **interleaved** in the same session (3 reps each) on a shared machine, so the absolute numbers are inflated but the comparison is fair. Best of 3, MLMG `Timers: Solve` [s]: | grids | `development` | this PR | | --- | --- | --- | | `max_grid_size=32` | 4.284 | 4.328 | | `max_grid_size=64` | 4.534 | 4.554 | GPU (NVIDIA RTX A5000, CUDA 13.2, `make -j8 COMP=gnu USE_MPI=FALSE USE_CUDA=TRUE CUDA_ARCH=86 DIM=3`), same test at `n_cell=256`, again built up front and timed interleaved, best of 3: | grids | `development` | this PR | | --- | --- | --- | | `max_grid_size=32` (512 boxes) | 1.085 | 1.083 | | `max_grid_size=64` (64 boxes) | 0.834 | 0.833 | MLMG converged in the same number of iterations in every one of these runs. ## Checklist The proposed changes: - [ ] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate P.S Generated using Claude Code
## Summary Rewrites `mlebabeclap_gsrb` (2D and 3D) from a `Box`-based kernel that loops internally into a per-cell `(i,j,k,n)` kernel, and calls it from `MLEBABecLap::Fsmooth` through `AMREX_HOST_DEVICE_PARALLEL_FOR_4D` instead of `AMREX_LAUNCH_HOST_DEVICE_LAMBDA`. In 3D this also removes the hand-written quadruple `for` loop and the `// amrex::Loop here causes gcc 8 to crash.` workaround that went with it, since there is no longer a loop inside the kernel at all. The arithmetic is untouched: only the loop wrapper is removed and the body re-indented. The kernel already took its EB geometry as an `EBData`, so no argument rework was needed beyond dropping `Box const& box` / `int ncomp` in favour of `(i,j,k,n)`. Red/black Gauss-Seidel only writes cells of one parity and only reads neighbours of the other, so the switch from a sequential `amrex::Loop` to a concurrent `ParallelFor` introduces no ordering dependence. Results are bitwise unchanged. ## Additional background This is part of the ongoing split of the stale WIP AMReX-Codes#4922 ("GPU specific kernels for MLEBABecLap"), as requested there; it supersedes the `mlebabeclap_gsrb` part of AMReX-Codes#4922. The templating on `T` that AMReX-Codes#4922 also introduced has been dropped, per the review comment on that PR, and AMReX-Codes#4922's change of `dhx` from `m_b_scalar/(h[0]*h[0])` to `m_b_scalar*dxinv[0]*dxinv[0]` has deliberately been left out so that the answers stay bit-for-bit the same. This PR is independent of the `mlebabeclap_adotx` / `mlebabeclap_adotx_centroid` PRs (different kernels, different function in `AMReX_MLEBABecLap_F.cpp`) and can be merged in any order relative to them. The `(i,j,k,n)` form is the prerequisite for a later PR that adds a fused (`MultiFab`-wide `ParallelFor`) GPU path for `Fsmooth`; that follow-up is not included here. Testing (all commands run from the repo root unless noted): * `cmake -S . -B build-split -DAMReX_SPACEDIM=3 -DAMReX_EB=ON -DAMReX_LINEAR_SOLVERS_EM=OFF -DAMReX_ENABLE_TESTS=ON -DAMReX_TEST_TYPE=Small -DAMReX_MPI=ON -DCMAKE_BUILD_TYPE=Release` then `cmake --build build-split -j8` and `ctest --test-dir build-split --output-on-failure` → builds clean, 9/9 tests pass. A 2D library build (`-DAMReX_SPACEDIM=2 -DAMReX_EB=ON`) also builds clean. * `Tests/LinearSolvers/CellEB`, `make -j8 COMP=llvm USE_MPI=FALSE DIM=3` and `DIM=2`, run over seven configurations (`sphere`, `sphere` + `eb_is_dirichlet=1`, `rotated_box`, `two_spheres`, `flower`, two-level `sphere`, periodic `sphere`; `n_cell=64`, `verbose=2`). This test drives the V-cycle, so `Fsmooth` is called on every level. The MLMG/BiCGStab residual histories and the initial/final max, 1- and 2-norm residuals are **bitwise identical** to `development` in both 2D and 3D. * CUDA build on an NVIDIA RTX A5000 (CUDA 13.2, gcc 11.4): `Tests/LinearSolvers/CellEB`, `make -j8 COMP=gnu USE_MPI=FALSE USE_CUDA=TRUE CUDA_ARCH=86 DIM=3`, same seven configurations. MLMG iteration counts are identical to `development` (9, 12, 11, 3, 11, 28, 11); the residual values differ only at the level of the run-to-run nondeterminism of the *unmodified* binary (GPU reductions are not bit-reproducible), and all solves converge to `resid/resid0 ~ 1e-13`. * The same CellEB matrix was re-run on Linux/gcc 11.4 (`make -j8 COMP=gnu USE_MPI=FALSE`, `DIM=3` and `DIM=2`) against a `development` build in a sibling worktree: residual histories again **bitwise identical** in both dimensions. ### Performance Neither timing changed measurably; this PR is a refactor, not an optimisation. CPU, `Tests/LinearSolvers/CellEB` `main3d.gnu.TEST.ex` (`make -j8 COMP=gnu USE_MPI=FALSE DIM=3`, gcc 11.4, single rank), `inputs n_cell=128 eb2.geom_type=sphere eb_is_dirichlet=1 verbose=1`. All binaries were built first and then timed **interleaved** in the same session (3 reps each) on a shared machine, so the absolute numbers are inflated but the comparison is fair. Best of 3, MLMG `Timers: Solve` [s]: | grids | `development` | this PR | | --- | --- | --- | | `max_grid_size=32` | 4.284 | 4.280 | | `max_grid_size=64` | 4.534 | 4.597 | GPU (NVIDIA RTX A5000, CUDA 13.2, `make -j8 COMP=gnu USE_MPI=FALSE USE_CUDA=TRUE CUDA_ARCH=86 DIM=3`), same test at `n_cell=256`, again built up front and timed interleaved, best of 3: | grids | `development` | this PR | | --- | --- | --- | | `max_grid_size=32` (512 boxes) | 1.085 | 1.062 | | `max_grid_size=64` (64 boxes) | 0.834 | 0.806 | MLMG converged in the same number of iterations in every one of these runs. Also worth noting for this one: the CPU path changes from a sequential `amrex::Loop` inside the kernel to `AMREX_HOST_DEVICE_PARALLEL_FOR_4D`, and `vlo`/`vhi` are now recomputed per cell rather than per box. Neither shows up in the timings above. ## Checklist The proposed changes: - [ ] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate P.S Generated using Claude Code
## Summary STL EB geometry depends on whether BVH acceleration is enabled. Although `eb2.stl_use_bvh` is intended purely as an optimization, toggling it changes the resulting geometry. For a small STL on a $$48^3$$ grid, the volume-fraction field differs by roughly 30%: the computed fluid volume is `0.0191364` with BVH enabled versus `0.0147104` without it, with 1110 versus 906 cut cells, respectively. Since BVH is enabled by default, users take the BVH path, which in this case disagrees with the non-BVH reference. The fixes address two sources of order dependence: 1. **The reference point depended on the BVH build.** `build_bvh` sorts the triangle array in place, while `prepare()` subsequently uses `a_tri_pts[0]` to seed `m_ptref`. Without BVH, this is the first triangle in the STL file; with BVH, it is whichever triangle sorts first. This changes the reference point and therefore the rays used for crossing tests, which can flip crossing parity at nodes lying exactly on the surface. `build_bvh` now sorts a copy of the triangle array instead. Simply hoisting the `tri0` capture is insufficient because the sidedness reduction skips index 0 to avoid a degenerate self-intersection, relying on the invariant `tri0 == tri_pts[0]`. 2. **`getIntercept` selected the first intersecting triangle rather than the nearest one.** This makes the result dependent on triangle ordering whenever an edge intersects multiple triangles; the test model contains 10 such x-edges. All six interception sites now select the nearest crossing, eliminating this source of order dependence. ## Additional background ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [x] changes answers in the test suite to more than roundoff level - [x] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate STL EB users will see their geometry change, since the BVH is on by default. The new result is the order-independent one matching the unaccelerated path. It is not identical to either previous variant, because the nearest-crossing rule also changes the non-BVH path on multi-crossing edges. P.S Template filled using an AI model
## Summary Both Clang and GCC warn when we use CXX23 features already. Supress them and/or use intrinsics for now to silence them. ## Additional background Seen downstream in WarpX with Clang 22 (counter) and a recent GCC (for assume). ## Checklist The proposed changes: - [ ] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…AMReX-Codes#5662) ResizeRandomSeed built the oneMKL engine_descriptor with an offset of 1. oneMKL initializes engine id as Engine{seed, id*offset}, so all N per-work-item engines were the same philox4x32x10 stream shifted by one element per engine: the k-th draw of work-item t equaled the (k-1)-th draw of work-item t+1, across kernel launches. Any kernel that draws more than one number per work-item therefore produced random fields that were correlated between neighboring work-items in a fixed space-time pattern. Tests/StochasticHeatEquation, which draws two numbers per face per step, measured the equilibrium temperature variance about 24% low on the Intel PVC CI job (64.2 vs 84.97), while CPU and CUDA runs with the same inputs land within a few percent for every seed. A CPU emulation of the offset-1 engine layout reproduces the deficit, and the same emulation with non-overlapping windows does not. Use an offset of 2^36 elements per engine, in the spirit of the distinct curand/hiprand subsequences used on CUDA and HIP, and check that N * offset cannot overflow. Philox skip-ahead is counter arithmetic, so the larger offset does not change the cost of initialization.
…s on nodal point (AMReX-Codes#5666) In that case, both "h" and "scale" are 0.0. The current code divides by h first, then multiplies by scale, resulting in a floating point exception, even though the limit is finite. This also adds new tests for NodalPoissonEB, using both implicit and STL geometries. The STL one fails without this fix.
The STL nodal EB Poisson test does not work in single precision. Add a NO_CTEST option to setup_test that builds the executable and stages its input files without registering a ctest, and use it for the STL test when AMReX_PRECISION is SINGLE.
…#5657) MLEBNodeFDLaplacian::Fapply picks its kernel per box: cut boxes use mlebndfdlap_adotx_eb, regular boxes mlebndfdlap_adotx. Same operator, different floating-point expression, so a node shared by a cut box and a regular box ended up with two copies differing by one ulp. Nothing reconciled them, and the bottom BiCGStab amplified the difference by more than an order of magnitude per iteration until it saturated at the bottom tolerance and the solve stopped converging. An EB nodal solve that converged with a single box failed once the domain was split into three or more. applyBC now calls FillBoundaryAndSync instead of FillBoundary, so every operator application starts from data whose duplicated nodes agree. This fixes all the nodal operators at once and is robust to any other source of asymmetry (FMA contraction, vectorization, GPU). That makes the trailing nodalSync in the smoothers redundant, so drop it. The one in MLNodeABecLaplacian's GPU sweep loop stays, since no applyBC separates those sweeps. Iteration counts are unchanged, but nodal results are no longer bitwise identical across box decompositions and rank counts. Also document that the mask is a weight rather than a boolean.
Closes AMReX-Codes#5767 The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate
get/set/enable/disableFPExcept were implemented for Linux/glibc only, and silently did nothing elsewhere. But we do enable trapping on macOS, on both x86_64 and arm64. So a caller that disables trapping to survive a failing computation was still killed by a signal there. Read and write the trap bits of the FP control register directly.
…s#5897) AMReX-Codes#5657 dropped these, on the grounds that every consumer of sol goes through MLNodeLinOp::applyBC and its FillBoundaryAndSync. Prolongation does not: interpolation() reads the smoothed coarse correction directly. The smoothers are Gauss-Seidel over validbox, so each box leaves a different value on a shared node, and prolongation consumed unsynchronized data. MLMG_NodalPoisson drifted 1.6e-6 and lost its box-decomposition independence. The applyBC change from AMReX-Codes#5657 is kept; it is the actual fix for the EB nodal solve.
The level subdirectories were always created as `Level_<i>`, so writing a plotfile with a non-default levelPrefix aborted in VisMF::Write because the directory it writes into was never created. Fixes AMReX-Codes#5788.
Move the max triangle count assert from the binary reader into prepare(), so ASCII STL files are bounded too and cannot build a BVH deeper than the fixed traversal stack. Also make setBVHOptimization abort when enabling BVH after an STL was read without it, which would leave fill/getBoxType traversing a null root. Fixes AMReX-Codes#5811.
Guard amr.regrid_file against more levels than amr.max_level, reset last_smallplotfile alongside last_plotfile on restart and stop_run, use std::floor so plot_log_per fires at the t=1 decade crossing, size the derive source MultiFab from the full nGrowVect, and propagate m_empty and m_initialized in the AuxBoundaryData copy constructor. Fixes AMReX-Codes#5808.
The GPU fallback passes nnz to a segmented radix sort whose item count is 32-bit, so assert on it as this function already does for nrows. The other item in the issue, the cfine lifetime in MLNodeTensorLaplacian, is handled by AMReX-Codes#4895 with The_Async_Arena. Fixes AMReX-Codes#5806.
…MReX-Codes#5888) Message::wait() now marks the message finished so count()/pid()/tag() work afterwards, and the pre-call comm-profile entries use the BeforeCall/NoTag sentinels instead of the not-yet-valid status. SYCL Handler::warpDim() returns the sub-group size, not the number of sub-groups. Gpu::fillAsync routes types whose size is not a multiple of 4 to the ParallelFor fallback instead of failing a static_assert. Fixes AMReX-Codes#5800.
advance() in MRI mode never set the fast fixed step, so the fast integrator ignored integration.fast_time_step. Step sizes and tolerances set after setup were also ignored: switching to adaptive stayed fixed, and set_tolerances / set_fast_tolerances had no effect. Pass step sizes (0 for adaptive) and tolerances to SUNDIALS whenever they change, in both advance() and evolve(). A fixed slow MRI step is kept, as before, because the fixed-step coupling table may lack an embedding. Abort with a clear message for adaptive slow MRI steps before SUNDIALS 7.2, which does not support them. Default fast_time_step to 0 and document the unset case. Closes AMReX-Codes#5795
…5987) The 3D grad_eb_of_phi_on_centroids{,_extdir} kernels used cells that are not connected to (i,j,k), e.g. across a wall thinner than a cell, unlike the 2D twins. Skip them. The 2D grad_eb_of_phi_on_centroids_extdir rejected only two of the four ghost corners. Reject all four as in the other extdir kernels, and skip cells outside phi before building rows. Closes AMReX-Codes#5777
…ReX-Codes#5989) eb_interp_centroid2facecent_{x,y} pick the auxiliary row/column on the side the face centroid points to, but the final tangential step always assumed it was at +1. On the -1 side this extrapolated. Use the signed offset, as the 3D version does. Verification (2D, CPU, before vs after, calling EB_interp_CellCentroid_to_FaceCentroid directly): a linear field 1 + 2x + 3y sampled at cell centroids around a cylinder. The interpolation should reproduce a linear field exactly. I tried three geometries: centred on 32², off-centre on 64² and 48². ``` ┌─────────────────────┬───────────────────┬───────────┐ │ Cut faces │ Before │ After │ ├─────────────────────┼───────────────────┼───────────┤ │ x-faces with fc < 0 │ error 0.010–0.015 │ ≤ 1.8e-15 │ ├─────────────────────┼───────────────────┼───────────┤ │ y-faces with fc < 0 │ error 0.007–0.010 │ ≤ 1.8e-15 │ ├─────────────────────┼───────────────────┼───────────┤ │ faces with fc ≥ 0 │ roundoff │ unchanged │ └─────────────────────┴───────────────────┴───────────┘ ``` Closes AMReX-Codes#5699
AMReX does not use errno from math functions. Without errno, math functions like sqrt are faster. Results are unchanged, but with FPE trapping clang may now trap in guarded code like x > 0 ? sqrt(x) : 0, as GCC already can. CMake: new option AMReX_MATH_ERRNO (default OFF) and exported target Flags_NO_MATH_ERRNO for GCC/Clang-family compilers and the CUDA host compiler. GNU make: new variable USE_MATH_ERRNO (default FALSE). Closes AMReX-Codes#5665 --------- Co-authored-by: Axel Huebl <axel.huebl@plasma.ninja>
…des#5991) buildFCData coarsens a 2x refined level built with eb2.max_grid_size. With an odd max_grid_size its grids may not coarsen by 2, giving an overlapping BoxArray and garbage face-centered data. Round the size down to even and assert the grids are coarsenable (AMReX-Codes#5635 item 3). A coarsening error was ignored and the data marked built. Leave the level's face-centered data unbuilt instead, so a face-centered factory on it fails (item 4). BuildFC covers every IndexSpace level, and coarse levels can fail where the finest level is fine, so this does not abort.
…-Codes#5994) With -fopenmp=libgomp, clang ignores OpenMP pragmas, so GNU Make HIP+OMP builds ran on one thread. Use plain -fopenmp. With real OpenMP, HIP RDC builds register some kernels launched in OpenMP regions under names missing from the device code, which aborts at run time. Default USE_GPU_RDC / AMReX_GPU_RDC to off when HIP and OpenMP are both on; an explicit setting still wins. Closes AMReX-Codes#5574.
set_amrex_profilers checked TP_PROFILE, which is never set, and the config header template had no entries for AMREX_CRAYPAT, AMREX_FORGE or AMREX_VTUNE. Check AMReX_TP_PROFILE and add the three macros to the template. Also install AMReX_ThirdPartyProfiling.H, as GNU Make does. Allow VTUNE with IntelLLVM (icpx) as well as classic Intel, checking the C++ compiler. Closes AMReX-Codes#5364.
- Add AlgMG as an always-available bottom solver (BottomSolver::algmg). - Add MLMG::setMultigridType (geometric/algebraic/hybrid); hybrid restarts from the best iterate on stall or divergence. - New MLAlgMG assembles the cell and nodal operators, including EB and overset, into a SpMatrix for AlgMG. - Preconditioner mode applies a single AlgMG V-cycle. - PCG accepts negative definite systems. - AlgMG's coarsest level now uses a direct (gathered, LU) solve by default. - Hypre: BoomerAMG setup runs once per assembled matrix; new GPU defaults (PMIS, extended+i interpolation, l1-Jacobi relaxation). - RZ/spherical metric factors are now carried into the assembled matrix. - Habec kernels moved to LinearSolvers/MLMG; template parameters renamed AlgInt/AlgGid. - Tests loop over the multigrid types; docs updated.
…MReX-Codes#6004) With its default -ffp-exception-behavior=ignore, clang may raise FP exceptions the source does not, e.g., by dividing in unused SIMD lanes. This made FFT_Poisson_2d/3d trap in the conda-forge single precision macOS builds in Geometry::computeRoundoffDomain, FFT::Poisson::solve and the test itself. Keep trapping on for other compilers and debug builds, where the test still catches real 0/0 in the solver. The command line can still turn trapping on. Document that trapping in optimized clang builds needs -ffp-exception-behavior=maytrap. Fix AMReX-Codes#6001.
…ting (AMReX-Codes#6006) In GPU builds, `HypreMLABecLap::setup` accumulates the fine-to-coarse (f2c) matrix entries into device memory that was never initialized: ```cpp Gpu::DeviceVector<Real> d_f2c_values(m_f2c_values[clev].size()); // not initialized ... hypmlabeclap_f2c_set_values(..., p_f2c_values + p_f2c_offset[icell], ...); // does (*values) += ... ``` On the host the same kernel writes into `m_f2c_values`, which is resized with value 0. On the device it adds to whatever the memory held before. This PR zero-initializes the device vector so that the GPU and CPU paths assemble the same matrix. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…X-Codes#5997) If the auxiliary row (x-faces) or column (y-faces) has a covered cell, use the other side, as the 3D version does. If both sides do, fall back to interpolating in the normal direction only, so covered values never enter.
…6003) The IJ-style matrices of AlgMG, hypre's IJ interface and PETSc now support the fourth-order Dirichlet boundary stencil, which couples to the cell two away from the boundary. A cell is never next to a boundary on both sides in one direction, so the stencil needs one extra entry per direction (3·dim+1 in total). MLMG no longer lowers maxorder to 3 for these solvers, except with hypre's structured and semi-structured interfaces. With embedded boundaries, hypre and PETSc now support maxorder 3, as AlgMG already did. They build EB rows with the same kernel. Previously MLMG lowered them to 2. The boundary order is now limited by the box length, as in MLMG's operator. This fixes a mismatch between matrix and operator for boxes of one or two cells at a Dirichlet face, which also affected hypre's structured interfaces at maxorder 3.
Host paths for the setup, with the same sums as before (bitwise-identical hierarchies and iterations): - amrex::RAP: R (A P) with A P built in chunks; with several ranks, only the rows of A P other processes need are formed for them. - SpGEMM: one-pass Gustavson on the split blocks without copies; optional per-row callback, used to truncate P while it is formed. - define_split, split_csr: per-row split, local part in place, no scan when all columns are local. - ext+i interpolation and truncation row by row; one-pass create_soc, also with remote columns. 128^3 Poisson setup: 1 rank 4.44 -> 2.33 s; 2/4/8 ranks 2.46/1.33/0.76 -> 1.50/0.82/0.46 s.
- Add comments to highlight sections on `MultiFab`, `iMultiFab`, and `MFIter` methods - Move `amrex_imultifab_build_owner_mask` from section on `MultiFab` methods to section on `iMultiFab` methods - Add method `amrex_imultifab_build_fine_mask`
…des#5993) In 2D, face types double as edge types, and build_faces promotes nearly full faces to regular. Edges cut at or near a node then stored the regular sentinel 1 while 3D stored the true centroid 0. Classify edges by level-set sign so 2D matches 3D. 3D output is unchanged. Also document the edge centroid convention at getEdgeCent. Closes AMReX-Codes#5625, closes AMReX-Codes#5624.
Each AlgMG case syncs the GPU hundreds of times, which is slow when several processes share one GPU (2 MPI ranks and ctest -j 2 on the GitLab H100 and MI300A jobs). On GPU builds, run the variations only for the checker problem (72 -> 21 cases) and mark the test RUN_SERIAL. The new input variation_problems selects the problems with variations. Repeat MLMG solves only if MLMG uses AlgMG, since the repeat checks that AlgMG is set up once.
…es#5656) ## Summary ParticleContainerToBlueprint historically walked only the compile-time SoA components, so any component added at runtime via `AddRealComp`/`AddIntComp` was silently dropped from the Blueprint output. This change allows runtime SoA components to be published when the caller provides names for all SoA components, while preserving the historical behavior for callers that provide names only for the compile-time components. Tested with MFIX-Exa against Ascent 0.9.5 / Conduit 0.9.7: runtime components are published with correct names and render correctly. ## Additional background ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate --------- Co-authored-by: Andrew Myers <atmyers2@gmail.com>
Arena::allocate_system now throws amrex::OutOfMemoryError (derived from std::bad_alloc) instead of printing and aborting. Its what() holds the old diagnostic. The sticky CUDA/HIP error left by a failed malloc is cleared. On GPUs, SpGEMM tries the vendor library first. If that runs out of memory, it falls back to multiplying blocks of rows, halving a block that still does not fit. The first pass only counts the nonzeros, C is allocated once, and the second pass recomputes each block into its slice of C. The peak is C plus one block's workspace. The vendor routines release their resources when an exception passes through.
…MReX-Codes#5980) ## Summary Adds an optional CUDA build setting, `GPU_MIN_BLOCKS=B` (GNU Make) or `-DAMReX_GPU_MIN_BLOCKS=B` (CMake), that limits registers per thread in AMReX kernels so that `B` blocks of `GPU_MAX_THREADS` threads fit on an SM. It is unset by default, and the default build is unchanged. `CUDA_MAXREGCOUNT` cannot do this: nvcc ignores `-maxrregcount` for kernels with `__launch_bounds__`, and all AMReX kernels launched through `launch_global<MT>` have launch bounds. When `B` is set: - `launch_global<MT>` uses `__launch_bounds__(MT, B*GPU_MAX_THREADS/MT)`, clamped to [1, 16]. Every launch path gets the cap: `ParallelFor` on boxes and MultiFabs, reductions, `Scan`, `CSR`, etc. Kernels whose block size differs from `GPU_MAX_THREADS` get the same register cap without exceeding the per-SM block limit. - `-maxrregcount` is lowered to `65536/(GPU_MAX_THREADS*B)`, rounded down to a multiple of 8 (or `CUDA_MAXREGCOUNT` if that is smaller). Device functions that are not inlined must fit in the cap of the kernels that call them; otherwise ptxas fails with RDC. CMake prints the effective value. For example, with `GPU_MAX_THREADS=256`, `B=2` caps registers at 128, `B=3` at 80 and `B=4` at 64. A larger `B` gives higher occupancy but can cause register spilling. `B*GPU_MAX_THREADS` should not exceed the GPU's limit of resident threads per SM (2048 on A100/H100, 1536 on sm_86/89/120); otherwise ptxas ignores the setting for kernels. CUDA only for now. HIP's `__launch_bounds__(MT, N)` means at least N waves per SIMD rather than N blocks per CU, so a HIP version needs to convert `B` and is left for a follow-up. Files: - `Src/Base/AMReX_GpuLaunchGlobal.H`: the cap in `launch_global<MT>`. - `Tools/CMake/AMReXCUDAOptions.cmake`, `AMReXSetDefines.cmake`, `AMReX_Config_ND.H.in`, `AMReXParallelBackends.cmake`: CMake option, input check, config define, effective `-maxrregcount`. - `Tools/GNUMake/Make.defs`: GNU Make define and `-maxrregcount`. - `Docs/sphinx_documentation/source/GPU.rst`: user documentation. ## Additional background Tested with CUDA 13.4 on an RTX 5070 (sm_120): - A register-heavy test kernel with `B=4` uses 64 registers through `ParallelFor(Box)`, MultiFab `ParallelFor` (128-thread blocks) and MultiFab `ParReduce`. Without the option it uses 254-255. - CMake build with `-DAMReX_GPU_MIN_BLOCKS=4`, particles, linear solvers and EB: no ptxas errors or warnings. `ctest` (Small): 5/5 pass. - CMake configure: the effective `-maxrregcount` is right for `B` unset, 1, 4, and 4 with `AMReX_CUDA_MAXREGCOUNT=32`. A non-integer `B` is rejected. - GNU Make flags checked with `print-NVCC_FLAGS` and `print-CPPFLAGS` for several combinations of `GPU_MIN_BLOCKS`, `GPU_MAX_THREADS` and `CUDA_MAXREGCOUNT`. The original version was also tested on an A100. ## Checklist The proposed changes: - [ ] fix a bug or incorrect behavior in AMReX - [x] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [x] include documentation in the code and/or rst files, if appropriate --------- Co-authored-by: Weiqun Zhang <weiqunzhang@lbl.gov>
## Summary Compile the single-Box GPU bottom solver of `MLEBNodeFDLaplacian` (AMReX-Codes#5654) only for `AMREX_SPACEDIM > 1`. In 1D, nvcc 12.2 (sm_80, GCC 11) hangs indefinitely on `AMReX_MLEBNodeFDLaplacian.cpp` (26s before AMReX-Codes#5654). The class does not support 1D anyway. ## Additional background Times out the CUDA CI in AMReX-Codes/pyamrex#620, which builds `AMReX_SPACEDIM="1;2;3"`. AMReX CUDA CI builds 3D only. nvcc 12.4 and 12.6 hang too; 12.9 is not affected. ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [ ] include documentation in the code and/or rst files, if appropriate 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…ds (AMReX-Codes#5999) With phi on centroids, a regular cell next to a cut cell needs the EB stencil, but tiles were classified on the tile box alone. Regular tiles bordering cut cells used the plain stencil, so results depended on max_grid_size and tile size (up to ~1.7% in a sphere test). Classify on the tile box grown by one cell instead.
PMIS drew its tie-breaking weights from the global random number stream, so repeated setups of the same matrix built different hierarchies. Use a hash of the global row index instead; the same matrix on the same ranks now always gives the same hierarchy. The AlgMG test's seed option is removed.
…odes#5996) ParallelForStrided(mf, stride, offset, f) visits the valid points whose index matches offset modulo stride, with one fused launch over all boxes. The GPU Gauss-Seidel smoothers launched every point once per color and masked out the others. MLNodeLaplacian and the red-black smoothers of MLPoisson and MLABecLaplacian now launch only the points of each color. Results are unchanged. On 3D GPU tests, NodalPoisson solves are 1.5-2x faster and cell-centered solves 7-21% faster.
MLNodeABecLaplacian: the GPU Jacobi sweep read neighbors of sol in the same kernel that updated sol. Replace it with the same 2^D-colored Gauss-Seidel the CPU path uses, via ParallelForStrided, and drop the Jacobi kernels. MLNodeTensorLaplacian: red-black Gauss-Seidel raced because the stencil couples same-parity diagonal neighbors. Sweep the 2^D index-parity classes with ParallelForStrided instead, on both GPU and CPU. Add color_offset for the parity offset of such sweeps; MLNodeLaplacian uses it too. Closes AMReX-Codes#5745
## Summary With the debug Microsoft C runtime (MSVC/clang-cl Debug builds), failed assertions and run-time check failures (`/RTC`) open a modal dialog box. In CI and batch jobs, nobody can close it: the run hangs forever without printing the error. WarpX disabled its MSVC CI for two years because of exactly that (BLAST-WarpX/warpx#7131). - User-facing: new runtime parameter `amrex.handle_crt_reports` (default: true), independent of `amrex.signal_handling`, which test suites (e.g., WarpX's and pyAMReX's pytest setups) often turn off. AMReX writes such reports with the rank to `stderr` and breaks into the debugger: an attached debugger stops at the failure; otherwise, by default, the process exits with `0x80000003` and Windows Error Reporting can write a crash dump. With AMReX's `SIGABRT` handler (`amrex.signal_handling` and `amrex.handle_sigabrt`), `abort()` raises `SIGABRT` without a CRT report. - Internal: new `AMReX_CrtReport.{H,cpp}` with CRT report hooks (`_CrtSetReportHook2`/`_CrtSetReportHookW2`) and `_set_abort_behavior`, set up in `amrex::Initialize` and undone in `amrex::Finalize`; the implementation is only compiled for `_MSC_VER && _DEBUG` (a no-op otherwise). ## Additional background <details><summary>Details, limitations, testing</summary> ### Details - Only errors and assertions whose report mode includes the dialog box (`_CRTDBG_MODE_WNDW`, the CRT default) are handled. Report modes without a dialog box that the application set, and warnings (`_CRT_WARN`), keep their handling. Handled reports replace the destinations configured together with the dialog box, and report hooks installed before `amrex::Initialize`. - The report is composed in a stack buffer and written with a single `WriteFile` to the standard error handle: no heap allocations (the debug heap itself reports, e.g., heap corruption), no CRT or AMReX calls that could report again (the rank is cached in `amrex::Initialize`), and no interleaving of reports from concurrent threads. If there is no standard error handle (e.g., GUI applications), only the break remains. Reports bypass `amrex::ErrorStream()` on purpose. Narrow reports are written as is if they are valid UTF-8 (incl. ASCII, e.g., `__FILE__` with `/utf-8`), else converted from the ANSI code page; wide-character reports (e.g., run-time check failures) are converted with `WideCharToMultiByte(CP_UTF8, ...)`. So the output is UTF-8, independent of the locale. Long messages are truncated at a code point boundary (`...`). - `<windows.h>` is only included in `AMReX_CrtReport.cpp`, like in `AMReX_OpenMP.cpp`. - While AMReX's `SIGABRT` handler is installed (`amrex.signal_handling` and `amrex.handle_sigabrt`), `_set_abort_behavior` disables the CRT report of `abort()`, so that `abort()` raises `SIGABRT` for that handler and `ParallelDescriptor::Abort`, as on other platforms. `amrex::Finalize` restores it together with the previous `SIGABRT` handler. - Like the signal handlers, the setup is global: repeated `amrex::Initialize` calls do not set it up again, and `amrex::Finalize` undoes it at its very end (after `MPI_Finalize`), also before a dynamically loaded AMReX library can be unloaded. - The CRT defaults apply with `amrex.handle_crt_reports = 0`, which can also be set via `AMREX_DEFAULT_INIT`. ### Known limitations - The handling is active from reading the runtime parameters in `amrex::Initialize` until the end of `amrex::Finalize`. Reports before or after, e.g., during static initialization or `MPI_Init`, use the report modes of the CRT. - Without a debugger, the system's post-mortem handling applies. On desktops, Windows Error Reporting can show a "has stopped working" dialog (if `DontShowUI` is not set), and a just-in-time debugger registered with `Auto=0` (e.g., by Visual Studio) can ask interactively. - Reports go to the standard error handle of the process. Where only AMReX's error stream is captured (a custom `a_oserr` in `amrex::Initialize`, or e.g. a Jupyter kernel), the report text is not shown. - A CRT report terminates the rank without `MPI_Abort`; the MPI launcher is expected to end the job. - The hooks act on the CRT that AMReX uses. With the shared debug CRT (`/MDd`, the CMake default), that is the CRT of all modules. With the static debug CRT (`/MTd`), modules other than AMReX have their own CRT and keep its report modes. - AMReX has no native backtrace on Windows (`Backtrace.N` files only contain the TinyProfiler call stack). The C++ call stack is available from the crash dump, e.g., `cdb -z <file.dmp> -lines -c ".ecxr; kpn; q"`. - AMReX's own Windows CI builds `RelWithDebInfo`, where this code is compiled out. It was tested via WarpX's MSVC Debug CI instead (see below). ### Testing WarpX's MSVC Debug CI job (VS 2026, MSVC 14.51) built against this PR (`ab79d9820e`, which differs from the head only in reading the parameter with `queryAdd`, ignoring failed report mode queries, the encoding handling of narrow reports, and documentation), without any other CRT workarounds, in BLAST-WarpX/warpx#7131: https://github.com/BLAST-WarpX/warpx/actions/runs/37186716581 - a regular WarpX test passes - `_CrtDbgReport(_CRT_ASSERT, ...)` after `amrex.initialize()`: `0::MSVC debug CRT assertion: Assertion failed: tmp test assertion !!!`, no dialog box - the same with `amrex.signal_handling=0`: handled the same way - the application set `_CRTDBG_MODE_FILE`: the CRT writes the report itself, AMReX does not interfere - `abort()`: AMReX's `SIGABRT` handler (`SIGABRT`, `Backtrace.0`, exit code 1) - `amrex.handle_crt_reports=0` (via `AMREX_DEFAULT_INIT`): dialog box (hangs), i.e., the CRT defaults apply - report after `amrex.finalize()`: dialog box (hangs), i.e., the hooks were removed - an injected `/RTCu` failure in WarpX prints `0::MSVC debug CRT error: ... The variable 'tmp_uninit' is being used without being initialized. !!!` and fails at the first time step. The crash dump shows the C++ stack (`_RTC_UninitUse` <- `FiniteDifferenceSolver::EvolveEPML` <- ... <- `WarpX::Evolve`). On Linux: - AMReX builds unchanged (CMake). - `AMReX_CrtReport.cpp` compiles warning-free with the MSVC Debug code enabled, against mocks of `crtdbg.h` and `windows.h` (GCC, `-Wall -Wextra -Wconversion -Wsign-conversion -Wshadow`). - A unit test of `AMReX_CrtReport.cpp` with these mocks, linked to AMReX, checks: only dialog-mode errors and assertions are handled; unknown report types are not handled; the rank prefix; UTF-8 output for narrow (UTF-8 and ANSI code page) and wide reports; long ASCII messages not truncated early; truncation (no split code points or surrogate pairs); no standard error handle; a report after `EndParallel`; repeated setup; and that finalizing restores the hooks and the `abort()` report. </details> ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [x] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [x] include documentation in the code and/or rst files, if appropriate 🤖 Generated with [Claude Code](https://claude.com/claude-code)
## Summary Make the `CArena` usage counters `m_used` and `m_actually_used` atomic, so `heap_space_used()` and `heap_space_actually_used()` can be called from another thread (e.g., a status report on a signal) without a data race. Before, they were read without `carena_mutex` while `alloc`/`free` update them. - User-facing: `heap_space_used()` and `heap_space_actually_used()` are thread-safe. - Internal: the counters are `std::atomic<std::size_t>`; updates still happen under `carena_mutex`. ## Additional background See awesome feature in BLAST-WarpX/warpx#7339 <details><summary>Testing</summary> ThreadSanitizer (GCC 13, CPU): a reader thread polls both getters while the main thread runs 20k `alloc`/`free`/`freeUnused` on a `CArena`: no reports. </details> ## Checklist The proposed changes: - [x] fix a bug or incorrect behavior in AMReX - [ ] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [x] include documentation in the code and/or rst files, if appropriate 🤖 Generated with [Claude Code](https://claude.com/claude-code)
Geometric MG failed for MLEBNodeFDLaplacian when cells were much shorter in one direction (e.g. dx/dz ~ 10). For operators that support it, MLMG now leaves uncoarsened every direction whose cells are at least 1.5 times as long as the shortest, until the cells are nearly cubic. No setting is needed; a fixed semicoarsening direction takes precedence. MLEBNodeFDLaplacian gets transfer kernels for levels coarsened in one direction, and restriction now works on the unscaled residual: rows next to the EB are scaled on every level, and restricting scaled residuals slowed MG near thin walls. The NodeEBFDLap test adds a case with 8:1 cells.
## Summary Add `TinyProfiler::CurrentName()`, the name of the innermost active profiler, which can be called from another thread while the profiled code runs (e.g., a watchdog, or a status report on a signal). Calling `PrintCallStack` from another thread is a data race on `ttstack` that can read freed memory. - User-facing: new `static const char* TinyProfiler::CurrentName () noexcept`, async-signal-safe, valid until `amrex::Finalize`; a stub in `AMReX_BLProfiler.H` returns `nullptr` without the tiny profiler. - Internal: `start()`/`stop()` publish a pointer to the (stable) `statsmap` key through a file-local lock-free atomic (release stores, acquire load). ## Additional background See awesome feature in BLAST-WarpX/warpx#7339 <details><summary>Known limitations</summary> - Profilers started by threads other than the OpenMP master thread are not tracked, as for `ttstack`. - Do not call `CurrentName()` concurrently with `amrex::Finalize`, which frees the names. - `nullptr` means "no active profiler", and also "tiny profiler disabled or not built". </details> <details><summary>Testing</summary> ThreadSanitizer (GCC 13, CPU, `AMReX_TINY_PROFILE=ON`), together with AMReX-Codes#6011 (merged; the reproducer also polls the `CArena` usage getters): a reader thread polls while the main thread runs 20k nested `BL_PROFILE` scopes and `CArena` allocations. Three runs each: - `CurrentName()`: no reports; returns `outer`, `inner`, and `nullptr` outside of profilers and after `Finalize`. - `PrintCallStack()` instead: 9-11 data race reports on `ttstack`. Stub: a caller of `amrex::TinyProfiler::CurrentName()` without `#ifdef` compiles (`-Wall -Wextra`) with the tiny profiler, without profiling, and with full profiling (`BL_PROFILING`). ```bash # AMReX with TSan cmake -S . -B build -DAMReX_SPACEDIM=3 -DAMReX_TINY_PROFILE=ON -DAMReX_MPI=OFF -DAMReX_OMP=OFF \ -DAMReX_FORTRAN=OFF -DAMReX_LINEAR_SOLVERS=OFF -DAMReX_PARTICLES=OFF -DAMReX_EB=OFF \ -DCMAKE_BUILD_TYPE=RelWithDebInfo -DCMAKE_CXX_FLAGS=-fsanitize=thread \ -DCMAKE_EXE_LINKER_FLAGS=-fsanitize=thread -DCMAKE_INSTALL_PREFIX=$PWD/install cmake --build build -j 6 && cmake --install build # reproducer (below), in its own directory cmake -S . -B build -DAMReX_DIR=<amrex>/install/lib/cmake/AMReX -DCMAKE_BUILD_TYPE=RelWithDebInfo \ -DCMAKE_CXX_FLAGS=-fsanitize=thread -DCMAKE_EXE_LINKER_FLAGS=-fsanitize=thread cmake --build build # setarch -R: TSan does not support high-entropy ASLR of newer kernels RACE_TEST_MODE=new TSAN_OPTIONS=halt_on_error=0 setarch $(uname -m) -R ./build/race_test amrex.verbose=0 RACE_TEST_MODE=old TSAN_OPTIONS=halt_on_error=0 setarch $(uname -m) -R ./build/race_test amrex.verbose=0 ``` <details><summary>Reproducer</summary> `CMakeLists.txt`: ```cmake cmake_minimum_required(VERSION 3.24) project(race_test CXX) find_package(AMReX REQUIRED COMPONENTS 3D TINYP) add_executable(race_test main.cpp) target_link_libraries(race_test PRIVATE AMReX::amrex_3d) ``` `main.cpp`: ```cpp // Reader thread polls profiler/arena state while the main thread runs. // Build with -fsanitize=thread; RACE_TEST_MODE=old uses PrintCallStack (racy baseline). #include <AMReX.H> #include <AMReX_Arena.H> #include <AMReX_BLProfiler.H> #include <AMReX_CArena.H> #include <AMReX_TinyProfiler.H> #include <algorithm> #include <atomic> #include <cstdlib> #include <iostream> #include <set> #include <sstream> #include <string> #include <thread> int main (int argc, char* argv[]) { const char* mode = std::getenv("RACE_TEST_MODE"); const bool old_api = (mode != nullptr && std::string(mode) == "old"); amrex::Initialize(argc, argv); { amrex::CArena carena; amrex::CArena const* arena = &carena; std::atomic<bool> stop{false}; std::set<std::string> seen; std::size_t max_used = 0; std::thread reader([&] () { while (!stop.load()) { if (old_api) { std::ostringstream os; amrex::TinyProfiler::PrintCallStack(os); } else { const char* name = amrex::TinyProfiler::CurrentName(); seen.insert(name ? name : "(none)"); } max_used = std::max(max_used, arena->heap_space_actually_used()); max_used = std::max(max_used, arena->heap_space_used()); } }); for (int i = 0; i < 20000; ++i) { BL_PROFILE("outer"); { BL_PROFILE("inner"); void* p = carena.alloc(1024 * (1 + i % 7)); carena.free(p); carena.freeUnused(); } } stop = true; reader.join(); std::cout << "names seen:"; for (auto const& s : seen) { std::cout << " " << s; } const char* after = amrex::TinyProfiler::CurrentName(); std::cout << "\nmax arena bytes seen: " << max_used << "\nafter loop: " << (after ? after : "(none)") << "\n"; } amrex::Finalize(); std::cout << "after Finalize: " << (amrex::TinyProfiler::CurrentName() ? "dangling!" : "(none)") << "\n"; } ``` </details> </details> ## Checklist The proposed changes: - [ ] fix a bug or incorrect behavior in AMReX - [x] add new capabilities to AMReX - [ ] changes answers in the test suite to more than roundoff level - [ ] are likely to significantly affect the results of downstream AMReX users - [x] include documentation in the code and/or rst files, if appropriate 🤖 Generated with [Claude Code](https://claude.com/claude-code)
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
14 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Additional background
Checklist
The proposed changes: