Skip to content

CI + NIGHTLY for pixi #518

Open
AndreiIulianMaftei wants to merge 184 commits into
ci-pixi-wheel-releasefrom
develop
Open

CI + NIGHTLY for pixi #518
AndreiIulianMaftei wants to merge 184 commits into
ci-pixi-wheel-releasefrom
develop

Conversation

@AndreiIulianMaftei

Copy link
Copy Markdown
Collaborator

Added ci pipeline for pixi wheels
Added nightly releases.

greole and others added 30 commits May 28, 2026 21:12
…sh-data

Fix the issue of duplicated boundary mesh data in unstructured mesh
Port the distributed logging functionality from stack/distributed:

- logging.hpp: make setNeonDefaultPattern MPI-environment aware,
  add debug()/error() convenience helpers and a terminate() that
  prints a stack trace and aborts MPI ranks.
- logging.cpp: mute log output on non-zero MPI ranks and implement
  terminate() using cpptrace + MPI_Abort under debug messaging.
- error.hpp: emit cpptrace stack traces from NF_ERROR_EXIT and
  include file/line in assertion failures.
- initialization.hpp: register the cpptrace terminate handler and
  pass the MPI environment to the default logging pattern setup.
- include/CMakeLists.txt: link MPI::MPI_CXX into the public API so
  the MPI-aware logging headers resolve.
greole and others added 30 commits June 21, 2026 12:56
Implement the linearUpwind surface-interpolation scheme, following
OpenFOAM's corrected-upwind scheme:

    phi_f = phi_U + (Cf - C_U) & grad(phi)_U

where U is the upwind cell (owner if faceFlux >= 0, neighbour otherwise).
interpolate() returns the gradient-corrected face value (used by the
explicit Gauss divergence) and weight() returns the upwind weights (used
by the implicit matrix assembly) - the same split OpenFOAM uses, where
the gradient correction is a deferred/explicit term.

The cell gradient is computed internally with Gauss-Green (Vec3 for a
scalar field, Tensor for a vector field). With NeoN's row-major tensor
convention (gradTensor(i,j)=d U_i/d x_j) the product `grad & d`
reproduces OpenFOAM's `(Cf - C) & gradVf` for both ranks.

GeometryScheme now caches the per-internal-face owner/neighbour-to-face
delta vectors (Cf - C), computed before reset() frees the mesh
cell/face centres.

Tests (CPU + GPU): reduces to upwind for a constant field; exact
reconstruction of a linear scalar and a linear vector field for both
upwind directions; cached geometry delta-vector magnitudes.
Wire the linearUpwind gradient correction into the implicit divergence,
matching OpenFOAM's fvmDiv (upwind matrix + the correction added to the
source, fvm += fvc::surfaceIntegrate(faceFlux*correction)).

- SurfaceInterpolationFactory gains corrected()/correction() hooks
  (default: uncorrected / zero), forwarded by the SurfaceInterpolation
  wrapper. linearUpwind overrides them; linear/upwind are unchanged.
- GaussGreenDiv::div(ls, ...) adds the deferred correction to the rhs
  for corrected schemes: -operatorScaling * sum_f F_f corr_f per cell.
- linearUpwind interpolate()/correction() now share one kernel via a
  runtime withSrcValue flag (interpolate = upwind value + correction,
  correction = correction only). A runtime flag is used instead of
  `if constexpr` because nvcc forbids first-capturing a variable inside
  an `if constexpr` block in an extended device lambda.

Validated against OpenFOAM fvmDiv: applyOperator(ls, x) == (M & x) * V
(test/test_implicitOperators, all executors).
Handle the two remaining linearUpwind gaps for distributed / fused solves:

Processor faces: a proc face is a cut internal face whose neighbour lives on
another rank. Each rank computes its owner-side correction grad[own] & (Cf-C[own])
locally and exchanges it (BoundaryData::communicate + waitAll); the neighbour's
owner-side correction is exactly this face's neighbour-side correction. The upwind
cell and the rhs sign use the outward flux normalSign*F (boundaryMesh weights), so
non-owner-side (flipped-normal) proc faces are handled correctly. correction() fills
the proc slots and GaussGreenDiv::div(ls,...) adds them to the owner-cell rhs,
reproducing the global internal-face correction across the rank boundary.

Fused operator: GaussGreenDivLaplacian now accepts 'Gauss linearUpwind' (in addition
to upwind) and adds the same deferred correction to the rhs in both implicitOperation
overloads. addDivCorrectionToRhs is shared via gaussGreenDiv.hpp.

Validated with a distributed test (mpirun -n 3): the partitioned correction
rhs(linearUpwind)-rhs(upwind) matches the global serial slice on every rank.
Enable solver and off-diagonal block caching
… nvcc

nanobind adds -Os to the _neon module; nvcc cannot parse it and aborts with
"nvcc fatal : 's': expected a number" when the binding is compiled through
nvcc_wrapper/kokkos_launch_compiler in a CUDA build.

Passing NOMINSIZE to nanobind_add_module suppresses that -Os, so the
optimization level is left to the build type (CMAKE_CXX_FLAGS_<CONFIG>):
-O3 in Release, -O0 in Debug, etc. — no hard-coded -O3 override.

Claude-Session: https://claude.ai/code/session_01ATmK2Ms4yxdK5W6QtoEHMy
…ffer

pointer in O(1) without submitting any GPU kernel.

The copy operator=(const Vector&) uses setContainer which dispatches an
async parallelFor kernel. When the source is a temporary (e.g. the result
of fromFoamField in constructFrom), the SYCL runtime frees the source
buffer via sycl::free(ptr, context) at end-of-expression before the
pending kernel executes, causing a GPU page fault on Intel PVC.

With the move assignment, assigning a temporary Vector now resolves to a
pointer swap. C++ overload resolution automatically prefers operator=(&&)
for rvalues, so all constructFrom call sites in readers.hpp (internalVector
and boundaryData.value() assignments) are fixed without any changes there.

+ additional clean-up

Fix distributed linearUpwind deferred correction at processor boundaries

Two bugs caused the linearUpwind deferred correction to be wrong on cells
adjacent to processor patch faces.

1. BoundaryData::waitAll() CPU path reallocated value_, invalidating any
   View<T> obtained via value() before the MPI exchange. In
   applyLinearUpwind the selection kernel read/wrote through a dangling
   pointer to freed memory, so the proc-face correction always reflected
   the received (neighbour) value regardless of the upwind direction.
   Fix: write received data directly into the existing value_ buffer
   without reallocation, keeping all live views valid.

2. GaussGreenGrad::grad() used in computeLinearUpwind linearly interpolates
   proc-patch faces using ghost-cell values from src.boundaryData(). Without
   a preceding halo exchange those slots were stale, giving a wrong gradient
   at proc-adjacent cells and therefore a wrong correction everywhere.
   Fix: call correctBoundaryConditions() on src when the mesh has proc faces
   before computing the gradient.
Add slip and inletOutlet boundary conditions
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants