2 Commits

Include renames

Author	SHA1	Message	Date
jungfrauandClaude Opus 5	f36cd88795	Give a 16-bit image a saturation code when its limit is 65534 ... The narrow encoding picked NARROW_BAD as the saturation code whenever the saturation limit was above UINT16_MAX - 2. That is one too strict. A real value is STRICTLY below the limit, so a limit of UINT16_MAX - 1 still leaves UINT16_MAX - 1 free to be the code; only a limit of the whole range has nothing to spare, and that is exactly the case where no pixel can be saturated, because the "is error" test claims UINT16_MAX first. At a limit of exactly 65534 the old condition therefore stored a saturated pixel as the masked code, and it widened back to INT32_MIN instead of INT32_MAX. Masked and saturated are not interchangeable: the strong-pixel search flags a saturated pixel unconditionally and a masked one never, so the overloaded core of the strongest spots dropped out of the strong-pixel mask. Bragg integration treats the two alike and the image statistics are taken from the raw value, so nothing downstream of those moved - which is why a battery over 24 crystals showed nothing. That value is not a corner case. GetByteDepthImage()-driven writing stores saturation_value = GetSaturationLimit() - 1 and the readers take it back as-is, so a 16-bit acquisition whose detector cutoff is at or above the full range comes back with a limit of exactly 65534. The one 16-bit dataset in the rotation test set declares 11963, which is safe, so it could not have caught this. Found by review, not by testing, because nothing tested the narrow path at all: every GPU test writes into the wide buffer directly and never asks for the narrow one, and every preprocessor test copies the image back to the host, which forces the wide path. So the test comes with the fix. It runs the wide and the narrow preprocessor over the same synthetic frame - values around each boundary, masked on a stride coprime with the value cycle so every value appears both masked and unmasked - across saturation limits of none, 5000, 0xFFFD, 0xFFFE and 0xFFFF, on both the host-upload and the device-decode entry point, and compares the statistics and every pixel. Against the old condition it fails with 625860 differing pixels; against this one it passes. Also makes the header self-contained: it uses __host__/__device__ and the CUDA vector types and only compiled because every includer happened to pull in cuda_runtime.h first. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>	2026-08-18 18:41:53 -04:00
jungfrauandClaude Opus 5	9f49e5abb7	Keep 16-bit images 16-bit through the GPU pipeline ... A detector reading out 16 bits had its frame widened to int32 the moment it was decoded, and every per-pixel pass over that frame then moved four bytes a pixel to carry two. Those passes - the ring statistics three times over, the strong-pixel search, spot extraction, the azimuthal and ROI integrators, Bragg integration - are the bulk of the image loop's device traffic, and 16 bits is the mode a fast acquisition runs in, which is exactly where throughput matters. The preprocessed image now keeps the width of its source. Two codes at the top of the 16-bit range carry the two special states, and they cannot collide with a real value: 0xFFFF masked, or the source's own bad-pixel marker. saturation a pixel at or above the saturation limit. 0xFFFE where the limit code leaves room - a 16-bit EIGER declares a count-rate limit of a few thousand, so there is room to spare - and 0xFFFF where the limit is the whole range, in which case the "is error" test has already claimed 0xFFFF, nothing can be saturated, and 0xFFFE stays a real value. Either way a real value is strictly below the saturation limit and so below both codes. Nothing is clipped and nothing is lost, and which code is in force is carried with the image rather than assumed. No engine learns a second convention. PixelView widens on load, so a masked pixel still reads as INT32_MIN and a saturated one as INT32_MAX, and every existing `v != INT32_MIN && v != INT32_MAX` test keeps its meaning. One code path, not two instantiations that can drift apart; the branch is on a pointer that is the same for every thread of every block, on kernels whose time is the loads it selects between. The vector loads are kept - four pixels still arrive in one transaction, 16 bytes wide or 8, whichever the image is. The wide path is unchanged, and is still taken for anything that is not a 16-bit source, and for any caller that wants the preprocessed image copied back to the host - that mirror is int32 and the CPU engines know only that convention. Measured on the one 16-bit dataset in the rotation test set, which is also the smallest detector in it (2.5M pixels, where per-pixel work is a small part of the loop): image loop 1.025 s -> 1.005 s at one GPU, whole run 5.64 s -> 5.52 s. The gain scales with the frame, so a 16M-pixel detector - where six full-frame passes are 86 % of the loop's GPU time - has much more to gain, and nothing here can measure that: every other dataset in the test set is stored 32-bit. Correctness on that dataset is exact where it can be: indexing rate, first-pass validation score and the integrated partial count are identical to the wide path, and its whole battery row - reflections, observations, space group, R_meas, CC1/2, ISa, mosaicity - is unchanged. Battery 6m17s -> 6m16s, 21/24 space groups, no failures. Also logs, once per run, the width the images are stored in, since it decides how much of the frame moves through every pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>	2026-08-18 01:52:27 -04:00