marCCD upload: stage the stored pixels instead of page-locking the frame in place

A UINT16_WIDENED frame (marCCD) was page-locked where it lay with PinInputRegion, which keeps the
registration until the engine moves on to another frame. The pre-scan hands its kept frames to the
spot measurement and to the defective-pixel search at once, each with its own engines, so one
engine could unregister a frame while another engine's copy from it was still in flight. That
takes the context down ("unspecified launch failure"), seen as a rare crash at the end of the
pre-scan on a 3072x3072 marCCD sweep (1 run in 30 reproduced it).

The stored bytes now go through the engine's own page-locked staging buffer, as the CBF and BSLZ4
decoders already do. The pixels decoded are unchanged.

Test: ImagePreprocessorGPU_SharedStoredFramesManyEngines - four engines upload four shared marCCD
frames concurrently; it aborted with the launch failure before the change and passes after it.

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

image_analysis/image_preprocessing/ImagePreprocessorGPU.cu

@@ -2,6 +2,7 @@
 // SPDX-License-Identifier: GPL-3.0-only
 
 #include <algorithm>
+#include <cstring>
 #include <type_traits>
 
 #include "ImagePreprocessorGPU.h"
@@ -551,8 +552,14 @@ ImageStatistics ImagePreprocessorGPU::AnalyzeStoredInt32(ImagePreprocessorBuffer
         last_decode = LastDecode::None;
         if (image.GetCompressedSize() != npixels * sizeof(uint16_t))
             throw JFJochException(JFJochExceptionCategory::InputParameterInvalid, "Mismatch in size");
-        PinInputRegion(image.GetCompressed(), image.GetCompressedSize());
-        cuda_err(cudaMemcpyAsync(gpu_decompressed_image, image.GetCompressed(), image.GetCompressedSize(),
+        // Staged through this engine's own page-locked buffer, as the CBF and BSLZ4 decoders stage
+        // theirs, and not page-locked in place: the pre-scan hands one kept frame to several engines
+        // at once, and an engine that unregisters a frame another one is still copying from takes the
+        // whole context down. AnalyzeOnDevice waits for the stream, so the buffer is free again here.
+        if (!host_stored.get())
+            host_stored = CudaHostPtr<uint8_t>(npixels * sizeof(uint16_t));
+        std::memcpy(host_stored.get(), image.GetCompressed(), image.GetCompressedSize());
+        cuda_err(cudaMemcpyAsync(gpu_decompressed_image, host_stored.get(), image.GetCompressedSize(),
                                  cudaMemcpyHostToDevice, *stream));
         return AnalyzeOnDevice<uint16_t>(processed_image, INT32_MIN, INT32_MAX);
     }