- The resident ingest reads the reflections twice: first only what the sort
and the cut need (frame position, resolution, background, raw-hkl key),
then - once the GPU/CPU decision has been made on the exact counts - the
observations the cut kept, written straight to their place. The
observation arrays and the sort buffers are never on the card together,
and no observation field is cut (held twice) any more.
- A set too large even for that sort scales on the CPU before anything is
allocated on the card (ScaleOnCPUWhereSortTooLarge, same 80%-of-total rule).
- The predicted detector position never goes to the card: it is taken on the
host, and the combine emits the index of each full's peak partial instead
of its px/py (8 B less per observation and 4 per full).
- A full carries its raw-hkl run index instead of h, k, l (8 B less per
full); the constant-1 partiality array of the fulls is gone.
- The fulls' CSR sorts and the error-model rank sort run on cub::DoubleBuffer.
Measured with per-array device accounting on an RTX 5080, largest engine of
each set, before (c130c7923) -> after, MB:
set obs/hkl ingest peak (B/refl) post-decision peak
4nwv 32.7 M/15.4 M 3277 (100) -> 2336 (72) 4265 -> 3625
8qaw 122.7 M/ 8.6 M 10756 (88) -> 7421 (61) 9327 -> 8156
9v16 54.7 M/ 5.7 M 4838 (89) -> 3419 (63) 4283 -> 3774
7qij 78.1 M/ 2.2 M 6768 (87) -> 4707 (60) 5242 -> 4707
8sa8 50.0 M/ 5.0 M 4107 (82) -> 3134 (63) 4205 -> 3603
The sort before the decision now holds 32 B per reflection (46 on a P1-like
set with many runs); the engine's overall peak is after the decision on all
of them.
GPUScalingPeakBytes refitted on the same fourteen engines: 63 + 115 ->
56 B/obs + 99 B/hkl, within 9%. DeviceBytes as an upper bound: 66/94/52 ->
58/30/86 per obs/event/hkl, ingest 89/56 -> 60/42 per reflection/run.
p.hkl byte-identical to the rc175 baseline on 15 battery sets (4nwv, 8qaw,
9v16, 7qij, 8sa8, 11if, 5reo, 6oel, 9jzo, 9min, metformin, lyso_x06da_ref,
thau_bl1a_3p8keV, insu_I_x06da_5keV_2, myob_x06da); compare shows no change.
[RotationScale] (RotationIngestGPU now also checks the host positions) and
ErrorModel_DeviceFitIsTheHostFit pass.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
leonarski_fandClaude Opus 5.5
f46dc303a3