Merge branch 'fix-9fhc' into rc175: marCCD frames copied into the engine's own pinned buffer (fixes a pre-scan race); ModelValidation test follows the null on every run

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

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);
     }

image_analysis/image_preprocessing/ImagePreprocessorGPU.h

@@ -29,6 +29,8 @@ class ImagePreprocessorGPU : public ImagePreprocessor {
     // The uncompressed-image counterpart of input_reg: the region PinInputRegion last page-locked.
     const void *pinned_input = nullptr;
     size_t pinned_input_bytes = 0;
+    // Staging for a marCCD frame's stored 16-bit pixels (UINT16_WIDENED), allocated on first use.
+    CudaHostPtr<uint8_t> host_stored;
 
     std::vector<int32_t> cpu_image;
 

tests/ModelValidationTest.cpp

@@ -510,12 +510,14 @@ TEST_CASE("ModelValidation_ReindexesTheDataIntoTheModelsIndexing", "[ModelValida
     CHECK(twin.indexing_op == gemmi::Op::identity());
     CHECK(twin.change_of_basis_op == gemmi::Op::identity());
 
-    // Data already in the model's indexing: nothing moves, and there is nothing to arbitrate.
+    // Data already in the model's indexing: nothing moves. The null is run all the same, as on every
+    // validation that judges its model (ModelValidationSchedule::fit_test), and the model fits.
     const auto same = ValidateAgainstModel(own, cell, model, prefix, logger, p4, true, 1, 1.0);
     REQUIRE(same.ok);
     CHECK(same.change_of_basis_op == gemmi::Op::identity());
     CHECK(same.indexing_op == gemmi::Op::identity());
-    CHECK_FALSE(same.fit_tested);
+    CHECK(same.fit_tested);
+    CHECK(same.model_fits);
     CHECK(same.r_work < 0.15);
 
     for (const char *suffix : {"_2fofc.ccp4", "_fofc.ccp4", "_anom.ccp4", "_maps.mtz"})

tests/StoredImageDecodeGPUTest.cpp

@@ -7,7 +7,9 @@
 #ifdef JFJOCH_USE_CUDA
 
 #include <cstring>
+#include <atomic>
 #include <random>
+#include <thread>
 #include <vector>
 
 #include "../common/PixelMask.h"
@@ -184,4 +186,58 @@ TEST_CASE("ImagePreprocessorGPU_StoredEncodingsMatchCPU", "[ImagePreprocessorGPU
     }
 }
 
+// The pre-scan hands the frames it kept to two consumers at once - the spot measurement and the
+// defective-pixel search - each with its own engines, so one marCCD frame is uploaded by several
+// preprocessors concurrently while each of them is also moving on to other frames.
+TEST_CASE("ImagePreprocessorGPU_SharedStoredFramesManyEngines", "[ImagePreprocessorGPU]") {
+    if (get_gpu_count() == 0)
+        SKIP("No CUDA GPU present");
+
+    DetectorSetup detector = DetJF4M();
+    detector.SaturationLimit(65535);
+    DiffractionExperiment x(detector);
+    x.DetectorDistance_mm(80).BeamX_pxl(1030).BeamY_pxl(1080);
+    const size_t npixels = x.GetPixelsNum();
+    PixelMask mask(x);
+
+    std::mt19937 rng(23);
+    std::vector<std::vector<uint16_t>> frames(4, std::vector<uint16_t>(npixels));
+    for (auto &f : frames)
+        for (auto &v : f)
+            v = static_cast<uint16_t>(rng() % 60000);
+
+    std::vector<ImageStatistics> reference(frames.size());
+    {
+        ImagePreprocessorCPU cpu_pre(x, mask);
+        ImagePreprocessorBuffer cpu_buf(npixels);
+        for (size_t f = 0; f < frames.size(); f++) {
+            std::vector<int32_t> widened(frames[f].begin(), frames[f].end());
+            reference[f] = cpu_pre.Analyze(cpu_buf, reinterpret_cast<const uint8_t *>(widened.data()),
+                                           CompressedImageMode::Int32);
+        }
+    }
+
+    std::atomic<size_t> mismatches{0};
+    std::vector<std::thread> threads;
+    for (int t = 0; t < 4; t++)
+        threads.emplace_back([&, t] {
+            auto stream = std::make_shared<CudaStream>();
+            ImagePreprocessorGPU gpu_pre(x, mask, stream, /*copy_image_to_host=*/false);
+            ImagePreprocessorBufferGPU gpu_buf(npixels, /*host_mirror=*/false);
+            for (int i = 0; i < 200; i++) {
+                const size_t f = (i + t) % frames.size();
+                const CompressedImage image(frames[f].data(), npixels * sizeof(uint16_t), x.GetXPixelsNum(),
+                                            x.GetYPixelsNum(), CompressedImageMode::Int32,
+                                            CompressionAlgorithm::UINT16_WIDENED);
+                ImageStatistics stats{};
+                gpu_pre.AnalyzeCompressed(gpu_buf, image, stats);
+                if (!SameStats(stats, reference[f]))
+                    mismatches++;
+            }
+        });
+    for (auto &t : threads)
+        t.join();
+    CHECK(mismatches == 0);
+}
+
 #endif