rugnux pre-scan: projection and defective pixels per card, the sample read once, beam stop memoised

The beam-stop projection (ShadowAccumulatorGPU) and the defective-pixel sums (HotPixelFinderGPU) are
now kept one per card, fed by workers pinned t % gpus like the other pre-scan workers, and folded
into the first card when they are read (cudaMemcpyPeer + an integer add / max-with-count kernel).
All of them are integer sums or a maximum with the first-value rule, so the fold is order-free and
the result is identical to one card's. With one card nothing is folded.

The pre-scan sample is read once and kept (compressed as stored; up to 2 GiB, a frame past that is
read again where it is wanted) for the spot measurement and the defective pixels, which read every
frame again each before - a CBF/marCCD/SMV frame was decoded on the host three times.

A beam-stop detection that finds no shadow leaves the mask as it was, so every probe copy of the run
repeated it on the same inputs. Its answer (shadow, ring centre, mean projection) is now kept in a
memo shared with the copies, keyed on the sampled frames, the geometry the rings are drawn in, the
polarization and the binary mask; up to 4 entries.

Debug TIMING lines for the pre-scan's parts (read and projection, beam stop, defective pixels, spot
measurement). New tests: two accumulators/finders folded match one fed every frame.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3

tests/HotPixelFinderTest.cpp

@@ -4,7 +4,9 @@
 #include <catch2/catch_all.hpp>
 
 #include <cmath>
+#include <algorithm>
 #include <random>
+#include <tuple>
 #include <vector>
 
 #include "../common/CUDAWrapper.h"
@@ -161,4 +163,75 @@ TEST_CASE("HotPixelFinder_DeviceMatchesHost", "[HotPixelFinder]") {
     CHECK(result.error == expected.error);
     CHECK(result.mask == expected.mask);
 }
+
+// The pre-scan keeps one device half per card and folds the cards together when the mask is read. Two
+// halves, each given part of the frames, folded into one must give exactly the chance counts and the
+// candidates of one half given every frame.
+TEST_CASE("HotPixelFinderGPU_FoldedCardsMatchOneCard", "[HotPixelFinder]") {
+    if (get_gpu_count() == 0) {
+        WARN("No CUDA GPU present. Skipping HotPixelFinderGPU_FoldedCardsMatchOneCard");
+        return;
+    }
+    // Four rings of two sectors: a key by band of rows and half of the detector, one column masked.
+    constexpr int NRINGS = 4, SECTORS = 2;
+    const size_t npixels = static_cast<size_t>(W) * H;
+    std::vector<int32_t> key(npixels);
+    std::vector<uint32_t> key_begin(NRINGS * SECTORS + 1, 0);
+    for (int y = 0; y < H; y++)
+        for (int x_ = 0; x_ < W; x_++) {
+            const int k = x_ == 7 ? -1 : (y * NRINGS / H) * SECTORS + x_ * SECTORS / W;
+            key[I(x_, y)] = k;
+            if (k >= 0) key_begin[k + 1]++;
+        }
+    for (int k = 0; k < NRINGS * SECTORS; k++)
+        key_begin[k + 1] += key_begin[k];
+    const HotPixelLevelRules rules{HotPixelFinder::MIN_SECTOR_PIXELS, HotPixelFinder::MIN_RING_PIXELS,
+                                   HotPixelFinder::LIT_NSIGMA, HotPixelFinder::LIT_OFFSET};
+    HotPixelFinderGPU one(key.data(), npixels, key_begin, NRINGS, SECTORS, rules);
+    HotPixelFinderGPU first(key.data(), npixels, key_begin, NRINGS, SECTORS, rules);
+    HotPixelFinderGPU second(key.data(), npixels, key_begin, NRINGS, SECTORS, rules);
+    HotPixelFinderGPU::Frame frame(std::make_shared<CudaStream>());
+    CudaDevicePtr<int32_t> device_image(npixels);
+
+    std::mt19937 rng(11);
+    std::uniform_real_distribution<double> u(0.0, 1.0);
+    std::vector<int32_t> image(npixels);
+    for (int f = 0; f < NFRAMES; f++) {
+        for (size_t i = 0; i < npixels; i++) {
+            const double d = u(rng);
+            image[i] = d < 0.01 ? INT32_MIN : d < 0.015 ? INT32_MAX : std::poisson_distribution<int>(8.0)(rng);
+        }
+        for (int p = 0; p < 40; p++)
+            image[I(20 + p, 30 + p)] += f % 3 == 0 ? 0 : 40;   // lit on most frames
+        REQUIRE(cudaMemcpy(device_image, image.data(), npixels * sizeof(int32_t), cudaMemcpyHostToDevice)
+                == cudaSuccess);
+        one.Add(device_image, frame);
+        (f % 3 == 0 ? second : first).Add(device_image, frame);
+        REQUIRE(cudaStreamSynchronize(*frame.stream) == cudaSuccess);
+    }
+    first.Fold(second);
+
+    std::vector<int64_t> lit_one, seen_one, lit_folded, seen_folded;
+    one.ChanceCounts(lit_one, seen_one);
+    first.ChanceCounts(lit_folded, seen_folded);
+    CHECK(lit_folded == lit_one);
+    CHECK(seen_folded == seen_one);
+    const std::vector<int> k_chance(NRINGS, 5);
+    const auto expected = one.GetCandidates(NFRAMES, 10, true, k_chance);
+    const auto result = first.GetCandidates(NFRAMES, 10, true, k_chance);
+    CHECK(expected.index.size() >= 40);
+    // The candidates are written by atomic slot, so their order is the device's; compare them by pixel.
+    const auto by_pixel = [](const HotPixelFinderGPU::Candidates &c) {
+        std::vector<std::tuple<uint32_t, uint16_t, uint16_t, uint16_t, int64_t, int64_t>> v;
+        for (size_t j = 0; j < c.index.size(); j++)
+            v.emplace_back(c.index[j], c.n_lit[j], c.n_error[j], c.n_error_ring_ok[j], c.sum_value[j],
+                           c.error_level_sum[j]);
+        std::sort(v.begin(), v.end());
+        return v;
+    };
+    CHECK(by_pixel(result) == by_pixel(expected));
+    CHECK(result.key_frames == expected.key_frames);
+    CHECK(result.key_level_sum == expected.key_level_sum);
+    CHECK(result.key_local_level_sum == expected.key_local_level_sum);
+}
 #endif