Auditing src/chrono_fsi/sph turned up one root cause with two user-visible consequences,
plus an unrelated dead error path. All three have fixes with tests, linked below.
Everything here was verified against main at 05de72f4c, built with gcc 12.2 / ROCm 7.2
on an MI350X (gfx950), Release, in both CH_USE_SPH_DOUBLE=OFF and ON.
Root cause
Device kernels read a per-translation-unit __constant__ ChFsiParamsSPH paramsD. It is
uploaded only during Initialize() (SphFluidDynamics.cu, SphCollisionSystem.cu,
SphForceWCSPH.cu, SphForceISPH.cu, SphBceManager.cu), and no code path re-uploads it
afterwards. Anything that writes m_paramsH after Initialize() therefore has no effect
on the simulation.
Consequence 1: configuration setters are silent no-ops after Initialize()
About thirty public setters on ChFsiFluidSystemSPH write only m_paramsH. Called after
Initialize(), they are accepted and do nothing, with no warning.
Demonstration (periodic fluid box, zero gravity): request
SetGravitationalAcceleration(0, 0, -9.81) after Initialize() and step 200 times, then
compare against a run where nothing was requested. Particle positions are identical to the
last bit (max |dpos| = 0 m). The same value set before Initialize() produces free fall
(max |v| = 0.392401 m/s, against g*t = 0.3924).
Fix: #782.
Consequence 2: the CRMTerrain moving patch does not work
CRMTerrain::Synchronize shifts the computational domain mid-run through
SetComputationalDomain (CRMTerrain.cpp:166). That shift never reaches the device, so
neighbor-search binning keeps using the initial domain and calcHashD aborts the run once
the patch travels roughly 40 kernel lengths beyond it.
Demonstration (headless port of demo_VEH_CRMTerrain_MovingPatch: same terrain, same
tracked body at 5 km/h, just driven farther): the run dies at about 1.5 m of travel with
Error flag intercepted in .../SphCollisionSystem.cu:344 from calcHashD
The bundled demo does not surface this because its body never travels that far. With the
fix, the same test completes 4.17 m with 34 patch relocations.
Fix: #783.
Unrelated: the rheology-failure check can never fire
SphFluidDynamics.cu declares a host stack bool error_occurred, passes it by
reference into the __global__ kernels EulerStep_D / MidpointStep_D and on into
TauEulerStep, then reads it back on the host and calls
gpuThrowError("Rheology model failed"). Device code cannot write a host stack variable
and never assigns it, so that error is unreachable: within the m_check_errors block,
positions and densities have working thrust::any_of NaN scans, but the updated CRM
stress state has no check at all. That is the state that went non-finite in the MCC
overflow fixed by #766.
Fix: #784.
Other findings from the same audit (no PRs; happy to split these into their own issues)
FsiDataManager::AddSphParticle (SphDataManager.cu:220) derives the MCC specific
volume as Sv = mcc_v_lambda - mcc_lambda*log(pc/p1) + mcc_kappa*log(pc/confining).
Any CRM bed built through ChFsiProblemSPH with the default particle pressure has
pc = 0 and zero confining stress, so every particle stores a non-finite Sv; with
those defaults it is NaN, since the second term is log(0/0). Harmless under mu(I),
which never reads it, but it poisons the consolidation state for MCC beds built that
way. Two in-code TODOs already flag the fragility of this initialization.
- Related, in the MCC branch of
TauEulerStep: both consolidation components updated
there are clamped at the end of the update (pcEvSv.x = fmax(100, ...) and
pcEvSv.z = fmax(1, ...)). Since fmax returns the non-NaN operand, the non-finite
specific volume above is silently replaced by the clamp floor rather than reported, so
such a bed runs on with an arbitrary specific volume instead of a derived one.
Chrono_vehicle_cosim does not compile in HIP builds:
ChVehicleCosimTerrainNodeGranularSPH.cpp is compiled as plain C++ but includes
chrono_fsi/sph/utils/SphUtilsPrint.cuh, which pulls in rocprim/HIP headers without the
HIP platform defines ("Must define exactly one of __HIP_PLATFORM_AMD__ or
__HIP_PLATFORM_NVIDIA__"). Chrono_vehicle and ChronoModels_vehicle build fine.
- An exception thrown during the CFD phase of
ChFsiSystem::DoStepDynamics is masked: the
MBS std::thread is still joinable while the stack unwinds, so the process hits
std::terminate and the message is lost. In coupled runs every FSI error surfaces as a
bare Aborted (core dumped); we needed a debugger to read the messages quoted above.
Auditing
src/chrono_fsi/sphturned up one root cause with two user-visible consequences,plus an unrelated dead error path. All three have fixes with tests, linked below.
Everything here was verified against
mainat05de72f4c, built with gcc 12.2 / ROCm 7.2on an MI350X (gfx950), Release, in both
CH_USE_SPH_DOUBLE=OFFandON.Root cause
Device kernels read a per-translation-unit
__constant__ ChFsiParamsSPH paramsD. It isuploaded only during
Initialize()(SphFluidDynamics.cu,SphCollisionSystem.cu,SphForceWCSPH.cu,SphForceISPH.cu,SphBceManager.cu), and no code path re-uploads itafterwards. Anything that writes
m_paramsHafterInitialize()therefore has no effecton the simulation.
Consequence 1: configuration setters are silent no-ops after Initialize()
About thirty public setters on
ChFsiFluidSystemSPHwrite onlym_paramsH. Called afterInitialize(), they are accepted and do nothing, with no warning.Demonstration (periodic fluid box, zero gravity): request
SetGravitationalAcceleration(0, 0, -9.81)afterInitialize()and step 200 times, thencompare against a run where nothing was requested. Particle positions are identical to the
last bit (max |dpos| = 0 m). The same value set before
Initialize()produces free fall(max |v| = 0.392401 m/s, against g*t = 0.3924).
Fix: #782.
Consequence 2: the CRMTerrain moving patch does not work
CRMTerrain::Synchronizeshifts the computational domain mid-run throughSetComputationalDomain(CRMTerrain.cpp:166). That shift never reaches the device, soneighbor-search binning keeps using the initial domain and
calcHashDaborts the run oncethe patch travels roughly 40 kernel lengths beyond it.
Demonstration (headless port of
demo_VEH_CRMTerrain_MovingPatch: same terrain, sametracked body at 5 km/h, just driven farther): the run dies at about 1.5 m of travel with
The bundled demo does not surface this because its body never travels that far. With the
fix, the same test completes 4.17 m with 34 patch relocations.
Fix: #783.
Unrelated: the rheology-failure check can never fire
SphFluidDynamics.cudeclares a host stackbool error_occurred, passes it byreference into the
__global__kernelsEulerStep_D/MidpointStep_Dand on intoTauEulerStep, then reads it back on the host and callsgpuThrowError("Rheology model failed"). Device code cannot write a host stack variableand never assigns it, so that error is unreachable: within the
m_check_errorsblock,positions and densities have working
thrust::any_ofNaN scans, but the updated CRMstress state has no check at all. That is the state that went non-finite in the MCC
overflow fixed by #766.
Fix: #784.
Other findings from the same audit (no PRs; happy to split these into their own issues)
FsiDataManager::AddSphParticle(SphDataManager.cu:220) derives the MCC specificvolume as
Sv = mcc_v_lambda - mcc_lambda*log(pc/p1) + mcc_kappa*log(pc/confining).Any CRM bed built through
ChFsiProblemSPHwith the default particle pressure haspc = 0and zero confining stress, so every particle stores a non-finiteSv; withthose defaults it is NaN, since the second term is
log(0/0). Harmless under mu(I),which never reads it, but it poisons the consolidation state for MCC beds built that
way. Two in-code TODOs already flag the fragility of this initialization.
TauEulerStep: both consolidation components updatedthere are clamped at the end of the update (
pcEvSv.x = fmax(100, ...)andpcEvSv.z = fmax(1, ...)). Sincefmaxreturns the non-NaN operand, the non-finitespecific volume above is silently replaced by the clamp floor rather than reported, so
such a bed runs on with an arbitrary specific volume instead of a derived one.
Chrono_vehicle_cosimdoes not compile in HIP builds:ChVehicleCosimTerrainNodeGranularSPH.cppis compiled as plain C++ but includeschrono_fsi/sph/utils/SphUtilsPrint.cuh, which pulls in rocprim/HIP headers without theHIP platform defines ("Must define exactly one of
__HIP_PLATFORM_AMD__or__HIP_PLATFORM_NVIDIA__").Chrono_vehicleandChronoModels_vehiclebuild fine.ChFsiSystem::DoStepDynamicsis masked: theMBS
std::threadis still joinable while the stack unwinds, so the process hitsstd::terminateand the message is lost. In coupled runs every FSI error surfaces as abare
Aborted (core dumped); we needed a debugger to read the messages quoted above.