Stills keep an overload as a lower bound only where it has pixels at the limit; run-level gate removed

The lower bound counts each saturated pixel at the saturation limit, so it needs one. An observation
overloaded only by pixels unreadable on this frame (a per-frame bad-pixel flag, data value 0) was kept
as a box sum with those pixels counted as zero, which bounds nothing; it now stays overloaded and is
dropped, per observation.

With the stuck pixels masked by the pre-scan, the run-level gate (kept overloads must be enriched
coarser than 4 A) has nothing left to catch: without it every still set gives the gated result
(clipped CCD and Rayonix sets identical to the fourth decimal, photon-counting sets unchanged), so it
is removed with Reflection::overload_kept. A neighbour-brightness test per saturated pixel was tried
and not kept: after the mask it only dropped 9-20% of the Rayonix overloads and worsened R_split on
one set (0.245 -> 0.278).

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

tests/BraggIntegrationEngineGPUTest.cpp

@@ -364,6 +364,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");