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