v1.0.0.rc-161 (#71)

This is an UNSTABLE release. It includes many experimental features, as well as many AI generated fixes. We recommend using rc.152 for production use.

* **rugnux: significantly better quality of results, and faster.** A large rework of integration, scaling, merging, geometry refinement and space-group determination, together with measurements the program previously made no attempt at - the direct beam before indexing, the beam stop, the goniometer rotation scale, and the stretches of a sweep the crystal did not deliver. A rotation dataset typically gains observations at better <I/sigma> and R_meas, and every `mx` and `scale` run writes a `<prefix>_report.txt` results report modelled on XDS's `CORRECT.LP`. Many defaults moved with it: spot detection is self-calibrating, beam-stop detection and rotation geometry post-refinement are on, resolution limits default to as far as the detector reaches, and ice-ring handling engages only where the crystal is measured to have ice.
* **jfjoch_viewer:** the beam-stop shadow, the detector calibration and the beam-centre measurement are reachable from "Analyze dataset"; the settings panel reports how the sample moved and how polarized the beam was; image rendering and interaction are faster.
* **Performance:** bitshuffle+LZ4 images are decoded on the GPU rather than on the host, with the bitshuffle inverse fused into preprocessing so the decompressed frame is never held in device memory.
* **Broker, writer, packaging and build:** image-slot lifetime and locking fixes, per-image datasets sized by the images actually written, the Debian/Ubuntu broker package renamed to `jfjoch`, and `image_analysis` compiling under MSVC again.

**Breaking change to the rugnux command line:**
* `--azint-only` and `--scale` are **removed**, replaced by `--mode azint` and `--mode scale`; the full pipeline is `--mode mx` and remains the default. A script passing the old flags now fails with the list of valid modes rather than silently running the wrong one.
* `-t`/`--stride` is **refused on rotation data**: skipping frames cuts every reflection's rocking curve, so the combined fulls and their partiality would be measured over frames the sweep never recorded. Select a contiguous range with `-s`/`-e` instead. `--mode azint` and `--force-still` still take a stride.

**Breaking changes to OpenAPI** - regenerate the client (`jfjoch-client` 1.0.0-rc.161, `frontend/src/client`) or read the affected fields as optional:
* `image_scale_b` is removed from the `plot_type` enum, so a client requesting that plot now gets an error rather than a curve.
* `azim_int_settings.high_q_recipA`, `spot_finding_settings.high_resolution_limit` and `spot_finding_settings.low_resolution_limit` are no longer `required`. All three mean "no limit at that end" when unset and are omitted from the response instead of carrying a placeholder value, which raises in a client generated from an rc.160-or-earlier spec. A value of 0 is still accepted and means the same thing.

**Breaking changes to the stored formats** - a consumer reading these fields must treat them as optional:
* The per-image image-scale B factor is no longer computed, so `/entry/MX/imageScaleBFactor` is absent from newly written HDF5 files and the corresponding key is absent from the CBOR DataMessage and END blocks. Files written by rc.160 and earlier still contain it and still open; nothing in the pipeline reads it any more.
* `_reflns.jfjoch_diffrn_ISa` now carries the whole-range `1/sqrt(a*b)` that XDS's ISa denotes, and the error-model `a` and `b` are reported in XDS's convention; the strong-reflection asymptote moves to `_reflns.jfjoch_diffrn_ISa_asymptotic`. **A file written by an earlier version carries the asymptote under the plain `ISa` name.**

Reviewed-on: #71
Co-authored-by: Filip Leonarski <filip.leonarski@psi.ch>

image_analysis/image_preprocessing/ImagePreprocessorGPU.cu

@@ -1,7 +1,17 @@
 // SPDX-FileCopyrightText: 2026 Filip Leonarski, Paul Scherrer Institute <filip.leonarski@psi.ch>
 // SPDX-License-Identifier: GPL-3.0-only
 
+#include <type_traits>
+
 #include "ImagePreprocessorGPU.h"
+#include "../../common/JFJochException.h"
+
+namespace {
+    void cuda_err(cudaError_t val) {
+        if (val != cudaSuccess)
+            throw JFJochException(JFJochExceptionCategory::GPUCUDAError, cudaGetErrorString(val));
+    }
+}
 
 template<class T>
 __global__ void preprocess_kernel(
@@ -89,30 +99,173 @@ __global__ void preprocess_kernel(
     }
 }
 
+// The per-pixel decision preprocess_kernel makes, in a form the fused kernel can reuse so the two
+// cannot drift apart. Priority: masked > error > saturated.
+template<class T>
+struct PreprocessAccum {
+    unsigned long long masked = 0, saturated = 0, error = 0;
+    long long max_v = INT64_MIN, min_v = INT64_MAX;
+
+    __device__ __forceinline__ int32_t Apply(T v, bool is_masked, T sat_value, T err_value) {
+        const bool is_err = (v == err_value);
+        const bool is_sat = !is_err && (v >= sat_value);
+        masked += is_masked;
+        error += (!is_masked && is_err);
+        saturated += (!is_masked && !is_err && is_sat);
+        if (!(is_masked || is_sat || is_err)) {
+            const int64_t val = (int64_t) v;
+            if (val > max_v) max_v = val;
+            if (val < min_v) min_v = val;
+        }
+        return is_masked ? INT32_MIN : is_err ? INT32_MIN : is_sat ? INT32_MAX : (int32_t) v;
+    }
+};
+
+// Reduce a block's thread-local accumulators into the image-wide statistics. Every accumulator is an
+// integer, so the result does not depend on the order the blocks arrive in.
+template<class T>
+__device__ __forceinline__ void FlushStats(PreprocessAccum<T> &l, ImageStatistics *stats) {
+    __shared__ unsigned long long s_masked, s_saturated, s_error;
+    __shared__ long long s_max, s_min;
+    if (threadIdx.x == 0) {
+        s_masked = 0; s_saturated = 0; s_error = 0; s_max = INT64_MIN; s_min = INT64_MAX;
+    }
+    __syncthreads();
+
+    atomicAdd(&s_masked, l.masked);
+    atomicAdd(&s_saturated, l.saturated);
+    atomicAdd(&s_error, l.error);
+    if (l.min_v <= l.max_v) {
+        atomicMax(&s_max, l.max_v);
+        atomicMin(&s_min, l.min_v);
+    }
+    __syncthreads();
+
+    if (threadIdx.x == 0) {
+        atomicAdd(&stats->masked_pixel_count, s_masked);
+        atomicAdd(&stats->saturated_pixel_count, s_saturated);
+        atomicAdd(&stats->error_pixel_count, s_error);
+        atomicMax((long long *) &stats->max_value, s_max);
+        atomicMin((long long *) &stats->min_value, s_min);
+    }
+}
+
+__device__ __forceinline__ uint64_t transpose8_fused(uint64_t x) {
+    uint64_t t;
+    t = (x ^ (x >> 7)) & 0x00aa00aa00aa00aaULL;  x = x ^ t ^ (t << 7);
+    t = (x ^ (x >> 14)) & 0x0000cccc0000ccccULL; x = x ^ t ^ (t << 14);
+    t = (x ^ (x >> 28)) & 0x00000000f0f0f0f0ULL; x = x ^ t ^ (t << 28);
+    return x;
+}
+
+// The bitshuffle inverse and the preprocessing in ONE pass. One thread owns one group of 8 elements
+// across every byte-plane, so once it has transposed its 8 bytes out of each plane it holds 8
+// complete elements and can emit 8 finished int32 pixels - the decompressed image never has to exist
+// in device memory at all. That removes a full-frame buffer per worker and a full-frame write plus
+// read from the pipeline.
+//
+// The last CUDA block (blockIdx.x == nblocks) finishes the handful of elements bitshuffle stores
+// verbatim; they are already on the device inside the uploaded chunk.
+template<class T, int ES>
+__global__ __launch_bounds__(256) void untranspose_preprocess_kernel(
+        const uint8_t *__restrict__ shuffled,
+        const BSLZ4BlockDesc *__restrict__ desc,
+        const uint8_t *__restrict__ mask,
+        int32_t *__restrict__ out,
+        ImageStatistics *__restrict__ stats,
+        T sat_value, T err_value, int nblocks,
+        const uint8_t *__restrict__ tail_src, uint32_t tail_elems, uint32_t tail_elem0) {
+    PreprocessAccum<T> l;
+
+    // The bytes are assembled in the unsigned counterpart of T - shifting a byte into the top of a
+    // signed type overflows it - and converted back at the end, which C++20 defines as two's
+    // complement reinterpretation. That is exactly what the byte order in the file means.
+    using U = typename std::make_unsigned<T>::type;
+
+    if (blockIdx.x == nblocks) {
+        if (threadIdx.x < tail_elems) {
+            U uv = 0;
+            #pragma unroll
+            for (int p = 0; p < ES; p++)
+                uv |= (U)((U)tail_src[threadIdx.x * ES + p] << (8 * p));
+            const T v = (T) uv;
+            out[tail_elem0 + threadIdx.x] = l.Apply(v, mask[tail_elem0 + threadIdx.x] != 0, sat_value, err_value);
+        }
+        FlushStats<T>(l, stats);
+        return;
+    }
+
+    const int b = blockIdx.x;
+    const uint32_t size = desc[b].nelem;             // bytes per plane
+    const uint8_t *in = shuffled + desc[b].out_off;
+    const uint32_t elem0 = desc[b].out_off / ES;     // first pixel of this block
+    const uint32_t n = size / 8;
+
+    for (uint32_t i = threadIdx.x; i < n; i += blockDim.x) {
+        uint64_t x[ES];
+        #pragma unroll
+        for (int p = 0; p < ES; p++) {
+            const uint8_t *pin = in + p * size;
+            uint64_t a = 0;
+            #pragma unroll
+            for (int k = 0; k < 8; k++) a |= (uint64_t)pin[k * n + i] << (8 * k);
+            x[p] = transpose8_fused(a);
+        }
+        int32_t o[8];
+        #pragma unroll
+        for (int k = 0; k < 8; k++) {
+            U uv = 0;
+            #pragma unroll
+            for (int p = 0; p < ES; p++) uv |= (U)((U)((x[p] >> (8 * k)) & 0xff) << (8 * p));
+            o[k] = l.Apply((T) uv, mask[elem0 + i * 8 + k] != 0, sat_value, err_value);
+        }
+        // elem0 is a multiple of 8 (bitshuffle blocks are), so this is 32-byte aligned.
+        int4 *dst = reinterpret_cast<int4 *>(out + elem0 + i * 8);
+        dst[0] = make_int4(o[0], o[1], o[2], o[3]);
+        dst[1] = make_int4(o[4], o[5], o[6], o[7]);
+    }
+    FlushStats<T>(l, stats);
+}
+
 ImagePreprocessorGPU::ImagePreprocessorGPU(const DiffractionExperiment &experiment, const PixelMask &mask,
-                                           std::shared_ptr<CudaStream> stream)
+                                           std::shared_ptr<CudaStream> stream, bool copy_image_to_host)
     : ImagePreprocessor(experiment),
       stream(stream),
-      gpu_mask(npixels),
-      gpu_decompressed_image(npixels * sizeof(uint32_t)),  // Overshoot - if input image is 1- or 2-byte, then it is still fine, while memory loss is minimal
+      copy_image_to_host(copy_image_to_host),
       gpu_stats(1),
       cpu_stats(1),
       cpu_stats_reg(cpu_stats) {
-    // Setup mask
+    // Setup mask. The same for every worker, so it is uploaded once per GPU and shared; keyed on the
+    // PixelMask's own vector, which the derived table is a pure function of.
     std::vector<uint8_t> mask_vec(npixels);
     for (int i = 0; i < npixels; i++)
         mask_vec[i] = (mask.GetMask().at(i) != 0);
+    gpu_mask = SharedDeviceTable(mask.GetMask().data(), npixels, mask_vec.data(), *stream);
 
-    cudaMemcpy(gpu_mask, mask_vec.data(), npixels, cudaMemcpyHostToDevice);
-
-    // Setup GPU settings
+    // Setup GPU settings. The current device, not device 0: workers are pinned round-robin across GPUs,
+    // so device 0's SM count can belong to a different card than the one these kernels launch on.
+    int device = 0;
+    cuda_err(cudaGetDevice(&device));
     cudaDeviceProp prop{};
-    cudaGetDeviceProperties(&prop, 0);
+    cuda_err(cudaGetDeviceProperties(&prop, device));
 
     threads = 128;
     blocks = 4 * prop.multiProcessorCount;
 }
 
+float ImagePreprocessorGPU::GetLastDecompressionTime_s() const {
+    return bslz4_decoder ? bslz4_decoder->GetDecodeTime_s() : 0.0f;
+}
+
+void ImagePreprocessorGPU::PinInputBuffer(std::vector<uint8_t> &buffer, size_t size) {
+    if (buffer.size() == size)
+        return;
+    // Unregister before the resize, which can move the buffer.
+    input_reg.unregister();
+    buffer.resize(size);
+    input_reg.rebind(buffer);
+}
+
 ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed_image, const uint8_t *image_ptr,
                                               CompressedImageMode image_mode) {
     switch (image_mode) {
@@ -133,30 +286,127 @@ ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed
     }
 }
 
+bool ImagePreprocessorGPU::AnalyzeCompressed(ImagePreprocessorBuffer &processed_image,
+                                             const CompressedImage &image,
+                                             ImageStatistics &stats) {
+    if (!BSLZ4DecoderGPU::Supports(image))
+        return false;   // caller decompresses on the host and uses Analyze()
+    if (image.GetUncompressedSize() != npixels * image.GetByteDepth())
+        return false;
+
+    if (!bslz4_decoder)
+        bslz4_decoder = std::make_unique<BSLZ4DecoderGPU>(npixels * sizeof(uint32_t), stream);
+
+    // LZ4 on the device, then ONE kernel that un-transposes the bitshuffle blocks and preprocesses
+    // them as it goes. The decompressed image is never materialised: the fused kernel reads the
+    // shuffled bytes and writes finished int32 pixels.
+    const BSLZ4ShuffledImage shuffled = bslz4_decoder->DecodeShuffled(image);
+
+    switch (image.GetMode()) {
+        case CompressedImageMode::Int8:
+            stats = UntransposeAndAnalyze<int8_t, 1>(processed_image, shuffled, INT8_MIN, INT8_MAX); return true;
+        case CompressedImageMode::Uint8:
+            stats = UntransposeAndAnalyze<uint8_t, 1>(processed_image, shuffled, UINT8_MAX, UINT8_MAX); return true;
+        case CompressedImageMode::Int16:
+            stats = UntransposeAndAnalyze<int16_t, 2>(processed_image, shuffled, INT16_MIN, INT16_MAX); return true;
+        case CompressedImageMode::Uint16:
+            stats = UntransposeAndAnalyze<uint16_t, 2>(processed_image, shuffled, UINT16_MAX, UINT16_MAX); return true;
+        case CompressedImageMode::Int32:
+            stats = UntransposeAndAnalyze<int32_t, 4>(processed_image, shuffled, INT32_MIN, INT32_MAX); return true;
+        case CompressedImageMode::Uint32:
+            stats = UntransposeAndAnalyze<uint32_t, 4>(processed_image, shuffled, UINT32_MAX, UINT32_MAX); return true;
+        default:
+            return false;   // Supports() already excludes these; belt and braces
+    }
+}
+
+// The device-decode counterpart of AnalyzeOnDevice: same per-pixel decision, same statistics, but
+// fed from the bitshuffled bytes rather than from a decompressed image.
+template<class T, int ES>
+ImageStatistics ImagePreprocessorGPU::UntransposeAndAnalyze(ImagePreprocessorBuffer &processed_image,
+                                                            const BSLZ4ShuffledImage &shuffled,
+                                                            T err_value, T sat_value) {
+    if (sat_value > saturation_limit)
+        sat_value = static_cast<T>(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));
+
+    // One CUDA block per bitshuffle block, plus one for the verbatim tail when there is one.
+    const int nb = shuffled.nblocks + (shuffled.tail_elems > 0 ? 1 : 0);
+    untranspose_preprocess_kernel<T, ES> <<< nb, 256, 0, *stream >>>(
+        shuffled.shuffled,
+        shuffled.desc,
+        gpu_mask->get(),
+        processed_image.getGPUBuffer(),
+        gpu_stats,
+        sat_value,
+        err_value,
+        shuffled.nblocks,
+        shuffled.tail_src,
+        shuffled.tail_elems,
+        shuffled.tail_elem0);
+    cuda_err(cudaGetLastError());
+
+    if (copy_image_to_host)
+        cuda_err(cudaMemcpyAsync(processed_image.data(), processed_image.getGPUBuffer(), npixels * sizeof(int32_t), cudaMemcpyDeviceToHost, *stream));
+    cuda_err(cudaMemcpyAsync(cpu_stats.data(), gpu_stats, sizeof(ImageStatistics), cudaMemcpyDeviceToHost, *stream));
+
+    // Check the synchronise: this path decodes bytes we did not produce, and cpu_stats still holds the
+    // sentinel the host wrote above, so an unchecked failure here returns it as if it were a real
+    // measurement and the caller's fallback to the host decoder never fires.
+    cuda_err(cudaStreamSynchronize(*stream));
+
+    // Only now can the device tell us whether every block actually decoded.
+    bslz4_decoder->ThrowIfDecodeFailed();
+
+    return cpu_stats[0];
+}
+
 template<class T>
 ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed_image,
                                               const uint8_t *input,
                                               T err_value,
                                               T sat_value) {
+    // Allocated here rather than in the constructor: only the host-upload path needs it, and the
+    // device-decode path - which is what BSHUF_LZ4 images take - never touches it. Overshoot to
+    // 4 bytes per pixel so a 1- or 2-byte image fits the same buffer.
+    if (!gpu_decompressed_image.get())
+        gpu_decompressed_image = CudaDevicePtr<uint8_t>(npixels * sizeof(uint32_t));
+
+    // On this engine's own stream, not the NULL stream: a NULL-stream copy implicitly synchronises with
+    // every blocking stream in the process, which serialised all workers behind whichever one was
+    // uploading. The stream is synchronised at the end of this function, so the ordering is unchanged.
+    cuda_err(cudaMemcpyAsync(gpu_decompressed_image, input, npixels * sizeof(T), cudaMemcpyHostToDevice, *stream));
+    return AnalyzeOnDevice<T>(processed_image, err_value, sat_value);
+}
+
+// Everything after the image is on the device, shared by the host-upload and the device-decode
+// entry points so the two cannot drift apart.
+template<class T>
+ImageStatistics ImagePreprocessorGPU::AnalyzeOnDevice(ImagePreprocessorBuffer &processed_image,
+                                                      T err_value, T sat_value) {
     if (sat_value > saturation_limit)
         sat_value = static_cast<T>(saturation_limit);
 
-    cudaMemcpy(gpu_decompressed_image, input, npixels * sizeof(T), cudaMemcpyHostToDevice);
-
     cpu_stats[0] = ImageStatistics{.max_value = INT64_MIN, .min_value = INT64_MAX};
-    cudaMemcpyAsync(gpu_stats, cpu_stats.data(), sizeof(ImageStatistics), cudaMemcpyHostToDevice, *stream);
+    cuda_err(cudaMemcpyAsync(gpu_stats, cpu_stats.data(), sizeof(ImageStatistics), cudaMemcpyHostToDevice, *stream));
     preprocess_kernel<T> <<< blocks, threads, 0, *stream >>>(
         reinterpret_cast<const T *>(gpu_decompressed_image.get()),
-        gpu_mask,
+        gpu_mask->get(),
         processed_image.getGPUBuffer(),
         gpu_stats,
         sat_value,
         err_value,
         npixels);
-    cudaMemcpyAsync(processed_image.data(), processed_image.getGPUBuffer(), npixels * sizeof(int32_t), cudaMemcpyDeviceToHost, *stream);
-    cudaMemcpyAsync(cpu_stats.data(), gpu_stats, sizeof(ImageStatistics), cudaMemcpyDeviceToHost, *stream);
+    cuda_err(cudaGetLastError());
+    // The preprocessed image is 4 bytes per pixel - by far the largest transfer here - and every GPU
+    // engine reads it straight from the device buffer, so it only comes back when a CPU engine needs it.
+    if (copy_image_to_host)
+        cuda_err(cudaMemcpyAsync(processed_image.data(), processed_image.getGPUBuffer(), npixels * sizeof(int32_t), cudaMemcpyDeviceToHost, *stream));
+    cuda_err(cudaMemcpyAsync(cpu_stats.data(), gpu_stats, sizeof(ImageStatistics), cudaMemcpyDeviceToHost, *stream));
 
-    cudaStreamSynchronize(*stream);
+    cuda_err(cudaStreamSynchronize(*stream));
 
     return cpu_stats[0];
 }