leonarski_fandClaude Opus 5.5
672e182d6a
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) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SVmAWnzCmRKAXVUCdc4iNi
52 lines
2.6 KiB
C++
52 lines
2.6 KiB
C++
// SPDX-FileCopyrightText: 2025 Filip Leonarski, Paul Scherrer Institute <filip.leonarski@psi.ch>
|
|
// SPDX-License-Identifier: GPL-3.0-only
|
|
|
|
#pragma once
|
|
|
|
#include <vector>
|
|
|
|
#include "SpotFindingSettings.h"
|
|
#include "ImageSpotFinder.h"
|
|
#include "SpotExtractorGPU.h"
|
|
#include "../indexing/CUDAMemHelpers.h"
|
|
|
|
class ImageSpotFinderGPU : public ImageSpotFinder {
|
|
protected:
|
|
// Protected rather than private because AdaptiveSpotFinderGPU derives from this engine: it is
|
|
// this same local-box detection with the fixed photon floor replaced by a per-resolution-ring
|
|
// one, so it reuses the stream, the bit buffers and the extractor rather than owning a second
|
|
// set of them.
|
|
std::shared_ptr<CudaStream> stream;
|
|
|
|
CudaDevicePtr<uint32_t> gpu_out_0;
|
|
CudaDevicePtr<uint32_t> gpu_out_1; // holds the strong-pixel bits after Detect()
|
|
SpotExtractorGPU extractor;
|
|
|
|
private:
|
|
const int numberOfCudaThreads = 128; // #threads per block of analyze_pixel (one warp per 32 columns)
|
|
const int numberOfWaves = 32; // #row bands of analyze_pixel
|
|
const int windowSizeLimit = 32; // limit on the window size (2nby+1, 2nbx+1): a warp holds 2 x 32 columns
|
|
int candidateBlocks = 0; // grid of analyze_candidates (256 threads per block)
|
|
|
|
void RunDetect(const ImagePreprocessorBuffer &image, const SpotFindingSettings &settings,
|
|
const uint32_t *gpu_candidates);
|
|
protected:
|
|
// Detect, but with the result wanted only at the candidate pixels (device bit buffer, the layout
|
|
// of the output): gpu_out_1 ends up holding (Detect's result & candidates), and the second pass is
|
|
// evaluated at the candidates alone. Asynchronous - it does not wait for the stream.
|
|
void DetectAt(const ImagePreprocessorBuffer &image, const SpotFindingSettings &settings,
|
|
const uint32_t *gpu_candidates);
|
|
public:
|
|
ImageSpotFinderGPU(int32_t width, int32_t height, std::shared_ptr<CudaStream> stream);
|
|
// 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<uint32_t> &packed_mask) override;
|
|
[[nodiscard]] uint32_t StrongPixelCount() const override { return extractor.StrongPixelCount(); }
|
|
const std::vector<DiffractionSpot> &ExtractComponents(const ImagePreprocessorBuffer &image,
|
|
const SpotFindingSettings &settings) override;
|
|
};
|
|
|
|
|