dT: check integrated state is finite, covering angular velocity as well - #79
Merged
Conversation
…s well
kT has two NaN tests already, and neither can see angular velocity. That is
not a placement mistake: angular velocity never crosses to kT. kT is handed
the translational speed array and the orientation quaternions, so no check
placed in kT can cover this at any strength. Hence a check on dT.
The two existing tests, stated precisely so this change is not mistaken for
a duplicate of either:
1. unpackMyBuffer() tests isfinite() on a max-reduced velocity. This one
is weak even on the translational side, because cub::DeviceReduce::Max
compares with >, every comparison against NaN is false, so NaN operands
are skipped and an all-NaN array reduces to the initial value -FLT_MAX,
which is finite and passes. Measured: a run whose whole state was NaN
reported "kinetic energy is nan while max speed is -3.40282e+38".
2. computeMarginFromAbsv() does a per-element isfinite() with
DEME_ABORT_KERNEL. This one is strong, and it is the real guard on the
translational side. But it reads marginSize, i.e. translational speed,
and it is skipped entirely when isExpandFactorFixed, since the margin is
then not derived from speed at all.
Angular velocity is where the corruption arrives first, which is what makes
the gap matter rather than merely exist. The defect fixed by #72 and #74
produced NaN through cross(cntPnt, F) with F exactly zero, and that lands in
angular acceleration via atomicAdd on alphaX
(DEMCollectForceKernels_Compact.cu:51 and :97), which integrates into angular
velocity, then through the orientation quaternion into positions, contact
geometry, forces, and only last into translational velocity.
Measured on a pre-fix tree with this check in place, 40 repetitions, one
Blackwell GPU, CUDA 12.9, one process per repetition, machine otherwise idle.
The reproducer sets no user expand factor, so the automatic margin route was
active and test 2 above was on duty throughout:
36 ran clean
3 this check reported angular velocity (x) as first non-finite
0 this check reported translational velocity as first non-finite
1 kT wedged in the contact-detection livelock before dT got another turn
Test 2 did not abort in any of those 40 runs, nor in 20 further sanitized
runs on current main. It was awake and the failures went past it, because all
of them were angular. In each of the three reports the last completed step
was healthy, at 0.570, 0.746 and 0.589 m/s with finite kinetic energy, so the
abort lands on the first step where the state goes bad rather than once the
corruption has surfaced. All four affected runs failed inside the
contact-array growth window, at 4366 to 9043 contacts.
checkStateIsFinite() probes the translational speed array and the three
angular components with sum reductions, which propagate NaN where a max does
not, and names which array failed. cubSumReduce<float, float> is already
instantiated, so no new template is introduced. Cost is at most four device
reductions per contact-detection cycle rather than per time step,
short-circuiting on the first bad array.
Verified not to fire on healthy runs: 40 repetitions on current main, zero
firings and zero failures. That property matters because the check aborts. An
injected NaN in angular velocity is reported by name, where without this
change the same injection dies with an unspecified launch failure out of
DEMCubContactDetection. An injected NaN in translational velocity is not
improved by this change and does not need to be, since test 2 catches it
first.
Known limitation: if kT wedges in the contact-detection livelock first, dT
waits on it and this check is starved rather than blind. That is the one
repetition above. The abort on non-finite bin coordinates added at
DEMBinSphereKernels.cu:63 covers that path from the kT side; the two are
complementary.
Not done here, deliberately: giving cubDEMMax a NaN-propagating functor would
also fix the clump_max_absv inspector, which still returns a finite number
for an all-NaN system, but that changes public inspector behaviour.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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
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.
kT has two NaN tests already, and neither can see angular velocity. That is not a placement mistake: angular velocity never crosses to kT. kT is handed the translational speed array and the orientation quaternions, so no check placed in kT can cover this at any strength. Hence a check on dT.
The two existing tests, stated precisely so this change is not mistaken for a duplicate of either:
unpackMyBuffer() tests isfinite() on a max-reduced velocity. This one is weak even on the translational side, because cub::DeviceReduce::Max compares with >, every comparison against NaN is false, so NaN operands are skipped and an all-NaN array reduces to the initial value -FLT_MAX, which is finite and passes. Measured: a run whose whole state was NaN reported "kinetic energy is nan while max speed is -3.40282e+38".
computeMarginFromAbsv() does a per-element isfinite() with DEME_ABORT_KERNEL. This one is strong, and it is the real guard on the translational side. But it reads marginSize, i.e. translational speed, and it is skipped entirely when isExpandFactorFixed, since the margin is then not derived from speed at all.
Angular velocity is where the corruption arrives first, which is what makes the gap matter rather than merely exist. The defect fixed by #72 and #74 produced NaN through cross(cntPnt, F) with F exactly zero, and that lands in angular acceleration via atomicAdd on alphaX
(DEMCollectForceKernels_Compact.cu:51 and :97), which integrates into angular velocity, then through the orientation quaternion into positions, contact geometry, forces, and only last into translational velocity.
Measured on a pre-fix tree with this check in place, 40 repetitions, one Blackwell GPU, CUDA 12.9, one process per repetition, machine otherwise idle. The reproducer sets no user expand factor, so the automatic margin route was active and test 2 above was on duty throughout:
36 ran clean
3 this check reported angular velocity (x) as first non-finite
0 this check reported translational velocity as first non-finite
1 kT wedged in the contact-detection livelock before dT got another turn
Test 2 did not abort in any of those 40 runs, nor in 20 further sanitized runs on current main. It was awake and the failures went past it, because all of them were angular. In each of the three reports the last completed step was healthy, at 0.570, 0.746 and 0.589 m/s with finite kinetic energy, so the abort lands on the first step where the state goes bad rather than once the corruption has surfaced. All four affected runs failed inside the contact-array growth window, at 4366 to 9043 contacts.
checkStateIsFinite() probes the translational speed array and the three angular components with sum reductions, which propagate NaN where a max does not, and names which array failed. cubSumReduce<float, float> is already instantiated, so no new template is introduced. Cost is at most four device reductions per contact-detection cycle rather than per time step, short-circuiting on the first bad array.
Verified not to fire on healthy runs: 40 repetitions on current main, zero firings and zero failures. That property matters because the check aborts. An injected NaN in angular velocity is reported by name, where without this change the same injection dies with an unspecified launch failure out of DEMCubContactDetection. An injected NaN in translational velocity is not improved by this change and does not need to be, since test 2 catches it first.
Known limitation: if kT wedges in the contact-detection livelock first, dT waits on it and this check is starved rather than blind. That is the one repetition above. The abort on non-finite bin coordinates added at DEMBinSphereKernels.cu:63 covers that path from the kT side; the two are complementary.
Not done here, deliberately: giving cubDEMMax a NaN-propagating functor would also fix the clump_max_absv inspector, which still returns a finite number for an all-NaN system, but that changes public inspector behaviour.