From 34092a510f8e68954c447bf450e69cf0a6089479 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Tue, 13 Oct 2026 05:21:33 +0200 Subject: [PATCH] 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) Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3 --- .../ImagePreprocessorGPU.cu | 11 +++- .../ImagePreprocessorGPU.h | 2 + tests/StoredImageDecodeGPUTest.cpp | 56 +++++++++++++++++++ 3 files changed, 67 insertions(+), 2 deletions(-) diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu index 6c7c11ef5..41362f129 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu @@ -2,6 +2,7 @@ // SPDX-License-Identifier: GPL-3.0-only #include +#include #include #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(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(processed_image, INT32_MIN, INT32_MAX); } diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.h b/image_analysis/image_preprocessing/ImagePreprocessorGPU.h index 1f095c036..7b2b60c84 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.h +++ b/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 host_stored; std::vector cpu_image; diff --git a/tests/StoredImageDecodeGPUTest.cpp b/tests/StoredImageDecodeGPUTest.cpp index 12e0256b8..8663a8079 100644 --- a/tests/StoredImageDecodeGPUTest.cpp +++ b/tests/StoredImageDecodeGPUTest.cpp @@ -7,7 +7,9 @@ #ifdef JFJOCH_USE_CUDA #include +#include #include +#include #include #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> frames(4, std::vector(npixels)); + for (auto &f : frames) + for (auto &v : f) + v = static_cast(rng() % 60000); + + std::vector reference(frames.size()); + { + ImagePreprocessorCPU cpu_pre(x, mask); + ImagePreprocessorBuffer cpu_buf(npixels); + for (size_t f = 0; f < frames.size(); f++) { + std::vector widened(frames[f].begin(), frames[f].end()); + reference[f] = cpu_pre.Analyze(cpu_buf, reinterpret_cast(widened.data()), + CompressedImageMode::Int32); + } + } + + std::atomic mismatches{0}; + std::vector threads; + for (int t = 0; t < 4; t++) + threads.emplace_back([&, t] { + auto stream = std::make_shared(); + 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