CudaDevicePtr: trim the allocation pool and retry when cudaMalloc fails

ca1baffef moved RotationScaleMergeGPU to synchronous cudaMalloc. cudaMalloc
cannot take memory the default pool holds, and after the image loop the pool
still reserves what the per-worker engines freed with cudaFreeAsync (the 1 GB
release threshold only applies at a synchronisation of the pool's streams).
Measured at the merge's construction on a 123.6M-partial rotation set: pool
reserved 5.64 GB, used 0.83 GB - 4.8 GB of the 16.3 GB card idle but
unavailable. The merge's per-observation arrays (about 9.5 GB there) then ran
out at a 494 MB array with 316 MB free, and the run failed with "Failed to
allocate device memory" right after the post-refine gather. The 77710e1ce
binary allocated those arrays from the pool and passed.

A failed synchronous allocation now waits for the device, trims the pool to
zero and tries once more. Memory placement only: no arithmetic changes.
p.hkl, p.mtz, p_P1.mtz, p_unmerged.mtz byte-identical on myob, cytc, lyso
and sparse; the failing set now passes (R 3:H, 1.30 A, ISa 10.9).

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

image_analysis/indexing/CUDAMemHelpers.h

@@ -204,6 +204,20 @@ inline cudaStream_t cuda_allocation_stream() {
     }
 }
 
+// Return the current device's pool memory that no buffer holds to the driver. Frees ordered on a
+// stream count only once the stream has reached them, hence the device-wide wait first.
+inline void cuda_trim_allocation_pool() {
+    int device = 0;
+    cudaMemPool_t pool = nullptr;
+    if (cudaGetDevice(&device) != cudaSuccess || cudaDeviceGetDefaultMemPool(&pool, device) != cudaSuccess) {
+        cudaGetLastError();
+        return;
+    }
+    cudaDeviceSynchronize();
+    cudaMemPoolTrimTo(pool, 0);
+    cudaGetLastError();
+}
+
 // Which allocator a device buffer uses. Pooled is what every per-worker buffer wants. Synchronous is
 // 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
@@ -245,6 +259,14 @@ public:
             cudaGetLastError();
             stream_ = nullptr;
         }
+        if (cudaMalloc(&ptr_, count * sizeof(T)) == cudaSuccess)
+            return;
+        // cudaMalloc cannot take memory the pool is holding: what pooled buffers handed back stays
+        // reserved by the pool, up to gigabytes after the image loop's engines are gone, until the
+        // pool is trimmed. Hand it back to the driver and try once more; the failed attempt must not
+        // stay behind as the last error (see above).
+        cudaGetLastError();
+        cuda_trim_allocation_pool();
         if (cudaMalloc(&ptr_, count * sizeof(T)) != cudaSuccess)
             throw JFJochException(JFJochExceptionCategory::GPUCUDAError,
                                   "Failed to allocate device memory");