Skip to content

dT: check integrated state is finite, covering angular velocity as well - #79

Merged
Ruochun merged 1 commit into
mainfrom
fix-nan-blind-state-check
Sep 6, 2026
Merged

dT: check integrated state is finite, covering angular velocity as well#79
Ruochun merged 1 commit into
mainfrom
fix-nan-blind-state-check

Conversation

@DanNegrut

Copy link
Copy Markdown
Collaborator

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.

…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>
@DanNegrut
DanNegrut requested a review from Ruochun September 5, 2026 15:56
@Ruochun
Ruochun merged commit 12f13cb into main Sep 6, 2026
7 checks passed
@Ruochun
Ruochun deleted the fix-nan-blind-state-check branch September 6, 2026 07:20
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.

2 participants