solve_driver.hpp File Reference

dem — the shared contact-solve driver: the full modern velocity + position solve sequence (warm-started colored PGS with persistent contacts, gravity statics / stabilization passes, friction, colored-GS overlap projection, adaptive stops), extracted from demStep so the single-GPU and the distributed (MPI) steps run ONE driver instead of drifting copies. More...

#include <algorithm>
#include <cstdlib>
#include <Kokkos_Core.hpp>
#include <stdexcept>
#include <string>
#include <utility>
#include <vector>
#include "broadphase_arborx.hpp"
#include "contact_preprocessing.hpp"
#include "particles.hpp"
#include "solve_copies.hpp"
#include "solver_friction.hpp"
#include "solver_multilevel.hpp"
#include "solver_position.hpp"
#include "solver_velocity.hpp"

Go to the source code of this file.

Classes
struct	peclet::dem::CudaIterGraph
struct	peclet::dem::RankK
	Rank-level M (docs/contact_solve_framework.md §13.3) as the hooks expose it to the driver: whether this rank reconciles copies with neighbours at all (it sends or receives ghosts). More...
struct	peclet::dem::SoloSolveHooks
	Single-GPU hooks: no ghost refresh, residuals are already global. Everything inlines away. More...

Namespaces
namespace	peclet
namespace	peclet::dem

Macros
#define	PECLET_DEM_GRAPH_LOOP(useVar, graphVar, emitVar, slotVar)   (void)graphVar;

Functions
int	peclet::dem::readInt (Kokkos::View< int, CpMem > v)
float	peclet::dem::readFloat (Kokkos::View< float, CpMem > v)
void	peclet::dem::fillGidBaseKokkos (Vi gid, int n, int base)
	gid(i) = base + i over [0, n) — the per-rank global-id re-base of the distributed step (namespace scope: nvcc forbids KOKKOS_LAMBDA in member functions).
float	peclet::dem::maxOwnedRadius (const Particles &P)
	Largest effective particle radius over the owned set (= max scale × globalScale, growth included).
void	peclet::dem::fillWorldRadiiKokkos (Vf scale, Vf rad, float gs, float bR, int n)
	World radii rad(i) = scale(i) * globalScale * baseRadius over [0, n): the owned span before a step, owned + ghosts after periodic ghost generation or a halo gather.
int	peclet::dem::findCollisionsGrow (Particles &P, float margin)
	Broad phase with an automatically-grown pair buffer.
int	peclet::dem::findCollisionsVerlet (Particles &P, float margin, float maxRad)
	Verlet-cached impulse broadphase (single-GPU, non-periodic).
int *	peclet::dem::demIterCountSlot (Particles &P, int slot)
	Diagnostics (docs/contact_solve_framework.md §12 S12): the device scalar a fused main loop writes its iteration count to (slot 0 velocity, 1 position), allocated on first use.
int	peclet::dem::demReadIterCount (const Particles &P, int slot)
	The count a fused main loop wrote (one small readback, only with P.iterCounters on); -1 with the counters off (no fence, no copy).
template<class Hooks >
void	peclet::dem::demSolveContacts (Particles &P, int nc, int nm, int nBodies, Kokkos::View< const int *, CpMem > keyIdx, const Hooks &hooks)
	One full velocity + position contact solve over the already-built contacts/manifolds (see file comment).

Detailed Description

dem — the shared contact-solve driver: the full modern velocity + position solve sequence (warm-started colored PGS with persistent contacts, gravity statics / stabilization passes, friction, colored-GS overlap projection, adaptive stops), extracted from demStep so the single-GPU and the distributed (MPI) steps run ONE driver instead of drifting copies.

The two callers differ only through the Hooks policy:

• SoloSolveHooks (single-GPU): every hook is a no-op / identity — the driver compiles to exactly the pre-extraction demStep sequence (validated bit-for-bit on the Serial backend).
• MpiSolveHooks (step_solve_mpi.hpp, PECLET_DEM_MPI): processor-block Gauss–Seidel with a SINGLE owner per contact (docs/mpi_momentum_conservation.md) — nc / nm are the counts of the contacts / manifolds this rank owns, and only they are swept, over owned + ghost body slots; the partner half of every impulse lands in the ghost slot. beginSolve marks the ghosts' velocity baselines; syncVelocities / syncPositions reverse-accumulate each ghost's change since its baseline onto its owner, then refresh owner->ghost (reverse, then forward), every syncEvery iterations plus once after every solve phase; publishPositions opens the position phase with a forward only (the integration of a ghost is not an interaction); syncFrictionCounts makes the legacy-friction counts and syncContactCounts the Jacobi count-averaging counts the serial ones. visibleManifolds(nm) is the whole visible range, read only by the label passes (warm-ledger match, grounded / height levels). allMax turns each adaptive-stop residual into a global MPI_Allreduce(MAX) so all ranks take the same break (a rank-local break would desynchronise the collective syncs and deadlock); allMaxAny folds a rank-local vote (the colouring-invariant check, §12 S3) into that same Allreduce.

nBodies is the body-slot span of the solve graph: numReal on the single-GPU path (ghost slots are realIndices-mapped onto their owners), numReal + numGhost under MPI (ghosts are self-mapped slots solved in place). keyIdx maps a body slot to the identity used in the persistent-pair keys: realIndices on the single-GPU path, the global particle id under MPI (local slots are not stable across halo rebuilds / migration).

Definition in file solve_driver.hpp.

Macro Definition Documentation

PECLET_DEM_GRAPH_LOOP

#define PECLET_DEM_GRAPH_LOOP	(		useVar,
			graphVar,
			emitVar,
			slotVar
	)		(void)graphVar;

Definition at line 233 of file solve_driver.hpp.