analysis: no sort ever sees a NaN key, and FFT kernel launches are checked
Build Packages / Create release (push) Successful in 16s
Build Packages / build:rugnux:aarch64 (cross) (push) Successful in 7m4s
Build Packages / build:rugnux-tgz (x86_64) (push) Successful in 8m40s
Build Packages / build:viewer-tgz:cpu (push) Successful in 10m54s
Build Packages / build:viewer-tgz:cuda (push) Successful in 11m40s
Build Packages / build:rpm (rocky8_nocuda) (push) Successful in 15m30s
Build Packages / build:rpm (rocky9_nocuda) (push) Successful in 16m21s
Build Packages / build:windows:nocuda (push) Successful in 17m15s
Build Packages / build:windows:cuda (push) Successful in 19m56s
Build Packages / build:rpm (ubuntu2204_nocuda) (push) Successful in 15m10s
Build Packages / HDF5 consumer tests (DIALS, XDS) (push) Successful in 24m33s
Build Packages / build:rpm (ubuntu2404_nocuda) (push) Successful in 17m31s
Build Packages / Generate python client (push) Successful in 45s
Build Packages / build:rugnux:windows (push) Successful in 11m0s
Build Packages / Build documentation (push) Successful in 1m21s
Build Packages / build:rpm (rocky8_sls9) (push) Successful in 19m34s
Build Packages / build:rpm (rocky9_sls9) (push) Successful in 19m30s
Build Packages / build:rpm (rocky8) (push) Successful in 17m7s
Build Packages / build:rpm (rocky9) (push) Successful in 17m50s
Build Packages / build:rpm (ubuntu2204) (push) Successful in 14m50s
Build Packages / build:rpm (ubuntu2404) (push) Successful in 13m28s
Build Packages / Unit tests (push) Successful in 1h39m47s

A production broker segfaulted inside std::sort in FilterFFTResults: a NaN
length key violates strict weak ordering, and libstdc++'s unguarded partition
scan then walks off the array (confirmed from the deployed binary's faulting
instruction). No legitimate producer of that NaN exists - both FFT back-ends
emit finite lengths - so the row was corrupted, most plausibly via the one
gap in the path: neither kernel launch in FFTIndexerGPU::ExecuteFFT was
error-checked, so a failed launch silently hands back uninitialised device
memory. An audit of every sort/nth_element site then found two more places
where a NaN key is reachable by construction.

- FFTIndexerGPU: cudaGetLastError after both kernel launches (the idiom
  every other GPU translation unit already follows).
- FilterFFTResults: drop non-finite rows when building the magnitude map;
  bit-identical in normal operation.
- SearchSpaceGroup: PearsonCC deliberately returns NaN for an unscorable
  operator (n_pairs < 2, zero variance) and the score sort consumed it
  unfiltered - ~24 operators on a cubic holohedry is past the introsort
  threshold, the same crash waiting to happen. Unscorable operators now
  rank last under a well-defined comparator.
- PostIndexingRefinement: a singular QR-solved cell puts inf into
  cell.inverse() and 0*inf = NaN into the residual norms fed to
  nth_element; non-finite distances now map to +inf, which says exactly
  "this spot does not index" and orders consistently.
- ModelScaling: a NaN R factor from the first grid point latched into
  best_r and won every later comparison; it is now skipped.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
This commit is contained in:
2026-09-16 14:05:49 +02:00
co-authored by Claude Opus 5
parent bef0c6f390
commit bfe95b4ede
6 changed files with 25 additions and 3 deletions
+1
View File
@@ -3,6 +3,7 @@
### 1.0.0-rc.170
* Fixed a `jfjoch_broker` crash during indexing: a non-finite value reaching the FFT shortlist, space-group operator ranking or indexing-refinement sorts no longer causes undefined behaviour, and GPU FFT indexer kernel launches are now error-checked.
* rugnux needs about a third less peak memory to scale, merge and post-refine rotation data with tens of millions of partial observations, with identical results.
* `rugnux --model`: the placed coordinate file carries the space group its own coordinates obey, and says so when that is not the group the reflection files beside it carry.
* `jfjoch_viewer`: fixes in the dataset plots (the stacked plots line up their y axes, the horizontal crosshair follows the cursor, a run in progress labels the angles of its whole sweep), in the inspector (a long dataset name wraps instead of being cut off, a live session shows its sample-head angles) and in the layout (the resize handle beside the file panel is centred).
+8 -1
View File
@@ -2,6 +2,7 @@
// SPDX-License-Identifier: GPL-3.0-only
#include <algorithm>
#include <cmath>
#include <numeric>
#include "../../common/JFJochMath.h"
#include "FFTIndexer.h"
@@ -185,8 +186,14 @@ std::vector<Coord> FFTIndexer::FilterFFTResults(size_t max_vectors,
std::vector<float> *magnitudes) const {
std::multimap<float, FFTResult> fft_result_map;
for (int i = 0; i < direction_vectors.size(); i++)
// No legitimate row is non-finite (lengths are bin*coeff or -1, magnitudes guarded
// prominences), but a corrupted result buffer would put NaN into the length sort below,
// and std::sort with a NaN key walks out of bounds. Drop such rows instead.
for (int i = 0; i < direction_vectors.size(); i++) {
if (!std::isfinite(result_fft[i].magnitude) || !std::isfinite(result_fft[i].length))
continue;
fft_result_map.insert(std::make_pair(result_fft[i].magnitude, result_fft[i]));
}
std::vector<FFTResult> fft_result_filtered;
int count = 0;
+2
View File
@@ -226,6 +226,7 @@ void FFTIndexerGPU::ExecuteFFT(const std::vector<Coord> &coord, size_t nspots) {
direction_vectors.size(),
d_input_fft);
}
cuda_err(cudaGetLastError());
cuda_err(cufftExecR2C(plan, d_input_fft, d_output_fft));
@@ -238,6 +239,7 @@ void FFTIndexerGPU::ExecuteFFT(const std::vector<Coord> &coord, size_t nspots) {
max_length_A, min_length_A, histogram_size,
bg_half,
direction_vectors.size(), d_result_fft);
cuda_err(cudaGetLastError());
cuda_err(cudaMemcpyAsync(result_fft.data(), d_result_fft, direction_vectors.size() * sizeof(FFTResult),
cudaMemcpyDeviceToHost, stream));
@@ -5,6 +5,7 @@
#include "PostIndexingRefinement.h"
#include <iostream>
#include <limits>
#include <thread>
namespace {
@@ -118,6 +119,10 @@ namespace {
resid = CalculateResiduals(spots, cell);
ArrayX<float> dist = resid.rowwise().norm();
// A singular refined cell puts inf into cell.inverse() and 0*inf = NaN into the
// residuals; a NaN key is UB in nth_element. Such a spot does not index at all,
// which is exactly what +inf says - and inf, unlike NaN, orders consistently.
dist = dist.isFinite().select(dist, std::numeric_limits<float>::infinity());
auto nth = std::begin(dist) + (cifssr.min_spots - 1);
std::nth_element(std::begin(dist), nth, std::end(dist));
scores(j) = *nth;
@@ -1619,7 +1619,13 @@ SearchSpaceGroupResult SearchSpaceGroup(
for (const auto& [rk, s] : op_cache)
result.operator_scores.push_back(s);
std::sort(result.operator_scores.begin(), result.operator_scores.end(),
[](const auto& a, const auto& b) { return a.cc > b.cc; });
[](const auto& a, const auto& b) {
// cc is NaN for an operator PearsonCC could not score; a NaN key breaks
// std::sort's ordering (out-of-bounds UB), so rank those last explicitly.
if (std::isfinite(a.cc) != std::isfinite(b.cc))
return std::isfinite(a.cc);
return a.cc > b.cc;
});
// A caller that already decided the point group elsewhere overrides the choice here, keeping the
// operator scores and the refusal report Stage A just produced. Nothing else is bypassed: Stage B
+2 -1
View File
@@ -54,7 +54,8 @@ ModelScaleReport FitModelScale(gemmi::Scaling<float> &scaling, ModelScaleBox box
scaling.fit_parameters(); // k_overall + anisotropic B only
++report.n_grid;
const double r = RFactor(scaling);
if (best_r < 0 || r < best_r) {
// A diverged fit gives r = NaN; latched as best_r it wins every later r < best_r.
if (std::isfinite(r) && (best_r < 0 || r < best_r)) {
best_r = r;
best_k_sol = k_sol;
best_b_sol = b_sol;