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
37 lines
1.6 KiB
C++
37 lines
1.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 "AzIntEngine.h"
|
|
#include "../indexing/CUDAMemHelpers.h"
|
|
#include "../indexing/CudaSharedTables.h"
|
|
|
|
class AzIntEngineGPU : public AzIntEngine {
|
|
std::shared_ptr<CudaStream> stream;
|
|
int threads;
|
|
int blocks;
|
|
size_t shared_needed;
|
|
size_t shared_size;
|
|
|
|
// Geometry-only tables: one copy per GPU, shared with every other engine on it (see
|
|
// CudaSharedTables.h) instead of one copy per worker thread.
|
|
std::shared_ptr<CudaDevicePtr<float>> gpu_azint_correction;
|
|
std::shared_ptr<CudaDevicePtr<uint16_t>> gpu_pixel_to_bin;
|
|
|
|
// Fixed point (AZINT_FIXED in the .cu), so the sums do not depend on the order the atomics arrive in.
|
|
CudaDevicePtr<unsigned long long> gpu_sum;
|
|
CudaDevicePtr<unsigned long long> gpu_sum2;
|
|
CudaDevicePtr<uint32_t> gpu_count;
|
|
std::vector<unsigned long long> fixed_sum;
|
|
std::vector<unsigned long long> fixed_sum2;
|
|
CudaRegisteredVector<unsigned long long> cpu_sum_reg;
|
|
CudaRegisteredVector<unsigned long long> cpu_sum2_reg;
|
|
CudaRegisteredVector<uint32_t> cpu_count_reg;
|
|
public:
|
|
AzIntEngineGPU(const AzimuthalIntegrationMapping& integration, std::shared_ptr<CudaStream> 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;
|
|
};
|