rotation merge: helper buffers wait for the probe pass beside them, as the engine's own do

The intermittent "out of GPU memory" on the 3600-frame cbf.gz set (about one run in three on a 16 GB
card, base included) was caught with allocation backtraces: FitErrorModelGPU's 42.8 MB in
MergeAndStats failed while the speculative geometry probe (GPUWorkBeside) held its FFT indexers on
the same card. The engine's own arrays (RotationScaleMergeGPU Impl::Alloc) wait for such work to end
and ask again; the merge's helpers - ErrorModelGPU, WilsonOutliersGPU, PairSortGPU and the cub
scratch of the engine itself - allocated directly and failed at once.

CudaAlloc::SynchronousAfterWorkBeside: Synchronous, and where the card is full after the pool is
trimmed, it waits for the GPU work beside the main line (wait_for_gpu_work_beside, 10 min as for
the engine) and asks once more. The four merge files use it. Only the probe passes register
GPUWorkBeside and they index only, so none of them merges and nothing waits for itself. When a
buffer is allocated is all that changes; what is computed is not.

Catch2 CudaDevicePtr_AfterWorkBesideWaitsForItsMemory: with the card filled by work beside, a
Synchronous buffer fails and a SynchronousAfterWorkBeside one gets its memory once the work ends.

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

image_analysis/indexing/CUDAMemHelpers.h

@@ -293,7 +293,11 @@ inline bool cuda_memory_available(size_t bytes, size_t &free_bytes) {
 // for a buffer that is freed by a thread other than the one whose work read it: cudaFreeAsync orders
 // the free on the freeing side's allocation stream, which says nothing about kernels queued
 // elsewhere, while cudaFree synchronises the whole device and therefore cannot be early.
-enum class CudaAlloc { Pooled, Synchronous };
+// SynchronousAfterWorkBeside is Synchronous for a step of the main line - the merge's helpers - whose
+// buffers may be asked for while a probe pass runs beside it on the same card (GPUWorkBeside): one that
+// does not fit waits for that work to end and is asked for once more, as the merge's own arrays are
+// (RotationScaleMergeGPU). Never for the work beside itself, which would wait for itself.
+enum class CudaAlloc { Pooled, Synchronous, SynchronousAfterWorkBeside };
 
 template <typename T>
 class CudaDevicePtr {
@@ -338,9 +342,17 @@ public:
         // stay behind as the last error (see above).
         cudaGetLastError();
         cuda_trim_allocation_pool();
+        cudaError_t err = cudaMalloc(&ptr_, count * sizeof(T));
+        if (err == cudaErrorMemoryAllocation && alloc == CudaAlloc::SynchronousAfterWorkBeside) {
+            cudaGetLastError();
+            // That work is a probe pass of seconds, so this long means it is stuck.
+            wait_for_gpu_work_beside(std::chrono::minutes(10));
+            cuda_trim_allocation_pool();
+            err = cudaMalloc(&ptr_, count * sizeof(T));
+        }
         // With the reason: a context lost to another kernel's fault fails here too, and must not read
         // as a card that is merely full.
-        if (const cudaError_t err = cudaMalloc(&ptr_, count * sizeof(T)); err != cudaSuccess)
+        if (err != cudaSuccess)
             throw JFJochException(JFJochExceptionCategory::GPUCUDAError,
                                   GPUMemoryFailure("Allocating GPU memory", count * sizeof(T),
                                                    err == cudaErrorMemoryAllocation, cudaGetErrorString(err)));

image_analysis/scale_merge/ErrorModelGPU.cu

@@ -16,7 +16,7 @@ namespace {
     constexpr int THREADS = 256;
     constexpr int CHUNKS = 64;   // blocks per bin in the sum kernel; fixed, so the tree is too
     constexpr int NB = ERROR_MODEL_BINS;
-    constexpr CudaAlloc ALLOC = CudaAlloc::Synchronous;
+    constexpr CudaAlloc ALLOC = CudaAlloc::SynchronousAfterWorkBeside;
 
     void cuda_err(cudaError_t val) {
         if (val != cudaSuccess)

image_analysis/scale_merge/PairSortGPU.cu

@@ -13,7 +13,7 @@
 
 namespace {
     constexpr int THREADS = 256;
-    constexpr CudaAlloc ALLOC = CudaAlloc::Synchronous;
+    constexpr CudaAlloc ALLOC = CudaAlloc::SynchronousAfterWorkBeside;
 
     void cuda_err(cudaError_t val) {
         if (val != cudaSuccess)

image_analysis/scale_merge/RotationScaleMergeGPU.cu

@@ -36,7 +36,7 @@ namespace {
     // free while unwinding from a failed call; nor can compute-sanitizer --track-stream-ordered-races
     // see it, and it reported every reassigned merge buffer as a use-after-free. cudaFree
     // synchronises the device first.
-    constexpr CudaAlloc ALLOC = CudaAlloc::Synchronous;
+    constexpr CudaAlloc ALLOC = CudaAlloc::SynchronousAfterWorkBeside;
 
     __device__ __forceinline__ double SafeInvD(double x, double fallback) {
         return (isfinite(x) && x != 0.0) ? 1.0 / x : fallback;

image_analysis/scale_merge/WilsonOutliersGPU.cu

@@ -25,7 +25,7 @@
 namespace {
     constexpr int THREADS = 256;
     constexpr size_t OBS_PER_SHELL = 2000;   // WilsonOutliers.cpp
-    constexpr CudaAlloc ALLOC = CudaAlloc::Synchronous;
+    constexpr CudaAlloc ALLOC = CudaAlloc::SynchronousAfterWorkBeside;
 
     void cuda_err(cudaError_t val) {
         if (val != cudaSuccess)

tests/CUDAMemHelpersTest.cpp

@@ -206,6 +206,39 @@ TEST_CASE("IndexerThreadPool_FFTIndexerWaitsForGPUWorkBeside", "[CUDAMemHelpers]
     }
 }
 
+// A merge helper's buffer that finds the card full while a probe pass holds the rest of it waits for
+// that work to end and its memory to come back (CudaAlloc::SynchronousAfterWorkBeside); a plain
+// synchronous one fails at once.
+TEST_CASE("CudaDevicePtr_AfterWorkBesideWaitsForItsMemory", "[CUDAMemHelpers]") {
+    if (get_gpu_count() == 0)
+        SKIP("no GPU");
+    std::vector<CudaDevicePtr<uint8_t>> held;
+    const auto fill = [&] {
+        for (size_t chunk = size_t(256) << 20; chunk >= (size_t(1) << 20); chunk /= 2)
+            while (true) {
+                try {
+                    held.emplace_back(chunk, CudaAlloc::Synchronous);
+                } catch (const JFJochException &) {
+                    break;
+                }
+            }
+    };
+    constexpr size_t BYTES = size_t(64) << 20;
+    {
+        std::optional<GPUWorkBeside> work(std::in_place);
+        fill();
+        CHECK_THROWS_AS(CudaDevicePtr<uint8_t>(BYTES, CudaAlloc::Synchronous), JFJochException);
+        std::thread ends([&] {
+            std::this_thread::sleep_for(std::chrono::milliseconds(500));
+            held.clear();
+            work.reset();
+        });
+        CHECK_NOTHROW(CudaDevicePtr<uint8_t>(BYTES, CudaAlloc::SynchronousAfterWorkBeside));
+        ends.join();
+    }
+    cuda_trim_allocation_pool();
+}
+
 #endif
 
 // A merge that finds the card full waits for the GPU work running beside it: until the last one