The scaling engine held far more on the card than any step needed, and the
image passes' pooled memory was still reserved when it started.
- The image passes' allocation pool is trimmed before the scaling's arrays
are allocated (they are not pooled): 4.3-8.8 GB of the 16 GB card was held
by others when the scaling started, now 1.5-2.3 GB (context, cuFFT and the
run's FFT indexers, which live for the whole run).
- A new grouping or combine frees the previous one's fulls, CSRs and merge
results before allocating its own; the final merge of a large cell used to
hold the P1 search's per-group arrays beside its own.
- The partials' group CSR goes at the combine; the arrays only the fulls'
scaling and download read (records, zeta, frame CSR, group means, image
number, px/py, clipped) go once the fulls are on the host.
- Merge buffers live only in the call that needs them: the error-model sums
and sample places, the outlier medians, half-sets and CC factors, the R_meas
sums; MergeAccum's per-group results go once downloaded (MergeAccumRelease).
- The combine's per-run scratch is local to the combine.
- Ingest: both radix sorts run on cub::DoubleBuffer (no cub copies of keys and
values), the sort outputs are allocated after the raw-hkl key is gathered and
freed, the uncut runs and old order go before the observations are cut, and
the delta_phi re-solve takes the hkl from its run, so the per-observation
64-bit key no longer outlives the sort.
Measured with per-array device accounting on an RTX 5080, engine peak after
the GPU/CPU decision (MB, before -> after; largest engine of each set):
4nwv 32.7 M obs, 15.4 M hkl: 7536 -> 4265
8qaw 122.7 M obs, 8.6 M hkl: 11716 -> 9327 (ingest 12533 -> 10756)
9v16 54.7 M obs, 5.7 M hkl: 5831 -> 4283
7qij 78.1 M obs, 2.2 M hkl: 6108 -> 5242 (ingest 7978 -> 6768)
8sa8 50.0 M obs, 5.0 M hkl: 5490 -> 4205
The ingest's own peak went from ~104 to 82-100 bytes per reflection.
GPUScalingPeakBytes refitted on fourteen engines of these five sets to the
peak after the decision it makes (the ingest has passed by then): 81 B/obs +
291 B/hkl -> 63 + 115, within 12%. DeviceBytes refitted as an upper bound:
67/260/200 -> 66/94/52 per obs/event/hkl (up to 18% over), ingest
103/58 -> 89/56 per reflection/run. With it the all-observation arm now runs
beside the first merge on a 16 GB card for 4nwv, 9v16 and 8sa8.
p.hkl byte-identical to the rc175 baseline on 15 battery sets (the five above,
11if, 5reo, 6oel, 9jzo, 9min, metformin, lyso_x06da_ref, thau_bl1a_3p8keV,
insu_I_x06da_5keV_2, myob_x06da), p.mtz identical apart from the history;
battery compare shows no change. [RotationScale] and
ErrorModel_DeviceFitIsTheHostFit pass.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
leonarski_fandClaude Opus 5.5
c130c79236