Fold the beam-stop batch before its frame size changes

ShadowAccumulatorGPU::Add sized the raw buffer before closing the pending batch, and
EnsureRawCapacity assigned frame_bytes on entry. FoldPending strides `raw` by frame_bytes, so a
frame of a different size arriving mid-batch made the already-decoded frames fold with the new
stride: every pixel of the pending batch read from the wrong offset, silently, with no error. The
depth-change branch that exists to handle exactly this ran one step too late to help.

Fold first, then resize, then adopt the new stride. The batch also closes on a change of frame
size, not only of pixel mode - a batch is one layout, and the mode alone does not fix the layout.
The decoder was likewise built once from the first frame and never rebuilt, so it is now rebuilt
when the frame size changes; without that the mixed-size path this commit repairs would still
decode into a buffer of the wrong size.

Also calls Gpu() once in ShadowFinder::AddImage instead of twice - it takes and releases a mutex
each time, once per image.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_011n8riB6X59oRjkrSHzNPAU

image_analysis/beam_stop/ShadowAccumulatorGPU.cu

@@ -105,16 +105,18 @@ bool ShadowAccumulatorGPU::Supports(const CompressedImage &image) {
 }
 
 void ShadowAccumulatorGPU::EnsureRawCapacity(size_t bytes_per_frame) {
+    if (bytes_per_frame * BATCH > raw_capacity) {
+        // The constructor already sized this for the widest pixel type, so in practice this only
+        // runs if that guess was too small. cudaMalloc and cudaFree both synchronise the whole
+        // device, which is why it is never done per frame.
+        FoldPending();
+        cuda_err(cudaStreamSynchronize(*stream));
+        raw_capacity = bytes_per_frame * BATCH;
+        raw = CudaDevicePtr<uint8_t>(raw_capacity);
+    }
+    // Only once the pending batch has been folded: FoldPending strides the buffer by frame_bytes,
+    // so it has to keep describing the frames already in it until they are gone.
     frame_bytes = bytes_per_frame;
-    if (bytes_per_frame * BATCH <= raw_capacity)
-        return;
-    // The constructor already sized this for the widest pixel type, so in practice this only runs
-    // if that guess was too small. cudaMalloc and cudaFree both synchronise the whole device, which
-    // is why it is never done per frame.
-    FoldPending();
-    cuda_err(cudaStreamSynchronize(*stream));
-    raw_capacity = bytes_per_frame * BATCH;
-    raw = CudaDevicePtr<uint8_t>(raw_capacity);
 }
 
 // Fold the frames decoded so far into the projection. One pass over the accumulator for the whole
@@ -150,12 +152,15 @@ void ShadowAccumulatorGPU::Add(const CompressedImage &image) {
         throw JFJochException(JFJochExceptionCategory::InputParameterInvalid,
                               "ShadowAccumulatorGPU: image size does not match the detector");
 
-    if (!decoder)
+    // A batch holds one pixel type and one frame size; a change of either closes the batch first,
+    // while frame_bytes and pending_mode still describe the frames already in it.
+    if (pending > 0 && (image.GetMode() != pending_mode
+                        || image.GetUncompressedSize() != frame_bytes))
+        FoldPending();
+
+    if (!decoder || image.GetUncompressedSize() != frame_bytes)
         decoder = std::make_unique<BSLZ4DecoderGPU>(image.GetUncompressedSize(), stream);
     EnsureRawCapacity(image.GetUncompressedSize());
-    // A batch holds one pixel type; a change of depth mid-run closes the batch first.
-    if (pending > 0 && image.GetMode() != pending_mode)
-        FoldPending();
     pending_mode = image.GetMode();
 
     decoder->Decode(image, raw.get() + static_cast<size_t>(pending) * frame_bytes);

image_analysis/beam_stop/ShadowFinder.cpp

@@ -427,10 +427,9 @@ void ShadowFinder::AddImage(const DataMessage &data, std::vector<uint8_t> &buffe
                               "ShadowFinder: shard out of range");
 
 #ifdef JFJOCH_USE_CUDA
-    if (ShadowAccumulatorGPU::Supports(data.image) && Gpu()) {
-        // One device, so the frames queue here - but each is only a chunk upload plus two kernels,
-        // and nothing was decompressed on the host to get this far.
-        ShadowAccumulatorGPU *acc = Gpu();
+    // One device, so the frames queue here - but each is only a chunk upload plus two kernels, and
+    // nothing was decompressed on the host to get this far.
+    if (ShadowAccumulatorGPU *acc = ShadowAccumulatorGPU::Supports(data.image) ? Gpu() : nullptr) {
         std::unique_lock ul(gpu_mutex);
         try {
             acc->Add(data.image);