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()); } } });