From b6f34db258fd46163d8d05dbbaed4e50ecdc4be0 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Tue, 13 Oct 2026 16:07:33 +0200 Subject: [PATCH] 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) Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3 --- image_analysis/indexing/CUDAMemHelpers.h | 16 +++++++-- image_analysis/scale_merge/ErrorModelGPU.cu | 2 +- image_analysis/scale_merge/PairSortGPU.cu | 2 +- .../scale_merge/RotationScaleMergeGPU.cu | 2 +- .../scale_merge/WilsonOutliersGPU.cu | 2 +- tests/CUDAMemHelpersTest.cpp | 33 +++++++++++++++++++ 6 files changed, 51 insertions(+), 6 deletions(-) diff --git a/image_analysis/indexing/CUDAMemHelpers.h b/image_analysis/indexing/CUDAMemHelpers.h index 43ae2ca5d..0da8b4719 100644 --- a/image_analysis/indexing/CUDAMemHelpers.h +++ b/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 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))); diff --git a/image_analysis/scale_merge/ErrorModelGPU.cu b/image_analysis/scale_merge/ErrorModelGPU.cu index 3443d8e09..cd8c5a4c4 100644 --- a/image_analysis/scale_merge/ErrorModelGPU.cu +++ b/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) diff --git a/image_analysis/scale_merge/PairSortGPU.cu b/image_analysis/scale_merge/PairSortGPU.cu index 693821158..50ef96811 100644 --- a/image_analysis/scale_merge/PairSortGPU.cu +++ b/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) diff --git a/image_analysis/scale_merge/RotationScaleMergeGPU.cu b/image_analysis/scale_merge/RotationScaleMergeGPU.cu index 62805d660..4315e4b35 100644 --- a/image_analysis/scale_merge/RotationScaleMergeGPU.cu +++ b/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; diff --git a/image_analysis/scale_merge/WilsonOutliersGPU.cu b/image_analysis/scale_merge/WilsonOutliersGPU.cu index aeb84cc10..b877d07a7 100644 --- a/image_analysis/scale_merge/WilsonOutliersGPU.cu +++ b/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) diff --git a/tests/CUDAMemHelpersTest.cpp b/tests/CUDAMemHelpersTest.cpp index 2d0047319..c349d9d88 100644 --- a/tests/CUDAMemHelpersTest.cpp +++ b/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> 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 work(std::in_place); + fill(); + CHECK_THROWS_AS(CudaDevicePtr(BYTES, CudaAlloc::Synchronous), JFJochException); + std::thread ends([&] { + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + held.clear(); + work.reset(); + }); + CHECK_NOTHROW(CudaDevicePtr(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