242 Commits

Include renames

Author	SHA1	Message	Date
leonarski_fandClaude Opus 5	b22e1b6822	rugnux: raise the ice-ring mask margin to the measured null ... The mask drops a hexagonal-ice ring when its merged half-set CC1/2 falls a fixed 0.05 below its resolution shoulders. That margin is not a significance level: at the populations these rings actually have, 0.05 spans 1.1 to 7.3 sigma across firings, and a nominal Fisher-z error understates the real scatter of these heavy-tailed intensities by ~2.7x, so the null has to be measured rather than derived. Measured it with decoy bands - the identical ring/shoulder statistic evaluated at q positions carrying no ice ring - over the 37-crystal rotation battery: the gap's empirical null is p95 +0.032, p99 +0.095. So 0.05 sits near the 96th percentile, about 4% of ice-free bands clear it, and roughly half the 22 observed firings are indistinguishable from bands with no ice in them. The firing gaps are continuous, not bimodal, with 12 of 22 in [0.05, 0.10). Raise it to 0.10, the 99th percentile of that null. Firings 22 -> 10, crystals 12 -> 5, decoy false-positive rate 3.4% -> 0.8%. An independent check against XDS - which integrates through ice rings and so measures exactly what we delete - agrees: of the firings with a usable comparison, 9 true / 9 false becomes 7 true / 2 false. Battery: space groups 34 OK / 3 DIFF, the same three crystals as baseline, and no other discrete decision changes on 37/37. The heavily iced crystal keeps all five of its rings and its CC1/2 of 96.6; eight others recover 3.9-11.9% more unique reflections and up to 10.4 completeness points. Cost is CC1/2 -0.84 on one crystal, -0.35 on another, and agreement with XDS on the common reflections worse by a median 0.0004. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-05 19:00:10 +02:00
leonarski_fandClaude Opus 5	09fb8e0306	Bragg integration: clip the background ring high side instead of trimming it ... The r2..r3 background ring was averaged with a 10% SYMMETRIC trimmed mean. A symmetric trim is not a consistent estimator of the mean of a right-skewed (Poisson) sample: on a clean Poisson ring it sits ~0.1 ct/px BELOW the true mean at every level, and with ~50 signal pixels in the r1 disk that under-subtraction adds ~5 counts to every partial on every frame. Measured two independent ways on four rotation datasets - stored background_mean against a plain ring mean over the same pixels on reflection-free frames, and directly on apertures that provably hold no reflection. Empty-aperture pedestal, counts: plain mean -0.03..-0.20, 10% symmetric trim +5.05..+6.34, 4 sigma clip +0.02..+0.54. Replace it with a high-side-only sigma clip at mean + n*sqrt(mean), n = 4 for monochromatic data. It rejects the same one-sided contamination the trim was there for - better, in fact: a 40 px neighbour core at +100 ct shifts the trim by +10.1 ct/px, because a symmetric trim collapses once contamination exceeds ~10% of the ring, versus +0.009 ct/px at 4 sigma. False rejection on a clean ring is 0.04-0.39%. Broadband data keep their tuned 3 sigma clip unchanged. The trim stays reachable with --background-trim for back compatibility; setting either estimator clears the other, so they can never stack. --integrator boxsum does not take the clip (matching what the shipped clip already did), so it now uses the plain ring mean unless --background-trim is given. The intensities get measurably more accurate: per-shell agreement with an independent processing of the same images improves on 14 of 16 crystals (weighted -0.0347, outermost shell 12/4), the outermost-shell R_meas NUMERATOR - absolute scatter, not a denominator effect - falls 13.5% median on 16/5, and CC1/2 in the outer shell improves on 14/7. EXPECT <I/sigma> TO FALL AND EDGE R_meas TO RISE. Both are inflated by information-free counts, so both get worse when the bias is removed; neither is evidence against this change. That fingerprint is exactly how the trimmed mean was accepted in the first place. Known cost: over the 37-crystal rotation battery the de-novo space-group count goes 34 OK / 3 DIFF to 33 / 4. The single regression is a two-lattice crystal whose merge fails the absolute-sanity gate under either background (R_meas 63.5%, CC1/2 72.2%) and which carries an unresolved indexing ambiguity on the very operator being scored, so its operator CC is diluted by construction. No other crystal changes space group, and twin protection is not weakened - the H-ratio veto that refuses genuinely twinned crystals gets MORE decisive (1.63 -> 1.84, 2.83 -> 3.99). Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-05 18:03:11 +02:00
leonarski_fandClaude Opus 5	8ecd126e93	rugnux: stop the space-group search starving on a low-ISa merge ... The correlation stage kept only reflections with I/sigma >= present_i_over_sigma (3.0). That statistic is taken on the P1-MERGED intensities, whose sigma is floored at b|I| (Merge.h, SigmaWithSystematicFloor) so that ISa = 1/b is the asymptotic I/sigma ceiling - no reflection in a merge can read above it. Verified over the rotation battery: max I/sigma equals 1/b on every merge. So a fixed cut is not a per-reflection test at all. Every reflection sitting at the floor reads 1/b exactly, however strong, and on a merge whose ISa falls below the cut NOTHING passes: every operator is left with no pairs, its CC is NaN, and the point group collapses to 1. The predicate "search-merge ISa < 3" identifies the affected crystals exactly. It is latent today - no crystal in the battery starves on the shipped integration background - but it fires on four as soon as an additive intensity bias is removed, and it is not a data-quality verdict: the crystals it silences have final merges at ISa 19-22 while their low-multiplicity search merge sits at 3.5-4.0, just above the cut. Cap the cut at the merge's own I/sigma quantile so the correlation stage always keeps at least its strongest quarter. A no-op wherever the fixed cut already keeps that many - the cut stays exactly 3.000 on healthy merges. Battery unchanged at 34 space groups matching XDS / 3 differing, with merged observations identical to 0.000% on all 37 crystals. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-05 16:48:40 +02:00
leonarski_fandClaude Opus 5	fb55645b81	Revert "rugnux: fit the profile radius from the strongest spots too" ... Reverts the profile-radius part of 457b1bfd1; the comparison-script and mosaicity-column changes from that commit are kept. The cap was validated on the rotation battery, which cannot test it: the profile radius feeds `ewald_dist_cutoff` in IndexAndRefine, and that is read only by the STILLS predictors (BraggPrediction/BraggPredictionGPU). The rotation predictors gate on the mosaicity window instead and never look at it. So "no space-group changes, 36 of 37 crystals bit-identical" showed the quantity is inert for rotation, not that capping it is safe - and the one regime where it does act was never exercised. Validating it needs the serial-stills battery, which is a much larger exercise. Until then the arbitrary constant is not worth carrying in a code path nobody measured. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-04 21:35:17 +02:00
leonarski_fandClaude Opus 5	017f64690c	rugnux: smooth the per-frame geometry before scaling ... Geometry is re-refined independently on every frame, against that frame's spots alone - as few as a dozen on a sparse crystal, where XDS fits its equivalent to about sixty times more data. Measured over ten datasets the per-frame orientation carries two components: a slow drift that is real, with rugnux and XDS agreeing to R^2 0.83-0.88 on the two crystals that genuinely slip by 1.5 and 0.54 degrees, and a fast jitter that is fit noise, scaling with spots-per-frame at exponent -0.79 where counting noise alone would give -0.5. The jitter is worth 1-8% on merged intensities, 24% on the sparsest crystal. It cannot be fixed by refining less. Turning per-image refinement off entirely loses six space groups and a whole crystal, and even a 624-spot-per-frame crystal collapses; dropping the beam-centre terms holds the space groups but is worse on 31 of 37 crystals. The freedom is earning its keep, so keep it and suppress only the band that cannot be physical - a crystal does not re-orient and snap back from one frame to the next. So smooth the orientation in frame order after integration and recompute each partial's delta_phi, and hence its partiality, from the smoothed lattice. Batching at integration time was not an option: frames are processed independently and the online path depends on that. This runs before the GPU upload, so the device path picks it up with no separate kernel. The window is chosen per dataset by leave-one-out cross-validation, because the two components vary far too much for one number - drift spans 0.018 to 1.288 degrees and jitter 0.005 to 0.221, so any fixed window over-smooths one crystal while under-smoothing another. Chosen windows range from +-1 to +-20 frames. It is capped: cross-validation scores how well neighbours predict a frame's orientation, which on a barely-drifting crystal keeps improving with width, but the per-frame fit is also absorbing a real per-frame systematic and smoothing too wide destroys it - uncapped, one crystal chose +-60 and lost 16% of its ISa. Battery over 37 crystals: space groups unchanged at 34 matching XDS, R_meas better on 31 and worse on 6, low-resolution R_meas 30/7, ISa 26/10, high-resolution CC1/2 23/12. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-04 20:06:41 +02:00
leonarski_fandClaude Opus 5	457b1bfd1d	rugnux: fit the profile radius from the strongest spots too ... Same defect as the mosaicity in 2c94f3013, in the same file's sibling fit. The profile radius is an RMS of the excitation error over whatever spots were kept, and weaker spots sit further off the Ewald sphere, so it grows with the depth of the list: measured over a 150 -> unlimited spot budget it rises ~60%, and on a clean dataset as much as on a hard one, so this is general rather than something one awkward crystal provoked. That made it a function of --max-spots, which is an indexing budget, rather than of the crystal. Its one consumer treats it as a membership gate (ewald_dist_cutoff is twice it) where reflections at the cutoff carry near-zero partiality and are removed downstream anyway, so the integrated data barely notices: with the mosaicity already pinned, the partial count moves 0.2% across a 34% change in the radius. It is also reported per image as a diagnostic, though, and a number that slides with an unrelated setting is misleading to anyone comparing two runs - and it would stop being benign the moment anything used it as a width rather than a gate. Battery over 37 crystals: no space group changes, 36 of 37 bit-identical, no failures, one crystal marginally better. Also in the comparison script: report XDS's mosaicity next to rugnux's. XDS has two and they are not interchangeable - CORRECT.LP's REFLECTING_RANGE_E.S.D. is post-refined, while INTEGRATE.LP's per-batch SIGMAR is its integration-stage estimate, and the two differ by up to 2.3x. The XDS cell now prints both as postrefined|MLE so a per-image estimate is compared against the one measured the same way. Fixes a scoping bug in the same addition where every crystal read the last directory's INTEGRATE.LP. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-04 10:52:30 +02:00
leonarski_fandClaude Opus 5	2c94f3013e	rugnux: fit the mosaicity from the strongest spots only ... The per-image mosaicity MLE ran over the whole indexed spot list, so it rode on --max-spots, which is an indexing budget. A spot is detected when I_full * R(tau) clears the finder threshold, so selecting by intensity censors on R(tau): a deeper list holds proportionally more large-|tau| partially recorded spots and the fit widens with it. Raising the budget 250 -> 1000 widened sigma_M 0.059 -> 0.075 deg on a rotation dataset whose measured rocking width says 0.054. That is not cosmetic. An over-wide mosaicity mis-states every partiality in scaling: forcing the mosaicity across that range moved the merge error model from b 0.039 / ISa 26 to b 0.167 / ISa 6, and the space-group search lost a genuine 422 with it, merging the crystal in 222 instead. Cap the fit at the strongest 250 spots. FilterSpotsByCount leaves the list strongest-first, so this selects exactly the spots a smaller --max-spots would, and the mosaicity becomes invariant: 0.0538 deg at 250, 500, 1000 and 2000 spots, with the correct space group at each. Trimming or down-weighting the tau tail does not work - the censoring is multiplicative in R(tau), so it widens the whole distribution rather than adding a tail. Battery over 37 crystals: exactly one change, the demoted crystal repaired (33 space groups matching XDS -> 34). 23 of 37 are bit-identical, never reaching 250 spots. Unaffected elsewhere: the default spot count is 250, and stills have no goniometer so they return before the fit. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-04 07:34:40 +02:00
leonarski_fandClaude Opus 5	d79b20e268	indexing: key the shared device tables on their content, not only on an address ... The cache returned a device copy for a (device, host address) pair and cast it to whatever the caller asked for, with nothing checking that the bytes behind that address were still the same bytes. A host buffer can be mutated in place - PixelMask::LoadMask does exactly that - or freed and reallocated at the same address, and either hands the caller a device copy of something else. Nothing would report it: the tables are read-only geometry, so the engine would simply mask the wrong pixels for the rest of the run while the azimuthal mapping, the written pixel_mask dataset and the viewer overlay used the new one. Today that is unreachable, but only because of two guards in unrelated files that neither state nor assert the requirement. The byte length and an FNV-1a checksum of the bytes being uploaded are now part of the key. Both are computed once per engine construction, over a buffer that is about to be copied to the device anyway, so the cost does not show. Expired entries are pruned on insert, since distinct content now means distinct entries. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 16:33:24 +02:00
leonarski_fandClaude Opus 5	5727cb68a4	rotation_indexer: write down why the supercell bar is unreachable, and what failed to fix it ... `frac > RATIO * best_frac` cannot be satisfied once best_frac passes 1/RATIO - above 0.667 for a ratio of 1.5, which is ordinary for good rotation data. Above that the two guards do not raise the bar, they close the branch: no axis multiple and no lower-symmetry setting can displace the incumbent however much better it fits, so a genuine superstructure is kept as its sub-cell and its satellite rows go unindexed, silently. The obvious repair - restate the bar on the fraction left UNINDEXED, which is well defined over the whole range - was implemented and measured. It regressed the 37-crystal battery from 34/37 to 32/37 correct space groups: a C2 lattice fell to P1, and a P2 case went to C222 keeping 2923 of 22440 reflections with CC1/2 in the last shell at -35%. The indexed fraction is too noisy to carry a looser test. So the unreachable-but-safe form stays, and the limitation is recorded at the comparison rather than left to be rediscovered. Fixing it properly needs the selection to be decided on something better than the indexed fraction. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 15:32:55 +02:00
leonarski_fandClaude Opus 5	abb94ca450	spot_finding: accumulate the adaptive ring statistics in integers ... The per-ring sums were floats reduced by atomics, so the ring sigma - and with it the detection threshold - depended on the order the blocks happened to arrive in. Detection compares an INTEGER pixel value against that threshold, so a threshold that drifts across an integer flips every pixel of that value in the ring at once, which is how a last-bit difference turned into a different spot list. A preprocessed pixel is an exact int32 and the masked and saturated sentinels are skipped, so v and v*v are exact in 64 bits, and integer addition is associative: the sums no longer care about arrival order. Both engines now accumulate the same way, so they agree exactly rather than approximately, and the GPU spot list is bit-identical across runs. The corrected sums that feed the reported azimuthal profile stay float - a pixel value times a float correction has no exact integer form - but they do not enter the detection decision. Cost: the ring reduction needs 28 bytes per bin instead of 20 in the plain pass, which drops it from eight co-resident blocks per SM to seven and costs about 11% of that kernel (0.582 -> 0.650 ms/frame on a 4.5 Mpx frame). End to end it does not show: alternating runs on three rotation crystals came out the same or slightly faster, and the battery is unchanged in every number. The CPU engine got 30% faster (32.2 -> 22.6 ms/frame), integers being cheaper than doubles. Tests: exact CPU/GPU agreement on the spot list, and 50 repeats of bit-identical output where there were four. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 15:19:11 +02:00
leonarski_fandClaude Opus 5	212fbf9bab	rugnux: make the geometry-refinement sample deterministic, and spread it over the run ... The stills first pass drew frames from a shared cursor and stopped when a shared counter reached its target, which got two things wrong at once. The cursor walked the equally spaced sample in ascending order, so stopping early read only its leading PREFIX - the beam centre, distance and cell were fitted to the beginning of the run, not across it, and the comment claiming otherwise was wrong. And where the stop landed depended on how the workers happened to interleave, so the set of frames varied run to run: on the same data at -N 32 and -N 8 the pass examined 483 and 457 frames and refined the detector distance to 168.0481 and 168.0530 mm. The sample is now cut into a fixed number of interleaved stripes, each stopping once it has contributed its share. Every stripe spans the whole run, so an early stop no longer biases the fit, and a stripe is processed identically whichever worker claims it - so what gets examined depends only on the data, not on timing and not on -N. The same three runs now give 451 frames examined and 168.0452 mm, identically. The bundle selection was order-dependent too: frames are collected in worker-completion order and sorted by spot count with a non-stable sort, so equally strong frames swapped places between runs. They carry their image ordinal now and it breaks the tie. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 15:19:11 +02:00
leonarski_fandClaude Opus 5	2d3c39c9dd	image_preprocessing: write whole elements out of the un-transpose ... The raw-bytes path assembled each element a byte at a time, which on a full frame cost about 4x against writing the 8 contiguous elements a thread owns through an element-typed pointer. They are 8*ES-byte aligned, so the compiler merges them. 72.4 MB frame: 1.524 -> 0.406 ms for upload plus both kernels. The test now also times the LZ4 pass on its own, so the bounds and validity checks in the hot loop can be costed rather than guessed at. They are free: 0.231 ms against 0.2297 ms measured for the kernel before any of them existed - the restored offset == 1 and power-of-two fast paths pay for them. compute-sanitizer memcheck reports no error over 400 single-bit-corrupted payloads and nine malformed containers. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 14:16:58 +02:00
leonarski_fandClaude Opus 5	4b1c611bdf	image_analysis: query the current device, and upload the resolution mask on the engine stream ... Two leftovers from earlier fixes of the same shape. BraggIntegrationEngineGPU still read device 0's shared-memory size to decide whether its profile grid fits; workers are pinned round-robin across GPUs, so on a heterogeneous node that check can pass on a different card than the one the kernel launches on. SpotExtractorGPU still uploaded its default resolution mask with a pageable copy on the NULL stream, which is not ordered against the engine stream now that streams are created non-blocking. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 14:14:24 +02:00
leonarski_fandClaude Opus 5	bec7e2e922	image_preprocessing: fuse the bitshuffle inverse with preprocessing, and verify the decode ... The device decoder was byte-exact on every valid input - 994 production-compressed images, 927 hand-built LZ4 blocks covering engineered (offset, matchlen) pairs across the overlap branch boundary, 18000 repeat decodes, sanitizer-clean - and an audit against LZ4_decompress_generic could not construct a valid block it mis-decodes. What it did not do was notice when the input was NOT valid, and that mattered more than it looks: the decode buffers are reused frame to frame, so a block that stopped early left the PREVIOUS image in place, and in the bitshuffled layout the untouched tail is the most significant byte-plane. A corrupt chunk therefore did not look like a missing corner. It looked like thousands of real pixels several powers of two too bright, fed to spot finding with no diagnostic, where the host decoder had raised an error. So the kernel now flags a block that fails to reach its declared length while consuming exactly its payload, and the host turns that into an exception once the caller has synchronised. Reads are clamped against the end of the payload as well as the output, both length chains are bounded exactly as read_variable_length bounds them, the two offset bytes are bounded, and LZ4's parsing restrictions are enforced. On the host side a block size that is not a multiple of 8 elements is rejected (it made the un-transpose read uninitialised shared memory), the block count is bounded by what the chunk could hold before it becomes an allocation (twelve header bytes could demand hundreds of MB of pinned memory, permanently, per worker), trailing bytes are rejected, and the stream is synchronised before any throw that happens after work is queued. An image of fewer than 8 elements is all verbatim tail and now decodes rather than throwing. When the device route fails for any reason the host decoder gets its turn, so it costs speed rather than the acquisition. The lanes cooperate on the copies and a later match can read bytes another lane wrote, which since Volta needs an explicit __syncwarp(); it worked only because ptxas happened to reconverge at the post-dominator. The prototype's offset == 1 and power-of-two fast paths are also restored - the shipped kernel ran a runtime modulo, an emulated 32-bit division per output byte, on the path its own comment calls the common case. The un-transpose is now fused with preprocessing. One thread owns one group of 8 elements across every byte-plane, so once it has transposed its 8 bytes out of each plane it holds 8 complete elements and emits 8 finished int32 pixels with the mask, the error marker, the saturation cap and the statistics applied. The decompressed image is never materialised: 0.623 -> 0.411 ms/frame at 18 Mpx, 0.523 -> 0.340 with 8 concurrent workers. Staging nothing in shared memory also drops the 48 kB ceiling, which had made any file whose bitshuffle blocks exceed it a hard failure; 64 kB blocks now decode. gpu_compressed is sized from the chunk with grow-on-demand instead of from the uncompressed size - it was reserving ~73 MB per worker to hold ~4 MB. Measured on a 1630x1553 uint32 rotation set at -N 32, peak GPU memory falls 3756 -> 3084 MiB; the same model gives ~144 MB per worker on an 18 Mpx frame. Decoding on the device also stopped reporting a decompression time, which blanked the broker's compression plot trace and filled /entry/profiling/compressionTime with NaN. The decoder brackets the decode with CUDA events and reports it again. Tests: a differential fuzz suite against the CPU decoder - incompressible and highly compressible data, engineered offsets, a size sweep hitting every rem%8 value twice, all six element sizes, an 18 Mpx frame, decoder reuse, concurrency, hand-built LZ4 blocks across the overlap boundary, 26 foreign bitshuffle block sizes from 128 B to 64 kB, corrupt payloads and malformed containers, with a coverage report that proves which LZ4 paths were reached rather than assuming it. Plus the fused path held byte for byte against ImagePreprocessorCPU, statistics included, and against the host-upload path on the same frame. Battery: 37 crystals, every merged number identical to the host-decode run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 14:13:55 +02:00
leonarski_fandClaude Opus 5	7e47afe47f	rugnux: parallelise candidate-cell refinement, and stop repeating work in the tail ... Three independent changes to the CPU-bound parts of an offline rotation run, none of which alters a result. Candidate-cell refinement now splits across threads. RefineCandidateCells already took a (block, nblocks) partition, but the only call site passed nblocks=1, so the whole first pass of a two-pass rotation run sat on one thread per scheme - two threads, unchanged at every -N, for a third of the run. A block touches only its own scores(j) and cells rows and holds its own scratch, so the split is exact. The budget is a new IndexingSettings::RefineThreads, left at 1 by default and set only where few indexer threads exist: raising it unconditionally would oversubscribe the paths that already run one indexer per image across all workers. The mmCIF writer built a std::ostringstream per formatted number, twelve per reflection. snprintf gives the same digits for 0.535 -> 0.220 s per file. The space-group search built the same orbit mapping twice per candidate point group - once for the merge chi^2 and once for the systematic-error b, an apply_to_hkl and Canonicalize per observation per operator each time. Build it once and hand it to both. 18 Mpx rotation set 24.6 -> 18.7 s, 2.5 Mpx 13.0 -> 10.7 s, and the 37-crystal battery 13m55s -> 10m47s with no failures, the same 34/37 space groups, and statistics unchanged on 30 of 37 (the rest drift within the run-to-run spread the binary already had, which a control build with the split disabled reproduces). Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 07:30:33 +02:00
leonarski_fandClaude Opus 5	13aa20a528	bragg_integration: grow the GPU reflection arrays with slack ... EnsureCapacity resized its 13 device arrays to exactly the current image's predicted-reflection count, so every image that set a new record freed and reallocated all of them. cudaMalloc and cudaFree take a device-wide lock in the CUDA driver, so those images stalled every other worker: sampling the worker threads during the per-image loop found 21-24 of 32 parked in cuMemAlloc_v2 or cuMemFree_v2, all called from this one function, and the running maximum makes 32 workers do far more allocator work than one does. Grow by half again instead. All transfers and kernel launches are sized by the per-image reflection count rather than by the capacity, and the member is already documented as holding at least that many, so over-allocating changes no result. On an 18 Mpx rotation set the integration stage drops from 1.37 to 1.25 ms per image at 32 workers; merged statistics, error model and adopted space group are unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 01:02:07 +02:00
leonarski_fandClaude Opus 5	6e4c0ce202	image_preprocessing: decode bitshuffle+LZ4 on the GPU ... The pipeline decompressed each image on the host and uploaded the result. On an 18 Mpx rotation dataset that made the host-to-device copy the bottleneck of the whole per-image loop: nsys puts the copies at 78% of the loop against 39% for every kernel combined - 3600 transfers of 72.4 MB - and they ran at only 12.5 GB/s of an available 27-28 because the host-side decompression was itself saturating host memory bandwidth. The GPU was mostly waiting. So the compressed chunk goes across instead, about 4 MB rather than 72 MB, and is decoded on the device. That removes the transfer and the host decompression that was throttling it, in one change. Measured on an idle machine, a run goes from 45.11 s to 24.97 s - 1.81x - with the merged output unchanged. THE APPROACH IS JON WRIGHT'S (ESRF): "Experiences with GPU decompression for bitshuffle + LZ4 data", HDF5 User Group 2021, and github.com/jonwright/ bslz4decoders. The kernels here are ours, but the idea and the demonstration that it is worth doing are his. Cited in docs/ACKNOWLEDGEMENT.md and in the new section 0 of docs/CPU_DATA_ANALYSIS.md. Two kernels mirror the CPU decoder. LZ4 runs one WARP per bitshuffle block: every lane parses the same sequence stream (a broadcast read, no divergence) and the literal and match copies are split across the 32 lanes so the stores coalesce; an overlapping match is treated as a pattern of period offset sourced from bytes that already precede the write position, which keeps it parallel rather than a serial byte loop. One thread per block instead measured 13x slower. The bitshuffle inverse then un-transposes each byte-plane through shared memory and interleaves the planes back into elements. Only BSHUF_LZ4 is decoded on the device. The zstd variants have no device decoder, and neither has an uncompressed or float image; Supports() returns false for those and the caller decompresses on the host exactly as before. The fallback is explicit, so a format we cannot decode on the device is a slower path and never a wrong answer. Tests hold the device decoder against the CPU one byte for byte, on data from the production compressor, for every element size the detectors emit - including the 8-bit DECTRIS modes, which take bitshuf_decode_block's separate elem_size == 1 branch - plus a many-block frame, the formats it must decline, and malformed containers, which must throw rather than run off a buffer. Battery: 37 crystals, no failures, identical to the host-decode run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 00:16:36 +02:00
leonarski_fandClaude Opus 5	59702b0123	spot_finding: give the ring reduction eight blocks per SM instead of four ... reduce_rings_shared is the largest kernel in the per-image loop - 73% of GPU kernel time on an 18 Mpx rotation run, launched three times per image - and it is bound by shared-memory atomic replay rather than by bandwidth: it reaches 156 GB/s against a measured 913 GB/s ceiling, and removing the atomics while keeping the same loads makes it five times faster. That is the case that wants resident warps to hide the serialisation, and four blocks per SM left only 512 of the 1536 threads an SM can hold. The per-block histogram is nbins * 20 B, about 9.6 kB at the default 0.01 1/A spacing, so eight blocks fit in shared memory with room to spare. Both kernels are grid-stride loops, so any grid is correct and a device that cannot co-schedule eight simply queues the rest. Measured: 9.21 s -> 5.33 s of kernel time over a run (852 -> 493 us per launch), cutting total kernel time from 12.57 s to about 8.85 s. flag_strong keeps four. It is bandwidth-shaped rather than atomic-bound and eight measured no better (181 vs 175 us). Wall clock is unchanged, and that is expected rather than disappointing: kernels are 39% of the image loop while the host-to-device copy is 78%, so faster kernels idle the GPU more without shortening the loop. This is groundwork for the transfer work, not a speedup on its own. The shared accumulators are float and summed with atomics, so the block count changes the summation order and with it the last bits. The 37-crystal battery is identical crystal for crystal except one observation in 925850 on a single dataset. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-03 00:02:34 +02:00
leonarski_fandClaude Opus 5	a047275760	spot_finding: fix the GPU finder's main loop, which the tests could not reach ... Two bugs in analyze_pixel, both confined to the middle stage of the wave. The kernel walks each wave's rows in three stages. The priming and drain loops read prev_out and substitute INT32_MAX for a pixel the previous pass found strong, exactly as the CPU finder's value_at() does on every read. The main loop did not - it read the image raw. So in the second pass the pixels the first pass found strong stayed in the background statistics, inflating the local mean and variance, and the halo of every broad spot failed the signal-to-noise test. The two engines therefore did not agree, despite 1a1e05ad1 having set out to make them. Separately, shared_sum2 is an int64 accumulator but val*val and old*old were computed in int32 at three sites. That wraps above 46340 while the detector overloads around 1e6, so any window containing a bright pixel got a corrupted variance. pixel_result already did the same arithmetic in 64 bits. The reason this survived is worth recording: numberOfWaves is fixed at 32, so a wave owns ceil(height/32) rows, and the main loop only runs while front < rmax with front starting NBX+1 = 16 rows ahead. At the existing test's 100 rows a wave owns 4 rows and THE MAIN LOOP NEVER EXECUTES - every row goes through priming or drain, and the parity test passed on the broken kernel. On a 4362-row detector frame a wave owns 137 rows and the main loop carries about 121 of them, so the bug covered roughly 88% of a real image. The new tall-image test is sized to the partitioning rather than to convenience: 1024 rows, spots placed inside the main-loop region, one core at 100000 to exercise the overflow. Against the unfixed kernel it reports GPU 25 pixels where the CPU finds 49 and fails six assertions. The default rugnux path is unaffected because it uses the adaptive finder, and the 37-crystal battery is identical crystal for crystal. The classic finder is what SpotFindingSettings defaults to, so this is the online receiver's path; measured there with --no-adaptive-spots, R-meas improves 11.2% -> 10.9% and CC1/2 98.8% -> 98.9% on one crystal, with multiplicity up on both tried. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 23:31:44 +02:00
leonarski_fandClaude Opus 5	83e95b0c5a	indexing: stop computing angles the candidate filter only compares ... Candidate cell filtering called acos three times per candidate to turn dot products into degrees, then compared those against the min/max angle bounds. acos is strictly decreasing on [-1, 1], so "angle outside [min, max]" is exactly "cosine outside [cos(max), cos(min)]" with the ends swapped - the bounds convert once, and the three acos calls per candidate disappear. The same loop also re-derived every already-accepted candidate's unit cell on each new triple, inside the duplicate scan: three more acos each, for every candidate accepted so far. Those cells are now kept alongside the candidates. Measured on de-novo serial stills, where the indexer runs once per image: 34.43 s -> 14.17 s on one dataset and 21.92 s -> 6.59 s on another, with the indexing rate and the merged reflection count unchanged (one gained 0.25 points of indexing rate). acos had been 40% of the whole process there. Scope is narrower than that number suggests, and worth stating: the win is on the de-novo path, which Auto selects for stills only when NO cell is known. With a known cell Auto picks ffbidx, which reaches the same filter but feeds it few candidates - measured neutral there (+0.5% instructions, -1.6% wall, identical output), and that path already runs 14x faster in absolute terms. Rotation runs the indexer twice per dataset rather than per image, so it is unaffected: the full 37-crystal battery is identical, crystal for crystal. Comparing cosines instead of angles can only move a candidate that sits on the bound, so the filter's behaviour is unchanged except at that measure-zero boundary. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 21:24:57 +02:00
leonarski_fandClaude Opus 5	ea667cb306	rugnux: handle ice rings in --scale as the full pipeline does ... --scale did none of the ice handling the run that wrote the _process.h5 had done, so re-scaling a stored dataset silently produced a different - and flatteringly more complete - answer than the pipeline it was meant to reproduce. Three separate gaps: * --detect-ice-rings was accepted and ignored. The --scale block returns before the line that applies it. * Reflections were never flagged as sitting on an ice ring, so the per-image scale fit included them. The flag is not stored per reflection, so it has to be recomputed from the resolution. * RotationScaleMerge was constructed with the ice half-width hardcoded to zero. That is what turns a resolution into a ring index, so every ice test inside the merge was a no-op whatever was passed to it. The CC1/2 ring test that decides which rings to drop moves into FindDecorrelatedIceRings, shared with the full pipeline so both reach the same verdict on the same data, and --scale now re-merges with the mask the way the pipeline does. The stills branch re-runs only the merge: the scaling has already been applied to the reflections and repeating it would compound it. Measured on a rotation dataset with three decorrelated rings, --scale went from 8765 unique / 36.3% completeness / R-meas 18.5% / <I/sig> 1.1 to 7638 / 31.6% / 18.0% / 1.3, against the full pipeline's 7692 / 31.8% / 17.9% / 1.3 - the reported completeness had been inflated by reflections the pipeline drops. The full pipeline is bit-identical across the refactor. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 19:49:55 +02:00
leonarski_fandClaude Opus 5	b6d3dcc6fe	rotation_indexer: demand a decisive margin before adopting an axis multiple ... Candidate selection promoted a later cell whenever it indexed 0.05 more of the accumulated spots. That margin is not meaningful when the candidate is a near-integer volume multiple of the incumbent: multiplying an axis halves the reciprocal spacing, so the multiple has a lattice point wherever its sub-cell has one and another in between, and it collects spots the sub-cell leaves unindexed for reasons that have nothing to do with the crystal. The indexed fraction is biased in its favour, and a small lead is not evidence. On one rotation dataset the true cell and a spurious 5x supercell were separated by 0.003 of indexed fraction against a bar of 0.05 - close enough that the -march flags the binary happened to be built with decided it. The baseline build kept the true cell and merged to an R-free of 0.24 against an external model; an -march=x86-64-v3 build (what CI uses) took the supercell, carried it into a doubled cell and a different space group, and merged to an R-free of 0.58, which is noise. Both were reproducible, five runs each, and independent of thread count. An integer multiple now has to index 1.5x the incumbent, the same shape the lower-symmetry-setting guard next to it already uses. A real superstructure's satellite rows are a large share of its spots and clear that comfortably. Both builds now agree on the true cell with a wide margin. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 19:03:47 +02:00
leonarski_fandClaude Opus 5	8f1b0b2281	spot_finding: accumulate spot centroids in integers ... The photon-weighted position sums were floats, so the centroid's last bit depended on the build rather than on the data: gcc contracts the multiply-add in AddPixel into an FMA under -march=x86-64-v3 and cannot at the baseline, and MSVC does not contract at all under /fp:precise. The GPU extractor had to match with __fmaf_rn, and the parity test still needed a two-ulp slack for hosts that do not fuse. Column, line and the per-pixel count are all integral, so the sums are exact in int64 and both implementations reach the same bits with nothing to match. The parity test now demands exact equality unconditionally and gets it, including on a baseline build. ConvertToImageCoordinates keeps the sums integral too: the raw -> image map is a signed axis swap plus an integer translation, so it is applied to the sums instead of to the centroid. Drops the SpotToSave constructor, which had no callers and could not have been converted without quantising the stored centroid. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 19:03:35 +02:00
leonarski_fandClaude Opus 5	e7be5447d3	receiver: stop copying every frame back from the device on the Lite path ... The Lite workflow built its analysis with the fused GPU engine disabled, which is also what decides whether the preprocessed image is copied device-to-host after every frame. So on a machine with a GPU the online path was moving the whole image back - 72 MB on a large detector, every frame, per worker - for a host reader that does not exist on that path. It was left off deliberately when the fused engine was added, to keep the online path unchanged in that commit, and never revisited. Nothing depends on it: the FPGA workflow uses a different analysis class, and strong-pixel values are read through a device gather rather than from the host image. Turning it on changes no result, and cannot: adaptive detection is unreachable online, because the REST schema exposes no way to enable it, so the classic GPU finder runs either way. Measured anyway, both engines on the same frames across five datasets including very weak ones: 2400 frames, 638260 spots, not one difference - identical lists, identical indexing rate, identical merge statistics to every printed digit. On a large detector with eight workers the median per-image cost falls from 94 to 59 ms and preprocessing from 21 to 6 ms; throughput rises from about 48 to 55 Hz. No percentile regresses, which is what matters for a service - the ninetieth improves from 128 to 74 ms and the tail with it. Spot finding gets faster too, because the large copy no longer contends with the device gather. Correct two statements while here. The flag's comment and the data-analysis document both said the online receiver uses the CPU adaptive finder; online never runs an adaptive finder at all, and the copy the flag really controls was not mentioned. That copy would be better expressed as what it is - whether a host engine will read the image, which the constructor already knows - rather than inferred from which spot finder is wanted. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 15:16:18 +02:00
leonarski_fandClaude Opus 5	4bdb229fb8	spot_finding: find connected components on the GPU ... The spot finder flagged strong pixels on the device and then labelled them on the host, so every frame sent the packed bitmask back - 2.26 MB on a large detector - and the host walked all of it to recover a few hundred pixels. Do the labelling on the device instead: compact the bitmask into a flat-index-sorted list, find each pixel's backward neighbours by binary search, union them lock-free with path halving, then label, accumulate and filter in one kernel. Only the spot list comes back, and only one stream synchronisation per frame. The gain in the ordinary case is modest - about a quarter off per-image spot finding - because the host algorithm is genuinely fast on a normal frame. What justifies it is the frame that is not ordinary. The host labels a sorted sparse list through a window spanning two detector lines, so its cost is quadratic in how many strong pixels share a line. A lit band of detector rows - a hot module, a panel edge - costs 33 ms at two rows and 377 ms at fifteen, all of it under the pixel cap that was supposed to bound this, and none of it maskable when the cause is a diffraction ring rather than a defect: a ring runs tangent to a row at its top and bottom, which is exactly the shape that hurts. The device version is flat at 0.05 to 0.64 ms across every geometry tried, so an online run no longer stalls a quarter of a second on an ice ring. Rejecting an over-cap frame is now free too, since the count is known before any pixel is written. Also label once and filter three times. The per-image minimum-pixel search runs the extraction at three settings, but that setting only decides which components are kept - it does not change the components - so the search itself need not be repeated. This helps the host path as much as the device one. The resolution mask moves to the device as a bit mask, uploaded when the limits change rather than per frame, since the compaction needs it there. Parity is asserted permanently rather than argued: five cases covering realistic frames, occupancy from a hundred pixels to past the cap, the pathological geometries including rings, the resolution mask, and a hundred-repeat determinism check - requiring the same partition, the same spot order, and identical counts. The centroid is a float sum and therefore order-dependent, so the device walks each component from its root in ascending order and fuses its multiply-add the way the host's does; note that whether the host fuses at all depends on the architecture flags, so exact centroid equality is asserted where the compiler fuses and a two-ulp bound otherwise. Making those accumulators integer would remove that dependence entirely and is worth doing separately. Regression set: all 37 crystals identical to the last printed digit. Unit suite passes with the new cases. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 13:34:43 +02:00
leonarski_fandClaude Opus 5	e4d70f0e55	image_preprocessing: inline the buffer accessors ... operator[], size(), data() and getBuffer() are one-line accessors that were defined in the .cpp. The build sets no link-time optimisation, so out of line each of them is a real call - once per pixel, from the CPU preprocessor, the CPU azimuthal integrator and the CPU spot finder - and they stop those loops vectorising at all. They show up in a profile directly: about six per cent of a whole azimuthal-integration-only run is spent in the call overhead of two accessors that do nothing but index a vector. Moving them into the header retires 30% fewer instructions on that run and takes the per-image CPU cost on a GPU-less pass from 34.6 to 24.2 ms, with the output bit for bit unchanged - same observation count, same cell, same merge statistics. It is worth nothing on the GPU path, where the image stays on the device, and everything on the paths that have no GPU to fall back on. This also explains a measurement that had been blamed on the pixel mask being a vector<bool>: a microbenchmark of that loop indexed a raw pointer and came out far faster than the same loop in the binary, and the difference was this call, not the mask. Measured properly the mask costs about 14% single-threaded rather than the 41% claimed, and at the thread counts this actually runs at the bit mask is FASTER than the byte mask it was proposed to become, because it moves eight times less traffic and the loop is bandwidth bound. That change should not be made. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 11:20:31 +02:00
leonarski_fandClaude Opus 5	639fbb3fbc	indexing: select predicted reflections by partiality, build indexers where it pays ... When more reflections are predicted for a frame than the output can hold, the surplus was dropped by keeping those closest to the Ewald sphere. On the rotation path that quantity is identically zero by construction - the rocking coordinate is chosen so the scattering vector lands exactly on the sphere - so the comparison fell through to h, k and l and the survivors were whichever came first in lexicographic order. Measured on a large cell: every value within one float ulp of zero, and the kept set had a MEAN PARTIALITY BELOW that of the full set, i.e. worse than choosing at random. Rank by partiality instead, which the predictor already computes and which is what the header always claimed was being kept. On the one regression crystal large enough to cross the cap this lifts completeness from 84.8% to 90.2% on the same observations; multiplicity and R_meas move the way they must when the same measurements cover more of reciprocal space. The online path asked for a cap of ten thousand but the truncation was hardcoded to the offline limit, so the broker predicted and integrated up to six times what it could transport and discarded the rest after paying for it. Honour the caller's limit, which also makes the post-integration re-truncation dead code. Indexer pool construction becomes a policy. The online service needs every indexer resident before data arrives, because a cuFFT plan built on the first frame is planning time inside the measurement; spending memory to be ready is the intended trade there and stays the default. Offline there is no such deadline, and a stills run with a known cell was holding a fully allocated FFT indexer per worker that the algorithm resolution can never dispatch - 2.8 GB where 0.4 GB is needed. rugnux and the viewer opt into building on first use; the broker, the receiver and the tests are untouched. This also removes a dangling reference that was latent: the worker held the settings by reference although the pool is routinely constructed from a temporary, which only survived because eager construction finished inside the constructor call. Finally, refuse a first-pass lattice that indexes fewer than a sixth of the validation frames. It fires on nothing in the regression set - the weakest real crystal sits at 22 of 60, more than twice the floor - so it is a backstop, but the failure it prevents is one the set does contain: a dataset with no crystal at all adopts a lattice from its powder rings, integrates every image against it, and dies much later inside the merge complaining about resolution. It now stops in the first pass and says what to try. Regression set: 36 of 37 crystals byte-identical, the exception being the completeness gain above; 34 of 37 space groups, no failures. Full unit suite passes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-02 09:12:27 +02:00
leonarski_fandClaude Opus 5	6e805f53c0	image_analysis: stop paying for work that is thrown away ... Three independent costs, each measured, none changing a result. Across the 37-crystal regression set the run time halves (median per crystal 2.0x, total 2.3x) and every crystal's merge statistics are unchanged. The image copy back from the device moved the whole preprocessed frame - 72 MB on a large detector, every frame, per worker - to serve a single host consumer that reads only the strong pixels, at most a few hundred kilobytes of it. Give the buffer a Gather() so that consumer asks for the values it actually wants (a host loop on the CPU, a small kernel on the GPU), and copy the frame back only when a CPU spot finder will genuinely read it. The copy the other way was worse: it came from an unregistered vector, so the driver staged it through its own pinned pool with a host-side memcpy on the calling thread, which does not overlap and collapses under concurrency - 11.6 GB/s at one worker, 1.6 GB/s at eight. That, not any hardware limit, is why throughput stopped improving past four to eight workers. Pinning the decompression buffer once per worker fixes it: on a 18 Mpx dataset the image loop goes from 13.6 to 7.9 ms per image at 32 workers, and 32 workers now beat 8 instead of losing to them. Ceres was computing seventeen partial derivatives where five are free. The per-image rotation refinement frees the beam and the orientation and holds distance, detector angles, rotation axis and cell constant, but the cost function declared all seven blocks, so every residual evaluated in Jet<17> arithmetic. A residual exposing only the two free blocks - the same arithmetic, the constants baked in - halves refinement, and it is exact rather than merely close: dual coordinates evolve independently, so the residuals and the free Jacobian columns are unchanged bit for bit. The merge sorted an index array with a comparator that dereferenced a 1.6 GB array of 72-byte records, i.e. a random walk over memory, single-threaded, twice per two-pass run. Sorting a packed key instead is 2.4x. French-Wilson allocated its integration scratch per reflection and ran serially; it now takes caller-owned scratch and runs over chunks, 4.2x. The correction surfaces re-tested every observation for usability and parity on each of ~22 passes and re-allocated their accumulators each time; bucket the indices once and hoist the buffers. Also convert std::round to std::rint where the rounded value only ever enters a squared residual. The tie rules differ - away from zero against to even - so this is safe exactly where a tie flips the sign but not the magnitude, and unsafe wherever the value becomes a Miller index; those sites keep std::round. Verified over all 2^32 float bit patterns: 8388608 exact ties exist, and the squared residual is bitwise equal for every one of them. Worth little on its own here, because the rounding that dominates is in candidate refinement, where the value is an index and the substitution is not available. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-01 21:35:28 +02:00
leonarski_fandClaude Opus 5	0ae1a307bc	indexing: complete a rank-deficient direction set, and keep the higher-symmetry setting ... The FFT shortlist could be rank-deficient, and then no cell could be formed at all. FilterFFTResults takes the strongest max_vectors RAW directions and only then prunes ones within 5 degrees of each other, but a single lattice row is sampled by many neighbouring directions of the 16k half-sphere, so thirty raw peaks routinely prune down to four or five distinct directions - the strongest, hence shortest, rows. When a crystal's densest rows share a plane, every surviving direction is coplanar, every triple the reduction forms is degenerate, and the indexer returns nothing. On such a crystal the weak third axis was the eighth distinct direction, at raw rank 78. Keep walking the same magnitude order for up to four more directions that are 5 degrees clear of everything kept, appended after the length sort so the earlier entries hold their positions and the reduction still forms every triple it formed before - the shortlist only gains candidates at its end. That exposed two ways a change of SETTING was mistaken for a different lattice. A centred conventional cell is an exact integer multiple of its primitive one, so the same lattice described two ways differs by that factor: comparing conventional volumes reads a setting change as a sub-cell or a supercell. Both the candidate selection in the rotation indexer and the pass-2 comparison in the driver did exactly that, and between them they discarded a correctly-classified cubic F cell in favour of the body-centred tetragonal description of the very same lattice. Compare primitive volumes in both, as the scheme comparison already did. Fixing the volumes alone was not enough, because the indexed fraction is also biased across crystal systems: a subgroup setting holds fewer cell parameters fixed than its supergroup, so it can never index fewer spots and will always look better by that measure. Where a candidate has a lower lattice point-group order at the same primitive volume - the signature of the same lattice in less symmetry - require it to index markedly better, not merely better, before it displaces the incumbent. A general metric-symmetry promotion was implemented and rejected on evidence: it raised a correct body-centred orthorhombic cell to triclinic and a monoclinic one to C-centred orthorhombic, and no threshold separates the cases, because a false pseudo-orthorhombic degeneracy measured tighter than a true cubic one on obliquity and on alternative-basis axis excess alike. Metric alone cannot decide this; only the intensities can, which is what the space-group search is for. Measured over the 37-crystal regression set: one crystal goes from failing outright to 91% indexed with 91% completeness and a better R_meas than the reference, one keeps the cubic setting it had before, and every other crystal is byte-identical. Full unit suite passes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-01 12:59:34 +02:00
leonarski_fandClaude Opus 5	0b1fb6c870	image_analysis: share the read-only GPU lookup tables per device ... One analysis engine is built per worker thread, and each uploaded its own copy of tables that are pure functions of the detector geometry: the pixel -> azimuthal bin map and the per-pixel corrections (both in AzIntEngineGPU AND again in AdaptiveSpotFinderGPU, from the same mapping), plus the pixel mask. On an 18 Mpx detector that is ~224 MB per worker; with 32 workers ~7 GB of device memory held 32 identical copies. Upload each table once per GPU instead and hand every engine on that device a shared pointer to it. The cache is keyed by (device, source-vector address) because workers are pinned round-robin across GPUs, so on a multi-GPU node each device keeps its own copy - a kernel may only read memory resident on the device it runs on - and the table is freed on the device that allocated it. Entries are held weakly, so a table goes away with the last engine using it. Measured on an 18 Mpx detector, 32 worker threads, 16 GB card: the stills path went from exhausting the card (OOM in de-novo indexing) to 8.6 GB peak, and a normal rotation run from 14.6 GB to 7.4 GB - it had been running within 1.6 GB of the limit, so any larger detector or second GPU consumer would have tipped it over. Per-worker footprint drops 403 -> 173 MB. Merge statistics are unchanged on a six-crystal regression subset, including two-pass runs where the second pass rebuilds the mapping on refined geometry, and wall time is unchanged (13.5-13.8 s vs 13.8-14.1 s). Also take the launch configuration from the current device rather than device 0 in AzIntEngineGPU and ImagePreprocessorGPU: with round-robin pinning, device 0's SM count and shared-memory size can belong to a different card than the one the kernels use. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 18:41:27 +02:00
leonarski_fandClaude Opus 5	1a1e05ad14	spot_finding: run the same two passes on the CPU as on the GPU ... ImageSpotFinderGPU::Detect launches its kernel twice, feeding the first pass's strong-pixel bitmap back in so the second recomputes each local background with those pixels excluded and keeps them strong. The CPU finder ran a single pass, so the two returned different spot lists for the same frame and a dataset processed without a GPU did not match one processed with it. It matters for any spot wide enough to reach into its own 31x31 background box: the spot inflates the mean and variance it is then tested against, so its outer pixels fail the SNR test. On the test image added here - a 5x5 core at 300 counts with a one-pixel ring at 25 - a single pass returns the 25-pixel core and 7500 counts where two passes return the full 49 pixels and 8100. pxl_val also becomes int64_t, matching the GPU's pixel_result signature. It was int32_t, so pxl_val * pxl_val overflowed above 46341 counts even though the surrounding sums were already 64-bit. The new parity test compares PixelCount and Count, not just the centroid, which does not move for a symmetric spot whether or not the ring was picked up; it was confirmed to fail against the old single-pass CPU. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 15:37:22 +02:00
leonarski_fandClaude Opus 5	b5b7cf2cf9	docs: say what the CPU prediction path actually does ... BraggPrediction.h claimed the buffer "GROWS to whatever a frame actually predicts, so a large cell is never truncated here". Only the two GPU Calc overrides call GrowCapacity; both CPU predictors stop at max_reflections. The cap is applied inside the h/k/l walk and before the resolution test, so what survives is the low-|h| block, not the reflections nearest the Ewald sphere - a cell large enough to overflow 20000 gives different merged reflections with and without a GPU. Documented rather than silently claimed otherwise. Also removed a paragraph describing a once-per-predictor overflow warning that no longer exists, and fixed the rugnux_cli.cpp path in HDF5.md. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 15:02:57 +02:00
leonarski_fandClaude Opus 5	164f15c903	geom_refinement: stop committing refinements that did not converge ... Four of the seven ceres::Solve calls in image_analysis obtained a Solver::Summary and never looked at it, so a solve that failed numerically had its parameters written back and was reported as success. StillsPartialityRefine and both PostRefine solves already gated on IsSolutionUsable(); this brings the rest to the same contract. IsSolutionUsable() is the right test rather than checking for CONVERGENCE: it accepts a solve that ran out of iterations or wall-clock time but still descended, which is exactly what the real-time callers depend on when they set max_solver_time instead of max_num_iterations. Only FAILURE and USER_FAILURE are rejected. XtalOptimizer checks before the write-back, so a failed refinement now leaves the caller's geom and latt untouched instead of half-updated. GeometryRefiner folds it into result.ok, which previously reported success from spot and frame counts alone. RingOptimizer returns a geometry by value that both callers assign straight back over their input, so it hands back the unchanged reference rather than a diverged beam centre. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 15:02:11 +02:00
leonarski_fandClaude Opus 5	fb0272023e	scale_merge: apply outlier rejection to the anomalous split on the GPU path ... The GPU merge kernel rejects outliers on the device and keeps a per-full flag there, but only returned the per-group counts. The host array the CPU path fills stayed all zero, and the anomalous I(+)/I(-) accumulator is host-side and unconditional - so with --reject-outliers and a GPU present, the observations the merged IMEAN dropped were still averaged into I(+) and I(-). The same command on a CPU-only host excluded them: the exported anomalous differences depended on whether a GPU was there. R_meas was unaffected, having its own device-side path that reads the flags in place. MergeAccum now hands the per-full flags back so every host-side reduction sees the same rejections. The comment claiming reject_outliers was excluded from the GPU path was never true. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 14:53:17 +02:00
leonarski_fandClaude Opus 5	d2d1d78545	spot_finding: keep the GPU wave inside the image ... rowsPerWave is rounded up, so with 32 waves the last waves can start at or past the last row: rmin was never clamped and only the drain loop checked front against height. On any detector below about 1500 rows - including the module-converted 500K and 1M geometries and the kernel's own unit tests - the priming and steady-state loops read whole rows past the end of the image buffer, and those garbage rows entered the sliding background window of the bottom rows. Blocks with no rows to write now return before the first __syncthreads (rmin depends only on blockIdx.y, so the block leaves together and the collective ops stay well formed), and both remaining reads are bounded by height. Rows past the end keep the INT32_MIN sentinel, which the window already treats as "not counted". The raw read in the steady-state loop is left as it is: making it apply the prev_out substitution that the other two read sites use would change which pixels are found, which is a separate question from this fix. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 14:43:49 +02:00
leonarski_fandClaude Opus 5	4a8a8e69bf	lattice_search: give Niggli character 40 its own reindex matrix ... Character 40 carried a verbatim copy of character 35's matrix (0-10 / -100 / 00-1), whose determinant is 1. A C-centred conventional cell needs determinant 2, so a genuine oC lattice was returned as its primitive monoclinic cell while still being labelled Orthorhombic 'C': the refiner then clamped a ~117 degree beta to 90 and prediction dropped half the reflections of a cell that has no centring. International Tables A 3.1.3.1 gives 0-10 / 012 / -100 for character 40. Character 35 is correct as it stands and is left alone. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 14:41:46 +02:00
leonarski_fandClaude Opus 5	5e3d580a0e	lattice_search: fix the sign of the Niggli character 9 reindex matrix ... International Tables A 3.1.3.1 gives 100 / -110 / -1-13 for character 9; the last element was -3. With a negative determinant the transform is left-handed and the "conventional" rhombohedral cell is not hexagonal - beta came out around 110-134 degrees instead of 90 and c was far too long. Any R lattice tall enough to reduce to character 9 was affected, and the downstream Trigonal->Hexagonal promotion then forced 90/90/120 onto that wrong cell. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 14:41:25 +02:00
leonarski_fandClaude Opus 5	0ca159449f	Bragg integration: integrate as far as the detector reaches, not to a fixed 1.0 A ... BraggIntegrationSettings::DMinLimit_A had a setter that nothing anywhere called, so it was always its 1.0 A default - in rugnux, the viewer and the broker alike, with no option or API field to change it. It feeds the predictor as high_res_A, which discards any reflection with |q| > 1/d_min, so integration simply stopped at 1.0 A however far the detector reached. Five of the 33 rotation test datasets have detectors reaching past it, down to 0.981 A. On one of them, run with no resolution limit, the shell table ended dead at 1.00 A with that shell still at CC1/2 55.6% and <I/sig> 3.4 - cut mid-shell rather than fading out. This branch had already made the sibling limits detector-driven (spot finding, scaling), so the pipeline was finding spots the detector could see and then refusing to integrate them. Make it a std::optional: unset means as far as the detector reaches, a value limits. The limit is only a bound on how far the lattice walk goes, never a second opinion on what is measurable - both predictors independently drop reflections that miss the detector (BraggPrediction.cpp, BraggPredictionRot.cpp) - which is what makes the detector's own reach the right default. rugnux gains --integration-high-resolution (0 = no limit, as for --spot-high-resolution); the derived per-axis prediction range resolves against the same number, so the two cannot drift. Full battery: 30/33 space groups, unchanged from before, 0 failures and the same three known mismatches; 22 of 32 crystals bit-identical and nothing worse than 5 observations in ~500k. The datasets that gain do so because their detector reached past 1.0 A - the effect is understated here because the harness caps each merge at the XDS resolution anyway. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 14:03:34 +02:00
leonarski_fandClaude Opus 5	406c406988	Bragg prediction: one limit per index, not one cube ... Each Miller index is bounded by its OWN axis - |h| <= a/d_min, |k| <= b/d_min, |l| <= c/d_min - so a single half-width has to be sized for the longest axis and then walks the short ones far past anything the resolution cut can keep. Give the predictor max_h, max_k and max_l instead, in all four implementations (CPU and GPU, stills and rotation), and derive each from its own axis. On a 149/83/226 A cell that is 23.1M candidates per frame instead of 94.2M, 4.1x fewer. Results are bit-identical, as they must be - the candidates removed are only ones the |q| <= 1/d_min cut rejected anyway: over six rotation crystals every merged observation count, high-shell CC1/2 and space group matches the cube exactly, 6/6 space groups correct. It buys almost no time, and the earlier claim that the cube cost 22% of that crystal's wall clock was wrong. Removing 4.1x of the candidates moves it 1m58s -> 1m57s, so the whole prediction sweep is ~1% of the run. The 22% that crystal costs relative to a fixed max_hkl of 100 is genuine extra work at max_l = 227: real reflections inside the resolution sphere along the long axis, predicted and integrated either way. Per-axis limits do not reduce that and cannot. The user-facing setting stays a single number: it exists to bound the work, not to describe the crystal, and applies to all three indices when set. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 13:31:34 +02:00
leonarski_fandClaude Opus 5	b0e315e73c	Bragg prediction: derive the lattice walk from the cell, and expose it in the API ... Follow-up to making max_hkl a setting: it is now an optional, and unset means "take it from this crystal". The predictor keeps only |q| <= 1/d_min and h = a.q for the real-space axis a, so |h| <= a/d_min exactly - and likewise |k| <= b/d_min and |l| <= c/d_min. max(a,b,c)/d_min therefore bounds all three at once: nothing that could be predicted lies outside it, and nothing inside it is reached by a shorter axis. It applies to rotation and stills alike, both going through the one place the prediction settings are built. Offline (rugnux, viewer) the default is unset, so every crystal gets its own range; --max-hkl overrides it. Online the broker holds a concrete number, because the cost is the cube of it per image and a live acquisition should not have its frame rate decided by whichever sample is mounted: max_hkl joins bragg_integration_settings in the OpenAPI with a default of 100, so an omitted field arrives as that default (the generated model carries it) rather than as "derive it", and the frontend exposes it next to the integration model. Measured against a fixed 100 on six rotation crystals: three are bit-identical, two were being truncated and recover 419k and 5.8k observations with the high-shell CC1/2 going 15.1 -> 25.8% and 52.1 -> 55.3%, and the space group is unchanged 6/6. It reproduces a fixed 200 exactly, which is the bound being tight rather than merely safe. The sixth is worth recording: a 149/83/226 A cell derives 227, and because a single scalar has to cover the longest axis the cube is ~16x what a per-axis box would be - 22% wall clock, for a net 22 observations out of 364k (the per-frame 65536-reflection cap re-selects at the margin when more candidates are offered) and identical CC1/2, ISa and space group. Per-axis limits would remove that; the predictors already map a thread index to h, k and l separately. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 13:10:58 +02:00
leonarski_fandClaude Opus 5	a06c06931f	Bragg prediction: how far to walk the lattice is a setting, not a literal ... max_hkl was hardcoded to 100 at the one place production builds the prediction settings, so the only way to change it was to edit and rebuild - and it is not a constant of the method, it is a property of the cell. An axis is truncated once a/d_min exceeds it: 100 covers a 150 A axis at 1.5 A, but the same axis at 1.0 A, or a 250 A axis anywhere, loses its outermost reflections with nothing said. Move it into BraggIntegrationSettings next to the other prediction/integration parameters and add rugnux --max-hkl (1..511, default 100 - no behaviour change). Like the integration radii and the background trim it stays out of the OpenAPI, so the broker keeps the default it has today and live analysis cannot be handed a range that would not finish; the offline front end, which knows its cell, can ask for more. RugnuxCommandLine emits it when it is not the default. Measured on five rotation crystals at --max-hkl 200: two are bit-identical at no cost, and three were being truncated - one gains 419k observations (+17%) and takes its high-shell CC1/2 from 15.1% to 25.8% for +14% wall clock, the other two gain 12k and 5.8k observations with CC1/2 76.6->82.4% and 52.1->55.3% for +9% and +1%. ISa is unchanged throughout, and no frame overflowed the prediction buffer. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 12:00:45 +02:00
leonarski_fandClaude Opus 5	953302a9eb	Spot plot: the resolution axis comes from the detector, not from 1.5 A ... Dropping the fixed spot-finding limit left the reader still generating the spot-vs-resolution plot over shells that stop at 1.5 A, so a stored file reopened in the viewer showed a plot truncated at exactly the limit that was removed - GenerateSpotPlot drops every spot outside its shells. Pass the detector's own maximum resolution, as SpotAnalyze already does. That value is 0 when the geometry gives no scattering angle at all (no distance or no wavelength), and ResolutionShells throws on a non-positive d_min, once per image. There is no resolution axis to plot against in that case, so skip the plot. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 11:51:22 +02:00
leonarski_fandClaude Opus 5	bf866a0d4c	CUDA: the engines' setup copies belong on the engine's stream ... Making the worker streams non-blocking removed the implicit ordering that the constructors were still relying on. Each engine uploads its static inputs - the pixel mask, the pixel-to-bin map, the corrections, the ROI map - with a blocking NULL-stream cudaMemcpy, and then reads them from kernels on its own stream. A pageable host-to-device cudaMemcpy returns once the source has been staged, with the DMA still in flight, and a non-blocking stream no longer waits for the NULL stream. The failure mode is a silently unapplied mask or a stale mapping, not a crash, so it would not have announced itself. Put them on the stream the engine already owns, and synchronise once at the end of the constructor - that is required for the preprocessor, whose source is a local vector, and leaves the others settled rather than in flight for the cost of one one-time sync. The GPU spot-finder test uploaded its image the same way. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 11:51:22 +02:00
leonarski_fandClaude Opus 5	a2adc4e021	Stills scaling: an image whose scale collapsed is dropped, not merged unscaled ... Leaving it at G = 1 looked like the conservative choice and is the more damaging of the two errors. The per-image scale enters as rlp/(partiality*G) and multiplies intensity and sigma alike, so substituting 1 for a scale that was really 1/200 of the run median puts the intensities in 200x too low with sigmas 200x too low too - 1/G^2 times the weight they deserve. The merge cannot defend itself against that, because the number that is wrong is the number the weight is built from. And if the collapsed value was instead a failed fit, G = 1 merges the image mis-scaled by an unknown factor. Per-crystal scales on serial stills genuinely span orders of magnitude, unlike frames of one rotation sweep, so both readings are live. An image whose scale is not believable has no usable scale. Write NaN into its image_scale_corr, which every merge path already skips on, so it drops out of the merged intensities, the error model and the statistics consistently. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 11:51:08 +02:00
leonarski_fandClaude Opus 5	da74197dea	Stills partiality: an unmeasurable CC is not a reason to adopt the refined tilt ... The "keep what the crystal came in with" gate required std::isfinite(cc) before it would reject, so a refined model whose CC could not be measured at all was adopted. ImageReferenceCC returns NaN when fewer than 20 reflections clear the partiality cut - which is exactly what a refinement that collapsed the partialities produces, since the cut is on the partialities it just rewrote. The gate therefore failed open on precisely the crystals it exists to catch, and wrote the NaN into image_scale_cc, on which --min-image-cc then drops the image from the merge, the error model and the statistics. Treat a CC that cannot be measured as worse than one that can, so the crystal is put back exactly as it arrived. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 11:51:08 +02:00
leonarski_fandClaude Opus 5	196c72a7fe	Bragg prediction: the rotation GPU launch was one plane short in each direction ... The kernel guards against 2*max_hkl+1 and maps thread i to h = i - max_hkl, but the host launched a grid sized 2*max_hkl. The h = k = l = +max_hkl planes were therefore never launched while -max_hkl was, so the GPU predicted an asymmetric subset of what the CPU loop (inclusive on both ends) does. The same bug was fixed on the stills twin when the whole hkl range moved to the GPU; the rotation predictor kept the old expression. It only bites where the cell actually reaches |h| = 100 inside d_min - a ~150 A axis at 1.5 A - so most data never noticed. Over the 33-crystal rotation battery 29 crystals are bit-identical and 4 gain observations, all of them large-cell or high-resolution: +8519, +4693, +901 and +758 observations, with the high-shell CC1/2 up 15.0->15.1%, 52.0->52.2%, 76.3->76.6% and 51.6->52.1%. Nothing is lost anywhere, and R-meas and ISa move by at most 0.01. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-31 11:50:10 +02:00
leonarski_fandClaude Opus 5	2be8680422	CUDA: let worker streams run concurrently ... Every per-thread stream was created with cudaStreamDefault, and the 20 MB raw image upload went to the legacy NULL stream. A NULL-stream operation implicitly synchronises with every blocking stream in the process, so with one engine per worker thread no two workers' GPU work could ever overlap - the whole GPU pipeline ran serially however many threads were asked for. Create the streams non-blocking and put the upload on the engine's own stream. Measured on 2000 serial stills, interleaved, medians of three: 24.6 -> 19.2 s at -N 32 (-22%), 32.7 -> 21.0 s at -N 16 (-36%), CPU utilisation 436-570% -> 723-859%. Output bit-identical - same observations, uniques, completeness, R-meas, CC1/2, error model and cell. The stream is synchronised at the end of the same function, so the ordering the code relies on is unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-30 23:05:11 +02:00
leonarski_fandClaude Opus 5	07607d3d47	Keep the online reflection cap where the transport can carry it ... Raising the per-image reflection limit to 65536 for offline reprocessing also raised the image-buffer headroom derived from it, and that headroom divides a FIXED total buffer - so every slot grew from compressed+4 MB to compressed+16.7 MB and the receiver's slot count, i.e. how much of a burst it can absorb, fell by about three. Online never needed the raised limit: measured on three serial stills datasets the worst frame predicts 1380 reflections, 14% of even the old cap. So split them, the same way the geometry refinement's stopping rule is split: online keeps the transport-sized 10000, offline gets the full 65536, and the buffer headroom derives from the online one. Both still come from BraggPrediction so the cap, the prediction and the headroom cannot drift apart. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-30 22:49:16 +02:00
leonarski_fandClaude Opus 5	3e56d96921	Stills partiality: adopt the refined tilt only when it fits better ... RefineOne re-measured the image's correlation to the reference after writing the refined partialities - because --min-image-cc drops images by it - and then ignored what it measured. A crystal the tilt model suits worse than the fixed partiality it replaces kept the refined model anyway, and the refinement is on by default. Compare against the CC the crystal arrived with and put it back untouched when the refinement does not improve it, which is the same state a crystal with too few reflections to fit ends in. Also four things noted in review and left until now: AdaptiveThresholdTest.cpp was listed twice in the test target, AdaptiveThreshold.h was the one header in image_analysis/spot_finding not in its library's source list, CLAUDE.md said update_version.sh rewrites VERSION when it only reads it, and the CHANGELOG did not mention that image_scale_b is gone from the plot_type enum - which breaks a client that asks for that plot. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-30 22:03:18 +02:00
leonarski_fandClaude Opus 5	b161da05c1	Geometry refinement: bound offline reprocessing by iterations, not the clock ... The per-image refinement stopped on a wall-clock budget (40 ms, and 20 ms for the rotation-only extra pass). Online that is exactly right - the budget is real and an image that overruns it costs the acquisition. Offline it means the same file refines to a different lattice depending on what else the machine was doing at the time, which is not a property reprocessing should have. Bound it by iteration count instead when the caller is offline. IndexAndRefine takes the workflow as a constructor argument: the receiver asks for the wall-clock bound, rugnux and the viewer get the reproducible one. 50 iterations is Ceres' own default; the per-image problem converges well inside it, so it bounds the pathological case rather than the normal one - measured on five battery crystals, every number is unchanged from the timed version. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-07-30 20:42:04 +02:00