Adaptive spot finder: pin the threshold to the image, and the GPU to itself

The existing cases plant blobs at 200 on a background of 8..12, so any threshold
between 12 and 200 passes them - replacing RingThreshold with a constant leaves
them all green. Two cases that do not:

- the CPU threshold has to track the background: a frame and the same frame
  scaled ten times must give the same spots, with a pixel a few sigma above the
  background staying unfound in both. A constant threshold, or one that drops
  the sigma term, fails one scale or the other.
- the GPU engine has to agree with itself across runs, which is what the ring
  sums being order-independent buys.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

tests/AdaptiveSpotFinderGPUTest.cpp

@@ -71,8 +71,8 @@ std::vector<std::pair<int, int>> SortedCoords(const std::vector<DiffractionSpot>
 } // namespace
 
 // Spot-finding functionality: the fused GPU engine must reproduce the reference CPU adaptive finder's
-// spot list (the two share AdaptiveThreshold.h and the host connected-component extractor; the only
-// difference is the GPU's float atomic ring reduction, which is exact for a realistic background).
+// spot list. The two share AdaptiveThreshold.h and the host connected-component extractor, and both
+// sum the rings in double, so the only difference left is the order the ring sums are accumulated in.
 TEST_CASE("AdaptiveSpotFinderGPU_SpotFindingParity", "[AdaptiveSpotFinderGPU]") {
     if (get_gpu_count() == 0) {
         WARN("No CUDA GPU present. Skipping AdaptiveSpotFinderGPU_SpotFindingParity");
@@ -147,6 +147,40 @@ TEST_CASE("AdaptiveSpotFinderGPU_AzimuthalIntegration", "[AdaptiveSpotFinderGPU]
     }
 }
 
+// The ring sums are built by atomics, which arrive in an arbitrary order, so the same frame has to be
+// re-run to show the engine agrees with itself: detection is a hard "value >= threshold" on integer
+// counts, and a threshold that wobbles between runs flips pixels on the boundary and with them the size
+// of a connected component. Two runs, same spot list.
+TEST_CASE("AdaptiveSpotFinderGPU_RunToRunReproducible", "[AdaptiveSpotFinderGPU]") {
+    if (get_gpu_count() == 0) {
+        WARN("No CUDA GPU present. Skipping AdaptiveSpotFinderGPU_RunToRunReproducible");
+        return;
+    }
+
+    DiffractionExperiment x = MakeExperiment();
+    PixelMask pixel_mask(x);
+    AzimuthalIntegrationMapping mapping(x, pixel_mask);
+
+    ImagePreprocessorBufferGPU buffer(x.GetPixelsNum());
+    FillTestImage(buffer, x);
+    REQUIRE(cudaMemcpy(buffer.getGPUBuffer(), buffer.getBuffer().data(),
+                       x.GetPixelsNum() * sizeof(int32_t), cudaMemcpyHostToDevice) == cudaSuccess);
+
+    std::vector<bool> res_mask(x.GetPixelsNum(), false);
+    const SpotFindingSettings settings = AdaptiveSettings();
+
+    auto stream = std::make_shared<CudaStream>();
+    AdaptiveSpotFinderGPU gpu(mapping, stream);
+
+    const auto first = gpu.Run(buffer, settings, res_mask);
+    REQUIRE(first.size() > 0);
+    for (int repeat = 0; repeat < 4; repeat++) {
+        const auto again = gpu.Run(buffer, settings, res_mask);
+        REQUIRE(again.size() == first.size());
+        REQUIRE(SortedCoords(again) == SortedCoords(first));
+    }
+}
+
 TEST_CASE("AdaptiveSpotFinderGPU_Speed", "[AdaptiveSpotFinderGPU][.benchmark]") {
     if (get_gpu_count() == 0) {
         WARN("No CUDA GPU present. Skipping AdaptiveSpotFinderGPU_Speed");