From dfd69fcb9172229eaabefd1897c9e23d07caaca2 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Mon, 12 Oct 2026 22:37:47 +0200 Subject: [PATCH] marCCD frames are widened on the device The marCCD reader read each frame's 16-bit scanlines into scratch and widened them to int32 on the worker thread (marccd::ReadInto, ~20 % of the image-loop worker CPU on 4nwv), and the 4-byte image was uploaded. ReadRawImage now reads the scanlines straight into the image's buffer and hands them over as a new read-only CompressionAlgorithm::UINT16_WIDENED (mode Int32): two bytes a pixel cross PCIe, from page-locked memory, and the preprocessing kernel widens them itself. It is not CompressedImageMode::Uint16, which calls 65535 an error pixel; a marCCD's 65535 is its saturation marker. So the kernel compares each pixel with int32's error marker and the int32 saturation cap - the comparisons are made on the pixel's value, in int64, rather than in the storage type - and decides every pixel exactly as for the int32 image the host would have widened it to. The host widening (JFJochDecompressPtr) is the CPU path. Test: ImagePreprocessorGPU_StoredEncodingsMatchCPU, with a frame carrying 65535 at a saturation of 65535, against the CPU preprocessor on the host-widened image. Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3 --- compression/CompressionAlgorithmEnum.h | 10 ++++---- compression/JFJochDecompress.h | 9 +++++++ image_analysis/MXAnalysisWithoutFPGA.cpp | 8 ++++--- .../ImagePreprocessorGPU.cu | 24 +++++++++++++++---- .../ImagePreprocessorGPU.h | 11 +++++---- reader/JFJochMarCCDReader.cpp | 13 +++++++++- reader/MarCCD.cpp | 19 +++++++++++---- reader/MarCCD.h | 4 ++++ rugnux/RugnuxPipeline.h | 8 ++++--- tests/StoredImageDecodeGPUTest.cpp | 18 +++++++++----- 10 files changed, 94 insertions(+), 30 deletions(-) diff --git a/compression/CompressionAlgorithmEnum.h b/compression/CompressionAlgorithmEnum.h index 56406b6dc..f9e5aa78e 100644 --- a/compression/CompressionAlgorithmEnum.h +++ b/compression/CompressionAlgorithmEnum.h @@ -7,8 +7,10 @@ // Eiger firmware 1.x. Nothing here produces it. It is last so the values of the others do not move - // the enum crosses the CBOR stream. // -// CBF_BYTE_OFFSET is READ-ONLY too, and never leaves the process: it is how the miniCBF reader hands -// over a frame as the file stores it, so that the GPU path can upload those bytes and decode them on -// the device. The image's mode is what it decodes to (Int32). +// CBF_BYTE_OFFSET and UINT16_WIDENED are READ-ONLY too, and never leave the process: they are how the +// miniCBF and marCCD readers hand over a frame as the file stores it, so that the GPU path can upload +// those bytes and decode them on the device. The image's mode is what they decode to (Int32). CBF_BYTE_OFFSET +// is the x-CBF_BYTE_OFFSET payload; UINT16_WIDENED is unsigned 16-bit pixels, each one the int32 of the +// same value - which is not CompressedImageMode::Uint16, whose largest value means an error pixel. enum class CompressionAlgorithm {BSHUF_LZ4, BSHUF_ZSTD, BSHUF_ZSTD_RLE, BSHUF_ZSTD_RLE_HUFF, NO_COMPRESSION, - LZ4_NO_SHUFFLE, CBF_BYTE_OFFSET}; + LZ4_NO_SHUFFLE, CBF_BYTE_OFFSET, UINT16_WIDENED}; diff --git a/compression/JFJochDecompress.h b/compression/JFJochDecompress.h index cc027d330..bfaa97ac5 100644 --- a/compression/JFJochDecompress.h +++ b/compression/JFJochDecompress.h @@ -240,6 +240,15 @@ inline void JFJochDecompressPtr(uint8_t *output, DecodeCBFByteOffset(source, source_size, reinterpret_cast(output), nelements); return; } + if (algorithm == CompressionAlgorithm::UINT16_WIDENED) { + if (elem_size != sizeof(int32_t) || source_size != nelements * sizeof(uint16_t)) + throw JFJochException(JFJochExceptionCategory::Compression, "Mismatch in size"); + const auto *src = reinterpret_cast(source); + auto *dst = reinterpret_cast(output); + for (size_t i = 0; i < nelements; i++) + dst[i] = src[i]; + return; + } size_t block_size = 0; if (algorithm != CompressionAlgorithm::NO_COMPRESSION) diff --git a/image_analysis/MXAnalysisWithoutFPGA.cpp b/image_analysis/MXAnalysisWithoutFPGA.cpp index 3409013e9..e7efc5ccf 100644 --- a/image_analysis/MXAnalysisWithoutFPGA.cpp +++ b/image_analysis/MXAnalysisWithoutFPGA.cpp @@ -121,9 +121,11 @@ void MXAnalysisWithoutFPGA::Analyze(DataMessage &output, // sees the same error it saw before any of this existed - but a GPU-side problem costs speed // rather than the acquisition. Unless the context is lost, which the host route cannot survive. cuda_throw_if_context_lost(); - // The miniCBF encoding is decoded on the device whenever there is one; the host decoder is the - // CPU path, not a slower second try. A stream the device refuses, the host refuses as well. - if (output.image.GetCompressionAlgorithm() == CompressionAlgorithm::CBF_BYTE_OFFSET) + // The miniCBF and marCCD encodings are decoded on the device whenever there is one; the host + // decoder is the CPU path, not a slower second try. A stream the device refuses, the host + // refuses as well. + if (output.image.GetCompressionAlgorithm() == CompressionAlgorithm::CBF_BYTE_OFFSET + || output.image.GetCompressionAlgorithm() == CompressionAlgorithm::UINT16_WIDENED) throw; spdlog::warn("Device decoding failed ({}), falling back to host decompression", e.what()); decoded_on_device = false; diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu index 5ec3cf53d..6c7c11ef5 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu @@ -21,8 +21,8 @@ __global__ void preprocess_kernel( const uint8_t *__restrict__ mask, int32_t *__restrict__ output, ImageStatistics *__restrict__ stats, - T saturation_limit, - T err_value, + int64_t saturation_limit, + int64_t err_value, int npixels) { // Shared block accumulators __shared__ unsigned long long s_masked; @@ -410,7 +410,8 @@ ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed bool ImagePreprocessorGPU::AnalyzeCompressed(ImagePreprocessorBuffer &processed_image, const CompressedImage &image, ImageStatistics &stats) { - if (image.GetCompressionAlgorithm() == CompressionAlgorithm::CBF_BYTE_OFFSET) { + if (image.GetCompressionAlgorithm() == CompressionAlgorithm::CBF_BYTE_OFFSET + || image.GetCompressionAlgorithm() == CompressionAlgorithm::UINT16_WIDENED) { stats = AnalyzeStoredInt32(processed_image, image); return true; } @@ -543,6 +544,19 @@ ImageStatistics ImagePreprocessorGPU::AnalyzeStoredInt32(ImagePreprocessorBuffer if (!gpu_decompressed_image.get()) gpu_decompressed_image = CudaDevicePtr(npixels * sizeof(uint32_t)); + if (image.GetCompressionAlgorithm() == CompressionAlgorithm::UINT16_WIDENED) { + // Uploaded as stored, two bytes a pixel, and widened by the preprocessing kernel itself: with + // int32's error marker and int32's saturation, each pixel is decided exactly as the int32 image + // the host would have widened it to. + 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(), + cudaMemcpyHostToDevice, *stream)); + return AnalyzeOnDevice(processed_image, INT32_MIN, INT32_MAX); + } + last_decode = LastDecode::CBF; if (!cbf_decoder) cbf_decoder = std::make_unique(stream); @@ -556,9 +570,9 @@ ImageStatistics ImagePreprocessorGPU::AnalyzeStoredInt32(ImagePreprocessorBuffer // entry points so the two cannot drift apart. template ImageStatistics ImagePreprocessorGPU::AnalyzeOnDevice(ImagePreprocessorBuffer &processed_image, - T err_value, T sat_value) { + int64_t err_value, int64_t sat_value) { if (sat_value > saturation_limit) - sat_value = static_cast(saturation_limit); + sat_value = saturation_limit; cpu_stats[0] = ImageStatistics{.max_value = INT64_MIN, .min_value = INT64_MAX}; cuda_err(cudaMemcpyAsync(gpu_stats, cpu_stats.data(), sizeof(ImageStatistics), cudaMemcpyHostToDevice, *stream)); diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.h b/image_analysis/image_preprocessing/ImagePreprocessorGPU.h index 4e7c8e640..1f095c036 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.h +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.h @@ -38,14 +38,17 @@ class ImagePreprocessorGPU : public ImagePreprocessor { // Which decoder the last AnalyzeCompressed used, for GetLastDecompressionTime_s. enum class LastDecode { None, BSLZ4, CBF } last_decode = LastDecode::None; - // The file encoding of the miniCBF reader (CompressionAlgorithmEnum.h): the stored bytes are - // uploaded, decoded into gpu_decompressed_image as the int32 image the host would have produced, - // and preprocessed by the same kernel as an uploaded int32 image. + // The file encodings of the miniCBF and marCCD readers (CompressionAlgorithmEnum.h): the stored + // bytes are uploaded, decoded into gpu_decompressed_image as the int32 image the host would have + // produced, and preprocessed by the same kernel as an uploaded int32 image. ImageStatistics AnalyzeStoredInt32(ImagePreprocessorBuffer &processed_image, const CompressedImage &image); template ImageStatistics Analyze(ImagePreprocessorBuffer &processed_image, const uint8_t *input, T err_value, T sat_value); // Preprocess an image already sitting in gpu_decompressed_image (the host-upload path). - template ImageStatistics AnalyzeOnDevice(ImagePreprocessorBuffer &processed_image, T err_value, T sat_value); + // err_value and sat_value are compared with the pixel's value, not its bits, so a narrower T can + // carry the int32 semantics of the image it widens to. + template ImageStatistics AnalyzeOnDevice(ImagePreprocessorBuffer &processed_image, + int64_t err_value, int64_t sat_value); // Preprocess straight out of the bitshuffled bytes (the device-decode path). Same per-pixel // decision and same statistics as AnalyzeOnDevice, with the un-transpose folded in. template ImageStatistics UntransposeAndAnalyze(ImagePreprocessorBuffer &processed_image, diff --git a/reader/JFJochMarCCDReader.cpp b/reader/JFJochMarCCDReader.cpp index 317bd5e7b..504d6e410 100644 --- a/reader/JFJochMarCCDReader.cpp +++ b/reader/JFJochMarCCDReader.cpp @@ -203,7 +203,18 @@ bool JFJochMarCCDReader::HasImage(int64_t image_number) const { bool JFJochMarCCDReader::ReadRawImage(int64_t image_number, JFJochReaderRawImage &image) { if (!HasImage(image_number)) return false; - image.image = DecodeInto(image_number, image.image_buffer, image.read_buffer); + // The pixels as the file stores them, two bytes each: the consumer widens them, on the device where + // there is one, so the host neither widens nor uploads a four-byte image. + const size_t npixel = static_cast(header0_.nx) * static_cast(header0_.ny); + image.image_buffer.resize(npixel * sizeof(uint16_t)); + const auto h = marccd::ReadRaw(files_[image_number], reinterpret_cast(image.image_buffer.data()), + npixel); + if (h.nx != header0_.nx || h.ny != header0_.ny) + throw JFJochException(JFJochExceptionCategory::InputParameterInvalid, + "marCCD image size differs from the first image of the sweep"); + image.image = CompressedImage(image.image_buffer.data(), image.image_buffer.size(), + static_cast(header0_.nx), static_cast(header0_.ny), + CompressedImageMode::Int32, CompressionAlgorithm::UINT16_WIDENED); return true; } diff --git a/reader/MarCCD.cpp b/reader/MarCCD.cpp index 68ac7c171..114cf9a16 100644 --- a/reader/MarCCD.cpp +++ b/reader/MarCCD.cpp @@ -299,7 +299,7 @@ Header ReadHeader(const std::string &path) { return Parse(h, path); } -Header ReadInto(const std::string &path, int32_t *out, size_t capacity, std::vector &scratch) { +Header ReadRaw(const std::string &path, uint16_t *out, size_t capacity) { const Header header = ReadHeader(path); const size_t npixel = static_cast(header.nx) * static_cast(header.ny); if (npixel > capacity) @@ -331,14 +331,25 @@ Header ReadInto(const std::string &path, int32_t *out, size_t capacity, std::vec if (scanline_bytes != static_cast(TIFFScanlineSize(tiff))) throw JFJochException(JFJochExceptionCategory::TIFFGeneratorError, path + ": TIFFScanlineSize mismatch"); - scratch.resize(scanline_bytes * length); + auto *dst = reinterpret_cast(out); for (uint32_t row = 0; row < length; row++) { - if (TIFFReadScanline(tiff, scratch.data() + static_cast(row) * scanline_bytes, - row, 0) < 0) + if (TIFFReadScanline(tiff, dst + static_cast(row) * scanline_bytes, row, 0) < 0) throw JFJochException(JFJochExceptionCategory::TIFFGeneratorError, path + ": TIFFReadScanline error"); } + return header; +} + +Header ReadInto(const std::string &path, int32_t *out, size_t capacity, std::vector &scratch) { + const Header header = ReadHeader(path); + const size_t npixel = static_cast(header.nx) * static_cast(header.ny); + if (npixel > capacity) + throw JFJochException(JFJochExceptionCategory::ArrayOutOfBounds, + path + ": image is larger than the buffer given for it"); + scratch.resize(npixel * sizeof(uint16_t)); + ReadRaw(path, reinterpret_cast(scratch.data()), npixel); + // The stored pixels are unsigned 16-bit; the rest of the code works in signed 32-bit, so they // are widened here and nothing downstream has to know what the file held. const auto *src = reinterpret_cast(scratch.data()); diff --git a/reader/MarCCD.h b/reader/MarCCD.h index 809a84937..dbab48187 100644 --- a/reader/MarCCD.h +++ b/reader/MarCCD.h @@ -58,6 +58,10 @@ Header ReadHeader(const std::string &path); // Header + pixels of one file. out is resized to nx * ny. Header Read(const std::string &path, std::vector &out); +// Header + the stored unsigned 16-bit pixels, as the file holds them, into memory the caller already +// has. Throws if the image does not fit in capacity pixels. +Header ReadRaw(const std::string &path, uint16_t *out, size_t capacity); + // The same, widening the stored uint16 into memory the caller already has. scratch carries the // 16-bit strip between frames, so a worker walking a sweep allocates once. Throws if the image // does not fit in capacity pixels. diff --git a/rugnux/RugnuxPipeline.h b/rugnux/RugnuxPipeline.h index 73fe2158d..62c94127a 100644 --- a/rugnux/RugnuxPipeline.h +++ b/rugnux/RugnuxPipeline.h @@ -106,9 +106,11 @@ namespace rugnux_internal { } space_.notify_one(); try { - // A frame its reader handed over as the file stores it - miniCBF - is written as the - // int32 image it decodes to: that encoding never leaves the process. - if (job.msg.image.GetCompressionAlgorithm() == CompressionAlgorithm::CBF_BYTE_OFFSET) { + // A frame its reader handed over as the file stores it - miniCBF, marCCD - is written + // as the int32 image it decodes to: those encodings never leave the process. + const CompressionAlgorithm algorithm = job.msg.image.GetCompressionAlgorithm(); + if (algorithm == CompressionAlgorithm::CBF_BYTE_OFFSET + || algorithm == CompressionAlgorithm::UINT16_WIDENED) { const CompressedImage stored = job.msg.image; stored.GetUncompressed(decoded_); job.msg.image = CompressedImage(decoded_.data(), decoded_.size(), stored.GetWidth(), diff --git a/tests/StoredImageDecodeGPUTest.cpp b/tests/StoredImageDecodeGPUTest.cpp index 6e133845c..12e0256b8 100644 --- a/tests/StoredImageDecodeGPUTest.cpp +++ b/tests/StoredImageDecodeGPUTest.cpp @@ -17,9 +17,9 @@ #include "../image_analysis/image_preprocessing/ImagePreprocessorGPU.h" #include "../image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.h" -// The miniCBF frames reach the GPU as their files store them, and the device has to turn them into -// exactly the int32 image the host decoder would have: the same pixels, the same refusal of a stream -// that ends early. +// The miniCBF and marCCD frames reach the GPU as their files store them, and the device has to turn +// them into exactly the int32 image the host decoder would have: the same pixels, the same refusal of +// a stream that ends early. namespace { // x-CBF_BYTE_OFFSET with the smallest container for each delta; force64 puts a delta in the 15-byte @@ -132,8 +132,9 @@ TEST_CASE("CBFDecoderGPU_ShortStreamIsRefused", "[ImagePreprocessorGPU]") { } } -// Through the preprocessor, as the image loop calls it: the stored encoding against the CPU -// preprocessor fed what the host decodes it to. +// Through the preprocessor, as the image loop calls it: the stored encodings against the CPU +// preprocessor fed what the host decodes them to. The marCCD frame carries 65535, its saturation +// marker, which a Uint16 image would call an error pixel - widened to int32 it is saturated. TEST_CASE("ImagePreprocessorGPU_StoredEncodingsMatchCPU", "[ImagePreprocessorGPU]") { if (get_gpu_count() == 0) SKIP("No CUDA GPU present"); @@ -151,10 +152,15 @@ TEST_CASE("ImagePreprocessorGPU_StoredEncodingsMatchCPU", "[ImagePreprocessorGPU std::mt19937 rng(17); const auto pixels = MakePixels(npixels, 9); const auto cbf = EncodeByteOffset(pixels, rng); + std::vector marccd(npixels); + for (size_t i = 0; i < npixels; i++) + marccd[i] = (i % 1500 == 0) ? 65535 : static_cast(rng() % 70000); const std::vector images = { CompressedImage(cbf.data(), cbf.size(), x.GetXPixelsNum(), x.GetYPixelsNum(), CompressedImageMode::Int32, - CompressionAlgorithm::CBF_BYTE_OFFSET) + CompressionAlgorithm::CBF_BYTE_OFFSET), + CompressedImage(marccd.data(), marccd.size() * sizeof(uint16_t), x.GetXPixelsNum(), x.GetYPixelsNum(), + CompressedImageMode::Int32, CompressionAlgorithm::UINT16_WIDENED) }; auto stream = std::make_shared(); ImagePreprocessorGPU gpu_pre(x, mask, stream, /*copy_image_to_host=*/true);