From 672e182d6a3f06eb37e226cdf411313436b93e71 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Sat, 10 Oct 2026 10:59:24 +0200 Subject: [PATCH] GPU engines wait for their stream before their buffers go; a lost context fails where it is seen A pooled CudaDevicePtr frees on the thread's allocation stream, not on the engine's stream, and the pool may hand the memory to another engine - or, past its release threshold, unmap it - as soon as that free is reached, which on an idle allocation stream is at once. An engine destroyed with work still queued (FFTIndexerGPU after SearchCap's last DirectionsChanged upload, a spot finder between DetectAt and Extract, a shadow accumulator after a pending fold, any engine on an exception path) thus had kernels or copies writing memory that was someone else's or no longer mapped. Now: - CudaStream synchronises before cudaStreamDestroy (destructor and move-assignment), which covers engines whose own stream is declared after their buffers (FFTIndexerGPU, the gather buffer); - every engine holding pooled buffers and a stream (shared or own, declared before the buffers) synchronises it in its destructor; BraggIntegrationEngineGPU also before EnsureCapacity reallocates, where a Run that threw leaves work queued. A GPU failure that is handled no longer hides a lost context: ShadowFinder, BeamCenterFFT, the rigid-body pool and model validation call cuda_throw_if_context_lost() before cuda_clear_error(), as the device-decode fallbacks already did; RotationScaleMergeGPU's Alloc does so before waiting up to ten minutes for GPU work beside it and then reporting a lost device as out of memory; and a failed cudaMalloc says why. BeamCenterFFT logs the failure it used to drop silently, the speculative geometry probe logs the exception it swallowed (its GPU fault was otherwise reported by the merge beside it, under the merge's name), and RotationScaleMergeGPU's DeviceGuard no longer throws from its destructor. Only synchronisation and error paths change: p.hkl md5 and the MTZ data (gemmi) are identical to the b530c2d full battery on myob_x10sa, cytc_x10sa, 8a1a, 9gdj, 11if, kdp_x10sa_20keV and 6z9g. Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_01SVmAWnzCmRKAXVUCdc4iNi --- image_analysis/azint/AzIntEngineGPU.h | 2 ++ .../beam_stop/ShadowAccumulatorGPU.h | 2 ++ image_analysis/beam_stop/ShadowFinder.cpp | 1 + .../BraggIntegrationEngineGPU.cu | 2 ++ .../BraggIntegrationEngineGPU.h | 2 ++ .../bragg_prediction/BraggPredictionGPU.h | 2 ++ .../bragg_prediction/BraggPredictionRotGPU.h | 2 ++ .../BeamCenterBackgroundGPU.cu | 3 +++ .../geom_refinement/BeamCenterFFT.cpp | 7 +++-- .../image_preprocessing/BSLZ4DecoderGPU.h | 2 ++ .../ImagePreprocessorGPU.cu | 5 +++- image_analysis/indexing/CUDAMemHelpers.h | 27 ++++++++++++++++--- image_analysis/roi/ROIIntegrationGPU.h | 2 ++ .../scale_merge/RotationScaleMergeGPU.cu | 3 ++- .../spot_finding/AdaptiveSpotFinderGPU.h | 3 ++- .../spot_finding/ImageSpotFinderGPU.h | 3 ++- .../spot_finding/SpotExtractorGPU.h | 2 ++ .../structure_refinement/ModelValidation.cpp | 3 ++- .../structure_refinement/RigidBodyGPU.cpp | 3 ++- rugnux/Rugnux.cpp | 5 +++- 20 files changed, 68 insertions(+), 13 deletions(-) diff --git a/image_analysis/azint/AzIntEngineGPU.h b/image_analysis/azint/AzIntEngineGPU.h index 480022c94..7a8bf85a0 100644 --- a/image_analysis/azint/AzIntEngineGPU.h +++ b/image_analysis/azint/AzIntEngineGPU.h @@ -30,5 +30,7 @@ class AzIntEngineGPU : public AzIntEngine { CudaRegisteredVector cpu_count_reg; public: AzIntEngineGPU(const AzimuthalIntegrationMapping& integration, std::shared_ptr stream); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~AzIntEngineGPU() override { if (stream) cudaStreamSynchronize(*stream); } void Run(const ImagePreprocessorBuffer &image, AzimuthalIntegrationProfile &profile) override; }; diff --git a/image_analysis/beam_stop/ShadowAccumulatorGPU.h b/image_analysis/beam_stop/ShadowAccumulatorGPU.h index 71907a0be..4abeaf14f 100644 --- a/image_analysis/beam_stop/ShadowAccumulatorGPU.h +++ b/image_analysis/beam_stop/ShadowAccumulatorGPU.h @@ -50,6 +50,8 @@ class ShadowAccumulatorGPU { public: explicit ShadowAccumulatorGPU(size_t npixels); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~ShadowAccumulatorGPU() { if (stream) cudaStreamSynchronize(*stream); } // True when this image can be decoded and accumulated on the device. static bool Supports(const CompressedImage &image); diff --git a/image_analysis/beam_stop/ShadowFinder.cpp b/image_analysis/beam_stop/ShadowFinder.cpp index 879d3e6ec..d1e0170c1 100644 --- a/image_analysis/beam_stop/ShadowFinder.cpp +++ b/image_analysis/beam_stop/ShadowFinder.cpp @@ -496,6 +496,7 @@ void ShadowFinder::AddImage(const DataMessage &data, std::vector &buffe return; } catch (const std::exception &e) { spdlog::warn("Beam stop: GPU accumulate failed ({}), falling back to the host", e.what()); + cuda_throw_if_context_lost(); cuda_clear_error(); // handled - see cuda_clear_error() } } diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu index 62ecc50b5..08215325b 100644 --- a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu +++ b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu @@ -897,6 +897,8 @@ void BraggIntegrationEngineGPU::EnsureCapacity(size_t n) { // it. Twice the record: a few megabytes per worker, and a worker rarely grows twice. const size_t new_capacity = std::max(2 * n, 2 * capacity); + // A Run that threw part way left work queued on the stream; its buffers are about to be freed. + cuda_err(cudaStreamSynchronize(*stream)); d_px_x = CudaDevicePtr(new_capacity); d_px_y = CudaDevicePtr(new_capacity); d_d = CudaDevicePtr(new_capacity); diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.h b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.h index a86f968f2..3e0cdd2f3 100644 --- a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.h +++ b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.h @@ -87,6 +87,8 @@ class BraggIntegrationEngineGPU : public BraggIntegrationEngine { public: BraggIntegrationEngineGPU(const DiffractionExperiment &experiment, std::shared_ptr stream, const PixelMask &mask); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~BraggIntegrationEngineGPU() override { if (stream) cudaStreamSynchronize(*stream); } std::vector Run(const ImagePreprocessorBuffer &image, const std::vector &predicted, size_t npredicted, int64_t image_number) override; diff --git a/image_analysis/bragg_prediction/BraggPredictionGPU.h b/image_analysis/bragg_prediction/BraggPredictionGPU.h index 8621fc782..9fdc392d6 100644 --- a/image_analysis/bragg_prediction/BraggPredictionGPU.h +++ b/image_analysis/bragg_prediction/BraggPredictionGPU.h @@ -46,6 +46,8 @@ class BraggPredictionGPU : public BraggPrediction { public: explicit BraggPredictionGPU(int max_reflections = kPredictionCapacity); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~BraggPredictionGPU() override { cudaStreamSynchronize(stream); } protected: void GrowCapacity(int count) override; diff --git a/image_analysis/bragg_prediction/BraggPredictionRotGPU.h b/image_analysis/bragg_prediction/BraggPredictionRotGPU.h index 0ef47eb35..fc251785f 100644 --- a/image_analysis/bragg_prediction/BraggPredictionRotGPU.h +++ b/image_analysis/bragg_prediction/BraggPredictionRotGPU.h @@ -52,6 +52,8 @@ class BraggPredictionRotGPU : public BraggPrediction { public: explicit BraggPredictionRotGPU(int max_reflections = kPredictionCapacity); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~BraggPredictionRotGPU() override { cudaStreamSynchronize(stream); } protected: void GrowCapacity(int count) override; diff --git a/image_analysis/geom_refinement/BeamCenterBackgroundGPU.cu b/image_analysis/geom_refinement/BeamCenterBackgroundGPU.cu index 550593f1d..6dfb87329 100644 --- a/image_analysis/geom_refinement/BeamCenterBackgroundGPU.cu +++ b/image_analysis/geom_refinement/BeamCenterBackgroundGPU.cu @@ -121,6 +121,9 @@ struct BeamCenterBackgroundGPU::Impl { CudaDevicePtr sort_scratch; size_t sort_scratch_bytes = 0; + // Waits for the stream, which is destroyed after the buffers freed before it (see CudaDevicePtr). + ~Impl() { cudaStreamSynchronize(stream); } + Impl(int w, int h) : width(w), height(h), npixels(static_cast(w) * h), usable(npixels), mean(npixels), row_block(h), diff --git a/image_analysis/geom_refinement/BeamCenterFFT.cpp b/image_analysis/geom_refinement/BeamCenterFFT.cpp index 7df072664..ff8726b66 100644 --- a/image_analysis/geom_refinement/BeamCenterFFT.cpp +++ b/image_analysis/geom_refinement/BeamCenterFFT.cpp @@ -11,6 +11,7 @@ #include "BeamCenterFFTCPU.h" #include "BeamCenterFFTEngine.h" #include "../../common/CUDAWrapper.h" +#include #ifdef JFJOCH_USE_CUDA #include "BeamCenterFFTGPU.h" #include "../../common/JFJochException.h" @@ -253,10 +254,12 @@ BeamCenterFFTResult BeamCenterFFTScore(int64_t width, int64_t height, BeamCenterFFTGPU gpu; try { return BeamCenterFFTScore(width, height, mean, settings, gpu); - } catch (const JFJochException &) { + } catch (const JFJochException &e) { // The card is shared with this run's analysis workers, so the memory the check above // saw can be gone by the time it is asked for. A capture is not worth failing a run - // over: score it on the CPU instead. + // over: score it on the CPU instead - unless the device is gone. + spdlog::warn("Beam centre FFT: GPU scoring failed ({}), scoring on the CPU", e.what()); + cuda_throw_if_context_lost(); cuda_clear_error(); // handled - see cuda_clear_error() } } diff --git a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h index acfac8b0f..38ad7a74a 100644 --- a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h +++ b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h @@ -87,6 +87,8 @@ class BSLZ4DecoderGPU { public: BSLZ4DecoderGPU(size_t max_uncompressed_bytes, std::shared_ptr stream); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~BSLZ4DecoderGPU() { if (stream) cudaStreamSynchronize(*stream); } // True when this image can be decoded on the device. Everything else must go the host route. static bool Supports(const CompressedImage &image); diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu index a40ff94bb..1610a8658 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu @@ -340,7 +340,10 @@ float ImagePreprocessorGPU::GetLastDecompressionTime_s() const { ImagePreprocessorGPU::~ImagePreprocessorGPU() { // Unchecked, as everywhere else in this teardown path: a destructor is noexcept and the region - // may already be gone with the device. + // may already be gone with the device. The stream is waited for first: its pooled buffers are + // freed next (see CudaDevicePtr), and the region may still be the source of a copy. + if (stream) + cudaStreamSynchronize(*stream); if (pinned_input) cudaHostUnregister(const_cast(pinned_input)); } diff --git a/image_analysis/indexing/CUDAMemHelpers.h b/image_analysis/indexing/CUDAMemHelpers.h index 72961d765..27e167491 100644 --- a/image_analysis/indexing/CUDAMemHelpers.h +++ b/image_analysis/indexing/CUDAMemHelpers.h @@ -23,14 +23,24 @@ public: throw JFJochException(JFJochExceptionCategory::GPUCUDAError, "Failed to create CUDA stream"); } + // The stream's queued work is waited for before it goes. cudaStreamDestroy alone returns at once and + // lets that work run on, while the owner's buffers are freed next: a pooled buffer's free is ordered + // on the allocation stream (CudaDevicePtr), not on this one, so its memory could be handed to + // another engine, or returned to the driver, under a kernel or copy still writing it. ~CudaStream() { - if (stream_) cudaStreamDestroy(stream_); + if (stream_) { + cudaStreamSynchronize(stream_); + cudaStreamDestroy(stream_); + } } // Move-only type CudaStream(CudaStream&& other) noexcept : stream_(other.stream_) { other.stream_ = nullptr; } CudaStream& operator=(CudaStream&& other) noexcept { if (this != &other) { - if (stream_) cudaStreamDestroy(stream_); + if (stream_) { + cudaStreamSynchronize(stream_); + cudaStreamDestroy(stream_); + } stream_ = other.stream_; other.stream_ = nullptr; } @@ -244,6 +254,13 @@ inline bool cuda_memory_available(size_t bytes, size_t &free_bytes) { return read_free(); } +// A pooled buffer's free is ordered on the allocation stream, not on the stream of the engine that +// used it, and the pool may hand the memory to another engine - or return it to the driver - as soon +// as that free is reached, which on an idle allocation stream is at once. So an engine with pooled +// buffers waits for its own stream before they go: in its destructor (a CudaStream it owns does that +// itself when it is destroyed before the buffers, i.e. declared after them) and before it reallocates. +// Otherwise a kernel or copy still queued writes memory that is someone else's, or unmapped. +// // 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 @@ -293,9 +310,11 @@ public: // stay behind as the last error (see above). cudaGetLastError(); cuda_trim_allocation_pool(); - if (cudaMalloc(&ptr_, count * sizeof(T)) != cudaSuccess) + // 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) throw JFJochException(JFJochExceptionCategory::GPUCUDAError, - "Failed to allocate device memory"); + std::string("Failed to allocate device memory (") + cudaGetErrorString(err) + ")"); } ~CudaDevicePtr() { release(); diff --git a/image_analysis/roi/ROIIntegrationGPU.h b/image_analysis/roi/ROIIntegrationGPU.h index 5528b37bd..1e0696ce8 100644 --- a/image_analysis/roi/ROIIntegrationGPU.h +++ b/image_analysis/roi/ROIIntegrationGPU.h @@ -33,5 +33,7 @@ class ROIIntegrationGPU : public ROIIntegration { std::vector max_init; // INT_MIN seed copied into gpu_max each frame public: ROIIntegrationGPU(const DiffractionExperiment &experiment, std::shared_ptr stream); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~ROIIntegrationGPU() override { if (stream) cudaStreamSynchronize(*stream); } void Run(const ImagePreprocessorBuffer &image, std::map &out) override; }; diff --git a/image_analysis/scale_merge/RotationScaleMergeGPU.cu b/image_analysis/scale_merge/RotationScaleMergeGPU.cu index 531492632..b6aa638a5 100644 --- a/image_analysis/scale_merge/RotationScaleMergeGPU.cu +++ b/image_analysis/scale_merge/RotationScaleMergeGPU.cu @@ -1170,6 +1170,7 @@ struct RotationScaleMergeGPU::Impl { try { return CudaDevicePtr(n, ALLOC); } catch (const JFJochException &) { + cuda_throw_if_context_lost(); // not memory: waiting would only hide the real error cuda_clear_error(); } if (!wait_for_gpu_work_beside(GPU_BUSY_TIMEOUT)) @@ -1282,7 +1283,7 @@ struct DeviceGuard { } ~DeviceGuard() { if (active) - set_gpu(prev); + cudaSetDevice(prev); // not set_gpu: a destructor must not throw } }; } // namespace diff --git a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h index fec9840b6..e6ae453cc 100644 --- a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h +++ b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h @@ -122,7 +122,8 @@ public: static constexpr int32_t HIST_VALUES = 1024; // as AdaptiveSpotFinderCPU AdaptiveSpotFinderGPU(const AzimuthalIntegrationMapping &mapping, std::shared_ptr stream); - ~AdaptiveSpotFinderGPU() override = default; + // Waits for the stream: its pooled buffers are freed before the base class's (see CudaDevicePtr). + ~AdaptiveSpotFinderGPU() override { if (stream) cudaStreamSynchronize(*stream); } AdaptiveSpotFinderGPU(const AdaptiveSpotFinderGPU &) = delete; AdaptiveSpotFinderGPU &operator=(const AdaptiveSpotFinderGPU &) = delete; diff --git a/image_analysis/spot_finding/ImageSpotFinderGPU.h b/image_analysis/spot_finding/ImageSpotFinderGPU.h index c7da639ea..97c2812cf 100644 --- a/image_analysis/spot_finding/ImageSpotFinderGPU.h +++ b/image_analysis/spot_finding/ImageSpotFinderGPU.h @@ -38,7 +38,8 @@ protected: const uint32_t *gpu_candidates); public: ImageSpotFinderGPU(int32_t width, int32_t height, std::shared_ptr stream); - ~ImageSpotFinderGPU() override = default; + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~ImageSpotFinderGPU() override { if (stream) cudaStreamSynchronize(*stream); } void Detect(const ImagePreprocessorBuffer &image, const SpotFindingSettings &settings) override; void SetResolutionMaskBits(const std::vector &packed_mask) override; diff --git a/image_analysis/spot_finding/SpotExtractorGPU.h b/image_analysis/spot_finding/SpotExtractorGPU.h index 4bef4f31d..289582abf 100644 --- a/image_analysis/spot_finding/SpotExtractorGPU.h +++ b/image_analysis/spot_finding/SpotExtractorGPU.h @@ -84,6 +84,8 @@ class SpotExtractorGPU { public: SpotExtractorGPU(int32_t width, int32_t height, std::shared_ptr stream); + // Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr). + ~SpotExtractorGPU() { if (stream) cudaStreamSynchronize(*stream); } void SetResolutionMask(const std::vector &packed_mask); diff --git a/image_analysis/structure_refinement/ModelValidation.cpp b/image_analysis/structure_refinement/ModelValidation.cpp index ab8797727..4c5bfb1e6 100644 --- a/image_analysis/structure_refinement/ModelValidation.cpp +++ b/image_analysis/structure_refinement/ModelValidation.cpp @@ -1722,11 +1722,12 @@ ModelValidationResult ValidateAgainstModel(const std::vector & probe_indexing_ambiguity, nthreads, wavelength_A, report_shell_d_min, schedule, twin_law, true); } catch (const RigidBodyGPUFailure &e) { - cuda_clear_error(); ModelValidationResult failed; failed.model_path = model_path; failed.failure_reason = fmt::format("the GPU failed during the validation ({})", e.what()); logger.Error("Model validation: {}", failed.failure_reason); + cuda_throw_if_context_lost(); // a lost device ends the run here, not at its next user + cuda_clear_error(); return failed; } #else diff --git a/image_analysis/structure_refinement/RigidBodyGPU.cpp b/image_analysis/structure_refinement/RigidBodyGPU.cpp index 9ba4c2d7e..ea6b855a7 100644 --- a/image_analysis/structure_refinement/RigidBodyGPU.cpp +++ b/image_analysis/structure_refinement/RigidBodyGPU.cpp @@ -184,9 +184,10 @@ std::unique_ptr RigidBodyGPUPool::Create(const gemmi::Model &m bytes / 1e6); return pool; } catch (const JFJochException &e) { - cuda_clear_error(); logger.Warning("Model validation: the rigid body runs on the CPU - the GPU engines could not be set up ({})", e.what()); + cuda_throw_if_context_lost(); + cuda_clear_error(); return nullptr; } } diff --git a/rugnux/Rugnux.cpp b/rugnux/Rugnux.cpp index 659039ccc..05667ee9e 100644 --- a/rugnux/Rugnux.cpp +++ b/rugnux/Rugnux.cpp @@ -422,7 +422,10 @@ void Rugnux::StartSpeculativeGeometryProbe(std::shared_ptr run, const st for (const auto &r : {run, hand}) { try { r->RunPipeline(nullptr, /*write_output=*/false, /*geometry_prepass=*/false); - } catch (const std::exception &) { + } catch (const std::exception &e) { + // Said, not swallowed: a GPU fault here is reported next by whichever user of the card + // touches it first - the merge beside it - under that user's name. + Logger("Rugnux").Warning("Speculative geometry probe failed: {}", e.what()); } } });