rc179: everything on rc178 after its merge into main (#88), except the CHANGELOG

The rc178 branch continued after #88 was squash-merged; this is its content from b9ada7914 (the tree
#88 squashed) to 094eee31a, as one change on main. The individual commits stay on the rc178 branch:

- still battery: --cold evicts each set's input from the page cache; Alvra pixel masks; manifest fixes
- AGIPD: converter to one rigid panel per module, int16 photons, mask folded in; still battery reads the converted sets
- DetectorModules: panel look-up by id, plane -> panel, plane bounds and a per-pixel walk; mask size checked
- tests: [PanelLayout] equivalence harness, stacked vs assembled panel layout
- Dense per-pixel maps through the panel table (panel layout stage 1)
- Panel id on spots and the panel table in the start message
- Viewer and preview show a panel detector as it lies on the detector plane
- Panel layout stage 3: per-panel spot finding (CPU+GPU, CCL) and raster algorithms on the plane raster
- Panel layout stage 2: prediction and Bragg integration through the panel table
- XTC reader: frames handed up as stored, streams indexed in parallel
- XFEL reader: integer frames stay integer, one gain, placement by panel
- EXTERNAL_TEST_DATA_STILLS.md: images used vs the entry for every set; drop "Not in the battery"
- still battery: EuXFEL sets read the depositors' Cheetah CXI files directly; AGIPD converter removed
- float32_to_int16.py: a non-finite pixel becomes -32768, the int16 error marker, not a zero count
- CrystFEL geometry: read the pixel masks and bad regions instead of ignoring them
- Panel layout stage 5: the XFEL reader hands up the stored frame and its panel table
- Panel table persisted in NXmx; modulation grid and pre-scan memo key on the plane
- still battery: Alvra sets take their pixel mask from the CrystFEL geometry, not --mask-good
- still battery: take the deposited model as mmCIF where the RCSB has no PDB-format file
- still battery: dimple gets PDB-format coordinates converted from an mmCIF-only deposition
- stills fixed cost: geometry refinement and pre-scan do the same work in less wall time
- stills geometry sample and trial spot finding run as many workers as the image pass
- Stills integration: background ring reach measured on empty apertures, radial curvature from the whole frame
- StillsPartialityRefine: the partiality width is the run's, measured on the crystals, not the prediction window's
- MergeOnTheFly: capture term and multiplicative outlier band for the stills merge (both opt-in)
- Stills given a cell: the written cell is the mean over the frames that are merged
- Stills: the radial curvature correction decided per run on the same empty apertures
- StillsPartialityRefine: measure the run width only on shells with signal
- Stills pre-scan: report the paired z of the radial-correction comparison
- module placement: the five concurrent fits each keep every thread
- still battery: score R-free against the deposited data through the same dimple protocol
- still battery: the R-free target is the lower of the deposition's dimple and published values
- Stills: background clip follows the radial curve; signal disk from the lattice spacing
- Stills keep an overloaded observation as a lower bound instead of dropping it
- Stills: drop the lattice-spacing signal disk; the correction decision back on the plain clip
- battery: R-free against the deposition without refinement only, one protocol for rotation and stills
- still battery: EuXFEL sets run on a 2,000-image subset for now
- stills geometry refinement: fit the detector distance and the cell together
- stills geometry refinement: a distance given with --detector-distance is held in the joint fit
- Metal integration engine: back to its panel-s5 kernels; Mac stills keep the old clip, no saturation count
- Metal integration test: the CPU comparison runs on a rotation experiment
- Stills: kept overloads go back to the drop where they are not concentrated at low resolution
- stills distance search: an absolute 30-3000 mm scan, a sample found without d-limits, the fit iterated
- battery results: XFEL still battery 20261022-1358 on rc178 0cf62452c - 12 sets, cold, R_free unrefined
- battery results: still summary gains the resolution of our merge and of the deposition
- Pre-scan masks a pixel at the saturation limit on most frames, as one holding the error value
- Stills keep an overload as a lower bound only where it has pixels at the limit; run-level gate removed
- XFEL reader: a per-frame mask is read the way the frames are, its chunk inflated outside the HDF5 lock
- stills distance search: every centre scanned with a wider acceptance, a larger sample where nothing indexes
- rugnux usage: --beam-center-check describes the absolute 30-3000 mm stills distance scan, not the old 60-160 % grid
- stills: detector tilt refined in the geometry pre-pass, kept where held-out frames agree
- CLAUDE.md: the changelog is written once, just before the merge into main - condensed one-liners of what users get
- broker: serve the placed model and its maps of the last rugnux live run (/rugnux/model.cif, /rugnux/map.ccp4)
- frontend: Mol* viewer of the placed model and its density on the Rugnux live page
- writer: _model.cif from rugnux live keeps the input model's metadata, as the served model.cif does
- rugnux live maps: one unit cell centred on the model; the viewer opens on all of it
- Anomalous difference map: ANOM/PHANOM in _maps.mtz, /rugnux/map.ccp4?type=anom, a magenta mesh in the viewer
- Mol* viewer: open with the maps drawn within 10 A of the centre of view
- Mol* viewer: aaredb's structure-viewer setup, on a black stage
- rugnux: the gain of a stills file that states none is read off its single-photon peak where the background resolves photons
- Mol* viewer: in line with AareDB's StructureView at its current head (8ece2c2)
- ACKNOWLEDGEMENT: the web frontend's Mol* viewer follows AareDB's setup (G. Gotthard, PSI)
- stills tilt decision: the modules stay where the run placed them
- still battery: no distance pins and no per-set tuning flags; gain from the data except where a file's baseline hides it
- battery results: rc178 6f6f69 - full still battery and full rotation battery with REFMAC
- Bragg integration: a CCD rotation sweep's sigmas at the measured gain, intensities in photons
- BackgroundGain: the gain of a frame that resolves photons is the spacing of its photon peaks, so a baseline left in the frame no longer hides it
- ShadowFinder: measure the background's azimuthal modulation instead of predicting it from the polarization
- BackgroundGain: one gain estimate for rotation and stills; the still battery no longer gives EuXFEL 152 its gain
- ShadowFinder: measure the modulation robustly under a wide one-sided shadow
- ShadowFinderTest: a one-sided shadow is not fitted as part of the background's modulation
- ShadowFinder: back to the predicted polarization factor, in the plane the background shows
- MXAnalysisAfterFPGA: the Bragg engine's edge mask in the image's layout, so raw / untransformed FPGA runs start again
- battery results: still battery on rc178 e54f55 replaces the earlier rc178 still tables
- XFELReader: per-frame mask and photon conversion without a library call or a branch per pixel
- stills geometry trial: the hypotheses a search asks are indexed side by side
- stills tilt decision: the fits start beside the pre-scan and the geometry trial
- XFELReader: a deflated frame is inflated and converted a piece at a time, while it is in the cache
- XFELReader: a byte-shuffled frame is unshuffled a piece at a time, a whole pixel per step
- GeometryRefiner: residuals carry only the blocks a fit moves; beam and distance one block
- GeometryRefiner: a spot's residual differentiates its observed and predicted halves apart
- docs: EXTERNAL_TEST_DATA lists every open-arm battery set, with corrected values and counts
- battery: doc_check.py cross-checks the manifests against the dataset pages
- stills module placement: the placement fit gets half the threads, each held-out fit an eighth
- still battery: the EuXFEL sets run on all their frames, as the published still runs did
- battery results: still battery on rc178 a0ce6c, cold on a quiet machine, replaces the earlier rc178 table
- macOS build: explicit size_t in the XFEL reader's row bounds; parentheses in the CPU preprocessor's select

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_015CeEGgpKdfkCTJb3YNMJcu

tests/BraggIntegrationEngineGPUTest.cpp

@@ -15,6 +15,7 @@
 #include "../common/BraggIntegrationSettings.h"
 #include "../common/DetectorSetup.h"
 #include "../common/DiffractionExperiment.h"
+#include "../common/JFJochMath.h"
 #include "../common/Reflection.h"
 #include "../image_analysis/bragg_integration/BraggIntegrationEngineCPU.h"
 #include "../image_analysis/bragg_integration/BraggIntegrationEngineGPU.h"
@@ -149,13 +150,15 @@ double CompareCpuVsGpu(IntegratorMode mode, std::optional<float> bandwidth_fwhm,
                      float r1 = 0.0f, float r2 = 0.0f, float r3 = 0.0f,
                      OverlapMode overlap = OverlapMode::Off, float companion_dx = 0.0f,
                      bool clip_spots = false, float odd_partiality = 1.0f,
-                     const SpotFootprint *footprint = nullptr) {
+                     const SpotFootprint *footprint = nullptr, float pixel_gain = 1.0f) {
     DiffractionExperiment experiment =
             MakeExperiment(mode, bandwidth_fwhm, clip_nsigma, radial, DetJF(2), stencil_k, r1, r2, r3,
                            overlap);
-    if (footprint) {
+    {
         BraggIntegrationSettings settings = experiment.GetBraggIntegrationSettings();
-        settings.Footprint(*footprint);
+        if (footprint)
+            settings.Footprint(*footprint);
+        settings.PixelGain(pixel_gain);
         experiment.ImportBraggIntegrationSettings(settings);
     }
     const size_t width = experiment.GetXPixelsNum();
@@ -204,9 +207,10 @@ double CompareCpuVsGpu(IntegratorMode mode, std::optional<float> bandwidth_fwhm,
         const auto out_clean = clean_cpu.Run(clean_image, clean_scene.predicted,
                                              clean_scene.predicted.size(), 5);
         // An unreadable pixel the (empty) mask does not explain reads as an overload, which keeps the
-        // reflection, flagged - so the cost shows in the measured ones.
+        // reflection - flagged, or on a still as a lower bound, clipped either way - so the cost shows
+        // in the complete ones.
         const auto measured = std::count_if(out_cpu.begin(), out_cpu.end(),
-                                            [](const Reflection &r) { return !r.overloaded; });
+                                            [](const Reflection &r) { return !r.clipped; });
         CHECK(static_cast<size_t>(measured) < out_clean.size());
     }
     for (size_t i = 0; i < out_cpu.size(); ++i) {
@@ -322,6 +326,16 @@ TEST_CASE("BraggIntegrationEngineGPU_MatchesCPU") {
         CompareCpuVsGpu(IntegratorMode::ProfileGaussian, std::nullopt, 4.0f, false, 8, 0.0f,
                         0.0f, 0.0f, 0.0f, OverlapMode::Off, 0.0f, false, 0.01f);
     }
+    // A CCD's gain: every variance at it, so the strong flag and the guard decide in photons, and the
+    // reflections leave in photons - computed independently in the two engines.
+    SECTION("ProfileGaussian pixel gain") {
+        CompareCpuVsGpu(IntegratorMode::ProfileGaussian, std::nullopt, 4.0f, false, 60, 0.0f,
+                        0.0f, 0.0f, 0.0f, OverlapMode::Off, 0.0f, false, 1.0f, nullptr, 4.0f);
+    }
+    SECTION("BoxSum pixel gain") {
+        CompareCpuVsGpu(IntegratorMode::BoxSum, std::nullopt, 4.0f, false, 60, 0.0f,
+                        0.0f, 0.0f, 0.0f, OverlapMode::Off, 0.0f, false, 1.0f, nullptr, 4.0f);
+    }
     SECTION("ProfileGaussian mono trim") { CompareCpuVsGpu(IntegratorMode::ProfileGaussian, std::nullopt, 0.0f); }
     // Unreadable pixels inside the signal disks themselves: the MINPK rescue keeps the reflection and
     // fits it over what is left, and the peak-loss rule throws back the ones that lost the profile's
@@ -362,6 +376,58 @@ TEST_CASE("BraggIntegrationEngineGPU_MatchesCPU") {
 // next one, so a reused engine has to give the same answer as a fresh one. A first frame whose spots
 // are somewhere else entirely is what would show a leftover mark: a stale mask pixel is read as a
 // neighbour's signal and dropped from the background ring, a stale owner steals a pixel outright.
+// The still's overload decision on the device: a clipped spot is kept as a lower bound, an overload by
+// an unreadable pixel is not - as on the CPU (BraggIntegrationEngineCPU_StillKeepsOnlySaturatedOverloads).
+TEST_CASE("BraggIntegrationEngineGPU_StillOverloadsMatchCPU", "[Integration]") {
+    if (get_gpu_count() == 0) {
+        WARN("No CUDA GPU present. Skipping BraggIntegrationEngineGPU_StillOverloadsMatchCPU");
+        return;
+    }
+    DiffractionExperiment experiment(DetJF(2));
+    experiment.DetectorDistance_mm(100.0f).IncidentEnergy_keV(WVL_1A_IN_KEV).BeamX_pxl(400.0f).BeamY_pxl(400.0f);
+    experiment.ImportBraggIntegrationSettings(BraggIntegrationSettings());
+    const int32_t limit = 500, bkg = 12;
+    experiment.Detector().SaturationLimit(limit);
+    const size_t width = experiment.GetXPixelsNum(), npixel = experiment.GetPixelsNum();
+
+    std::vector<int32_t> image(npixel, bkg);
+    auto spot = [&](int cx, int cy, double amplitude) {
+        for (int dy = -6; dy <= 6; ++dy)
+            for (int dx = -6; dx <= 6; ++dx) {
+                const int32_t v = bkg + static_cast<int32_t>(std::lround(amplitude * std::exp(-(dx * dx + dy * dy) / (2 * 1.3 * 1.3))));
+                image[(cy + dy) * width + cx + dx] = v >= limit ? INT32_MAX : v;
+            }
+    };
+    spot(600, 300, 2000.0);
+    spot(300, 300, 40.0);
+    image[300 * width + 300] = INT32_MIN;
+    std::vector<Reflection> predicted;
+    for (const auto &[x, y] : {std::pair{600, 300}, std::pair{300, 300}})
+        predicted.push_back(MakeReflection(static_cast<float>(x), static_cast<float>(y), 2.0f, x));
+
+    ImagePreprocessorBuffer cpu_image(npixel);
+    for (size_t i = 0; i < npixel; ++i)
+        cpu_image[i] = image[i];
+    BraggIntegrationEngineCPU cpu(experiment, PixelMask(experiment));
+    const auto out_cpu = cpu.Run(cpu_image, predicted, predicted.size(), 0);
+
+    auto stream = std::make_shared<CudaStream>();
+    ImagePreprocessorBufferGPU gpu_image(npixel);
+    for (size_t i = 0; i < npixel; ++i)
+        gpu_image[i] = image[i];
+    REQUIRE(cudaMemcpyAsync(gpu_image.getGPUBuffer(), gpu_image.getBuffer().data(),
+                            npixel * sizeof(int32_t), cudaMemcpyHostToDevice, *stream) == cudaSuccess);
+    BraggIntegrationEngineGPU gpu(experiment, stream, PixelMask(experiment));
+    const auto out_gpu = gpu.Run(gpu_image, predicted, predicted.size(), 0);
+
+    REQUIRE(out_cpu.size() == 2);
+    REQUIRE(out_gpu.size() == 2);
+    CHECK_FALSE(out_gpu[0].overloaded);
+    CHECK(out_gpu[1].overloaded);
+    CHECK(gpu.Counts().overloads_kept == 1);
+    CHECK(out_gpu[0].I == Catch::Approx(out_cpu[0].I).epsilon(0.001));
+}
+
 TEST_CASE("BraggIntegrationEngineGPU_ReusedEngineMatchesFresh") {
     if (get_gpu_count() == 0) {
         WARN("No CUDA GPU present. Skipping BraggIntegrationEngineGPU_ReusedEngineMatchesFresh");
@@ -413,6 +479,61 @@ TEST_CASE("BraggIntegrationEngineGPU_ReusedEngineMatchesFresh") {
     }
 }
 
+// The radial correction with its curve handed over from a per-ring profile (RadialBackgroundFromRings), as
+// on stills: the GPU reads the same curve and corrects the same empty aperture as the CPU.
+TEST_CASE("BraggIntegrationEngineGPU_RadialCurveFromRingsMatchesCPU", "[Integration]") {
+    if (get_gpu_count() == 0) {
+        WARN("No CUDA GPU present. Skipping");
+        return;
+    }
+    DiffractionExperiment experiment(DetJF(2));
+    experiment.DetectorDistance_mm(100.0f).IncidentEnergy_keV(WVL_1A_IN_KEV).BeamX_pxl(400.0f).BeamY_pxl(400.0f);
+    const float bx = 400.0f, by = 400.0f;
+    // Curved, plus a slope steep enough that the ring's flat clip would cut its beam-side half: the
+    // GPU must follow the curve in the clip as the CPU does (RadialClipOffset).
+    const auto background = [](double r) { return 40000.0 - 100.0 * r + 0.05 * (r - 300.0) * (r - 300.0); };
+    const size_t width = experiment.GetXPixelsNum(), npixel = experiment.GetPixelsNum();
+    ImagePreprocessorBuffer cpu_image(npixel);
+    ImagePreprocessorBufferGPU gpu_image(npixel);
+    for (size_t i = 0; i < npixel; ++i) {
+        cpu_image[i] = static_cast<int32_t>(std::lround(
+                background(std::hypot(static_cast<double>(i % width) - bx, static_cast<double>(i / width) - by))));
+        gpu_image[i] = cpu_image[i];
+    }
+    auto stream = std::make_shared<CudaStream>();
+    REQUIRE(cudaMemcpyAsync(gpu_image.getGPUBuffer(), gpu_image.getBuffer().data(),
+                            npixel * sizeof(int32_t), cudaMemcpyHostToDevice, *stream) == cudaSuccess);
+
+    const float low_q = 0.05f, dq = 0.01f;
+    const double F_px = 100.0 / 0.075, lambda = 1.0;
+    std::vector<float> rings(500);
+    for (size_t i = 0; i < rings.size(); ++i) {
+        const double q = low_q + (i + 0.5) * dq;
+        rings[i] = static_cast<float>(background(F_px * std::tan(2.0 * std::asin(q * lambda / (4.0 * PI)))));
+    }
+
+    std::vector<Reflection> predicted;
+    for (int i = 0; i < 8; ++i)
+        predicted.push_back(MakeReflection(bx + 300.0f * std::cos(0.7f * i), by + 300.0f * std::sin(0.7f * i), 2.0f, i + 1));
+
+    for (const auto mode : {IntegratorMode::BoxSum, IntegratorMode::ProfileGaussian}) {
+        experiment.ImportBraggIntegrationSettings(BraggIntegrationSettings().Integrator(mode).BackgroundRadialCorrection(true));
+        BraggIntegrationEngineCPU cpu(experiment, PixelMask(experiment));
+        BraggIntegrationEngineGPU gpu(experiment, stream, PixelMask(experiment));
+        cpu.RadialBackgroundFromRings(rings, static_cast<int32_t>(rings.size()), low_q, dq);
+        gpu.RadialBackgroundFromRings(rings, static_cast<int32_t>(rings.size()), low_q, dq);
+        const auto out_cpu = cpu.Run(cpu_image, predicted, predicted.size(), 0);
+        const auto out_gpu = gpu.Run(gpu_image, predicted, predicted.size(), 0);
+        REQUIRE(out_cpu.size() == predicted.size());
+        REQUIRE(out_gpu.size() == out_cpu.size());
+        for (size_t i = 0; i < out_cpu.size(); ++i) {
+            INFO("mode " << static_cast<int>(mode) << " reflection " << i);
+            CHECK(out_gpu[i].bkg == Catch::Approx(out_cpu[i].bkg).epsilon(0.001).margin(0.05));
+            CHECK(out_gpu[i].I == Catch::Approx(out_cpu[i].I).epsilon(0.01).margin(2.0));
+        }
+    }
+}
+
 // Hidden ([.]) benchmark: the raison d'etre of the GPU port is < 2 ms/frame (vs ~142 ms on the CPU
 // for ProfileIntegrate2D). Run explicitly with:  ./jfjoch_test "[bragg_bench]"
 TEST_CASE("BraggIntegrationEngineGPU_Benchmark", "[.][bragg_bench]") {