diff --git a/image_analysis/bragg_prediction/BraggPredictionGPU.cu b/image_analysis/bragg_prediction/BraggPredictionGPU.cu index 457aa33b6..9d1689f95 100644 --- a/image_analysis/bragg_prediction/BraggPredictionGPU.cu +++ b/image_analysis/bragg_prediction/BraggPredictionGPU.cu @@ -221,15 +221,14 @@ namespace { } // namespace void BraggPredictionGPU::GrowCapacity(int count) { - reg_out = CudaRegisteredVector(); BraggPrediction::GrowCapacity(count); - reg_out = CudaRegisteredVector(reflections); + h_out = CudaHostPtr(count); d_out = CudaDevicePtr(count); } BraggPredictionGPU::BraggPredictionGPU(int max_reflections) : BraggPrediction(max_reflections), - reg_out(reflections), d_out(max_reflections), + h_out(max_reflections), d_out(max_reflections), dK(1), d_count(1), h_count(1) { } @@ -269,8 +268,9 @@ int BraggPredictionGPU::Calc(const DiffractionExperiment &experiment, if (count == 0) return {}; - cudaMemcpyAsync(reflections.data(), d_out, sizeof(Reflection) * count, cudaMemcpyDeviceToHost, stream); + cudaMemcpyAsync(h_out, d_out, sizeof(Reflection) * count, cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); + std::copy(h_out.get(), h_out.get() + count, reflections.begin()); OrderOutput(count); return TruncateToOutput(count); diff --git a/image_analysis/bragg_prediction/BraggPredictionGPU.h b/image_analysis/bragg_prediction/BraggPredictionGPU.h index 9fdc392d6..7f38ba8f6 100644 --- a/image_analysis/bragg_prediction/BraggPredictionGPU.h +++ b/image_analysis/bragg_prediction/BraggPredictionGPU.h @@ -31,7 +31,9 @@ struct KernelConsts { }; class BraggPredictionGPU : public BraggPrediction { - CudaRegisteredVector reg_out; // requires stable storage + // The device's reflections land here and are copied into `reflections`: pinned, so the download is a + // DMA, and from the pool, so the engine does not register and unregister a vector of its own. + CudaHostPtr h_out; CudaDevicePtr d_out; // Dedicated stream diff --git a/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu b/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu index 07ce14209..83dc1717d 100644 --- a/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu +++ b/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu @@ -334,15 +334,14 @@ namespace { BraggPredictionRotGPU::BraggPredictionRotGPU(int max_reflections) : BraggPrediction(max_reflections), - reg_out(reflections), d_out(max_reflections), + h_out(max_reflections), d_out(max_reflections), dK(1), d_count(1), h_count(1) { } // The host buffer is page-locked and the device buffer is sized to match it, so both are rebuilt. void BraggPredictionRotGPU::GrowCapacity(int count) { - reg_out = CudaRegisteredVector(); // unregister before the vector reallocates BraggPrediction::GrowCapacity(count); - reg_out = CudaRegisteredVector(reflections); + h_out = CudaHostPtr(count); d_out = CudaDevicePtr(count); } @@ -382,8 +381,9 @@ int BraggPredictionRotGPU::Calc(const DiffractionExperiment &experiment, } if (count == 0) return 0; - cudaMemcpyAsync(reflections.data(), d_out, sizeof(Reflection) * count, cudaMemcpyDeviceToHost, stream); + cudaMemcpyAsync(h_out, d_out, sizeof(Reflection) * count, cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); + std::copy(h_out.get(), h_out.get() + count, reflections.begin()); OrderOutput(count); return TruncateToOutput(count); diff --git a/image_analysis/bragg_prediction/BraggPredictionRotGPU.h b/image_analysis/bragg_prediction/BraggPredictionRotGPU.h index fc251785f..d56dd9c86 100644 --- a/image_analysis/bragg_prediction/BraggPredictionRotGPU.h +++ b/image_analysis/bragg_prediction/BraggPredictionRotGPU.h @@ -37,7 +37,9 @@ struct KernelConstsRot { }; class BraggPredictionRotGPU : public BraggPrediction { - CudaRegisteredVector reg_out; // requires stable storage + // The device's reflections land here and are copied into `reflections`: pinned, so the download is a + // DMA, and from the pool, so the engine does not register and unregister a vector of its own. + CudaHostPtr h_out; CudaDevicePtr d_out; // Dedicated stream diff --git a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu index df84b0fca..cbe760fcd 100644 --- a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu +++ b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu @@ -2,6 +2,7 @@ // SPDX-License-Identifier: GPL-3.0-only #include +#include #include #include "BSLZ4DecoderGPU.h" @@ -146,6 +147,7 @@ void BSLZ4DecoderGPU::EnsureCompressedCapacity(size_t bytes) { const size_t want = bytes + bytes / 4; cuda_err(cudaStreamSynchronize(*stream)); // nothing may still be reading the old buffer gpu_compressed = CudaDevicePtr(want); + host_compressed = CudaHostPtr(want); compressed_capacity = want; } @@ -238,7 +240,8 @@ BSLZ4ShuffledImage BSLZ4DecoderGPU::PrepareChunk(const CompressedImage &image) { host_status.get()[0] = 0; cuda_err(cudaMemcpyAsync(gpu_status.get(), host_status.get(), sizeof(uint32_t), cudaMemcpyHostToDevice, *stream)); - cuda_err(cudaMemcpyAsync(gpu_compressed.get(), src, clen, cudaMemcpyHostToDevice, *stream)); + memcpy(host_compressed.get(), src, clen); + cuda_err(cudaMemcpyAsync(gpu_compressed.get(), host_compressed.get(), clen, cudaMemcpyHostToDevice, *stream)); const int nb = static_cast(nblk); if (nb > 0) diff --git a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h index 38ad7a74a..f6c306a98 100644 --- a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h +++ b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.h @@ -62,6 +62,10 @@ class BSLZ4DecoderGPU { std::shared_ptr stream; CudaDevicePtr gpu_compressed; + // The chunk is copied here before it is uploaded. The reader's buffer is pageable, and the driver + // uploads pageable memory through small staging buffers of its own, a piece at a time - at a few + // GB/s on a host whose memory is far from the card. Copied here, the upload is one DMA at the link rate. + CudaHostPtr host_compressed; CudaDevicePtr gpu_shuffled; // LZ4 output, still bitshuffled CudaDevicePtr gpu_desc; CudaHostPtr host_desc; // pinned, so the descriptor upload is truly async diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu index 1610a8658..b6efa5e3f 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu @@ -314,8 +314,7 @@ ImagePreprocessorGPU::ImagePreprocessorGPU(const DiffractionExperiment &experime stream(stream), copy_image_to_host(copy_image_to_host), gpu_stats(1), - cpu_stats(1), - cpu_stats_reg(cpu_stats) { + cpu_stats(1) { // Setup mask. The byte-per-pixel form and its checksum come from the PixelMask, which derives them // whenever the mask changes: they are the same for every worker, and deriving them walks every // pixel of the detector - 18 million of them on a 16 Mpx one, per engine, with an engine built per diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.h b/image_analysis/image_preprocessing/ImagePreprocessorGPU.h index c28f4a48c..9a056de92 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.h +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.h @@ -23,8 +23,7 @@ class ImagePreprocessorGPU : public ImagePreprocessor { CudaDevicePtr gpu_decompressed_image; CudaDevicePtr gpu_stats; - std::vector cpu_stats; - CudaRegisteredVector cpu_stats_reg; + PinnedVector cpu_stats; CudaRegisteredVector input_reg; // page-locks the caller's decompression buffer // The uncompressed-image counterpart of input_reg: the region PinInputRegion last page-locked. const void *pinned_input = nullptr; diff --git a/image_analysis/indexing/CUDAMemHelpers.h b/image_analysis/indexing/CUDAMemHelpers.h index 6611beb05..43ae2ca5d 100644 --- a/image_analysis/indexing/CUDAMemHelpers.h +++ b/image_analysis/indexing/CUDAMemHelpers.h @@ -370,25 +370,114 @@ public: operator T*() const { return ptr_; } }; +// Page-locked host memory is slow to get and to give back: cudaHostAlloc and cudaFreeHost map and lock +// every page under the process's memory-map lock, so the calls of all threads queue behind each other - +// measured at 1-3 ms a call on a workstation and 20-30 ms on a four-socket host. The engines that hold it +// are built per worker per step, which made that thousands of calls a run. So a freed buffer is kept, in a +// power-of-two size class, for the next engine to take. Up to CUDA_PINNED_POOL_LIMIT bytes are kept, a +// buffer larger than CUDA_PINNED_POOL_MAX_BUFFER (a whole frame) is not kept at all, and +// cuda_release_pinned_pool hands everything back at the end of a run. +// +// A buffer taken from the pool holds whatever its last owner left in it. And it goes back to the pool the +// moment its owner frees it, so - as with a pooled device buffer (CudaDevicePtr) - the owner waits for +// its stream before it does: a copy still queued into the buffer would write someone else's memory. +constexpr size_t CUDA_PINNED_POOL_LIMIT = 512ull << 20; +constexpr size_t CUDA_PINNED_POOL_MAX_BUFFER = 64ull << 20; + +namespace jfjoch_cuda_pinned_pool { + struct Pool { + std::mutex m; + std::map> free; // by size class + size_t kept_bytes = 0; + }; + + // Never destroyed, as the allocation streams above: a buffer can be freed during static destruction. + inline Pool &pool() { + static Pool *p = new Pool; + return *p; + } + + // What is actually allocated for `bytes`: the next power of two from 4 kB, so that an engine whose + // buffers grew a little still finds them; and exactly `bytes` for a buffer that is never kept. + inline size_t allocated_bytes(size_t bytes) { + if (bytes > CUDA_PINNED_POOL_MAX_BUFFER) + return bytes; + size_t c = 4096; + while (c < bytes) + c *= 2; + return c; + } +} + +inline void *cuda_pinned_alloc(size_t bytes) { + using namespace jfjoch_cuda_pinned_pool; + const size_t size = allocated_bytes(bytes); + { + std::lock_guard lock(pool().m); + auto it = pool().free.find(size); + if (it != pool().free.end() && !it->second.empty()) { + void *ptr = it->second.back(); + it->second.pop_back(); + pool().kept_bytes -= size; + return ptr; + } + } + // Portable: a buffer freed by a worker on one card is taken next by a worker on another. + void *ptr = nullptr; + if (cudaHostAlloc(&ptr, size, cudaHostAllocPortable) != cudaSuccess) + throw JFJochException(JFJochExceptionCategory::MemAllocFailed, "Failed to allocate pinned host memory"); + return ptr; +} + +inline void cuda_pinned_free(void *ptr, size_t bytes) noexcept { + using namespace jfjoch_cuda_pinned_pool; + const size_t size = allocated_bytes(bytes); + if (size <= CUDA_PINNED_POOL_MAX_BUFFER) { + std::lock_guard lock(pool().m); + if (pool().kept_bytes + size <= CUDA_PINNED_POOL_LIMIT) { + pool().free[size].push_back(ptr); + pool().kept_bytes += size; + return; + } + } + cudaFreeHost(ptr); +} + +// Hand every kept buffer back to the driver. Returns how many bytes that was. +inline size_t cuda_release_pinned_pool() { + using namespace jfjoch_cuda_pinned_pool; + std::map> free; + size_t bytes = 0; + { + std::lock_guard lock(pool().m); + free.swap(pool().free); + bytes = pool().kept_bytes; + pool().kept_bytes = 0; + } + for (auto &[size, ptrs] : free) + for (void *ptr : ptrs) + cudaFreeHost(ptr); + return bytes; +} + template class CudaHostPtr { T* ptr_ = nullptr; + size_t count_ = 0; public: CudaHostPtr() = default; - explicit CudaHostPtr(size_t count) { - if (cudaMallocHost(&ptr_, count * sizeof(T)) != cudaSuccess) - throw JFJochException(JFJochExceptionCategory::MemAllocFailed, - "Failed to allocate pinned host memory"); - } + explicit CudaHostPtr(size_t count) + : ptr_(static_cast(cuda_pinned_alloc(count * sizeof(T)))), count_(count) {} ~CudaHostPtr() { - if (ptr_) cudaFreeHost(ptr_); + if (ptr_) cuda_pinned_free(ptr_, count_ * sizeof(T)); } // Move-only type - CudaHostPtr(CudaHostPtr&& other) noexcept : ptr_(other.ptr_) { other.ptr_ = nullptr; } + CudaHostPtr(CudaHostPtr&& other) noexcept : ptr_(other.ptr_), count_(other.count_) { other.ptr_ = nullptr; } CudaHostPtr& operator=(CudaHostPtr&& other) noexcept { if (this != &other) { - if (ptr_) cudaFreeHost(ptr_); + if (ptr_) cuda_pinned_free(ptr_, count_ * sizeof(T)); ptr_ = other.ptr_; + count_ = other.count_; other.ptr_ = nullptr; } return *this; @@ -400,6 +489,22 @@ public: operator T*() const { return ptr_; } }; +// A std::vector in pooled pinned memory, for the host side of an engine's per-frame copies: the copies +// are genuinely asynchronous, and the memory comes from the pool instead of being registered and +// unregistered with every engine (see cuda_pinned_alloc). +template +struct CudaPinnedAllocator { + using value_type = T; + CudaPinnedAllocator() = default; + template CudaPinnedAllocator(const CudaPinnedAllocator &) {} + T *allocate(size_t n) { return static_cast(cuda_pinned_alloc(n * sizeof(T))); } + void deallocate(T *p, size_t n) noexcept { cuda_pinned_free(p, n * sizeof(T)); } + template bool operator==(const CudaPinnedAllocator &) const { return true; } +}; + +template +using PinnedVector = std::vector>; + template class CudaRegisteredVector { std::vector* vec_ = nullptr; diff --git a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu index eae6cc795..1c9123f23 100644 --- a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu +++ b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu @@ -454,12 +454,6 @@ AdaptiveSpotFinderGPU::AdaptiveSpotFinderGPU(const AzimuthalIntegrationMapping & prof_sum(nbins), prof_sum2(nbins), prof_count(nbins), - host_sum_reg(host_sum), - host_sum2_reg(host_sum2), - host_count_reg(host_count), - prof_sum_reg(prof_sum), - prof_sum2_reg(prof_sum2), - prof_count_reg(prof_count), last_profile(in_mapping) { // The current device, not device 0: callers round-robin engines across GPUs, so device 0's shared @@ -645,7 +639,7 @@ void AdaptiveSpotFinderGPU::Detect(const ImagePreprocessorBuffer &image, sum2_corr[b] = static_cast(static_cast(static_cast(prof_sum2[b])) / CORR_FIXED); } last_profile.Clear(mapping); - last_profile.Add(sum_corr, sum2_corr, prof_count); + last_profile.Add(sum_corr, sum2_corr, std::vector(prof_count.begin(), prof_count.end())); if (host_thr.empty()) { // Nothing valid to threshold against: leave no strong pixels for the extractor to build on. diff --git a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h index 2e0f077f5..7ed9ad68e 100644 --- a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h +++ b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.h @@ -92,26 +92,19 @@ class AdaptiveSpotFinderGPU : public ImageSpotFinderGPU { CudaDevicePtr gpu_word_range; // Host mirrors of the small per-ring transfers. - std::vector host_sum; // clipped raw sum } input to the host threshold computation - std::vector host_sum2; // clipped raw sum^2 } (exact integers - see the kernel) - std::vector host_count; // clipped raw count } + PinnedVector host_sum; // clipped raw sum } input to the host threshold computation + PinnedVector host_sum2; // clipped raw sum^2 } (exact integers - see the kernel) + PinnedVector host_count; // clipped raw count } std::vector host_thr; // per-ring threshold (empty -> frame had no valid pixels) std::vector host_bkg; // clipped per-ring mean, NaN where the ring is too sparse to trust - std::vector prof_sum; // plain corrected sum } azimuthal-integration profile - std::vector prof_sum2; // plain corrected sum^2 } (fixed point) - std::vector prof_count; // plain pixel count } + PinnedVector prof_sum; // plain corrected sum } azimuthal-integration profile + PinnedVector prof_sum2; // plain corrected sum^2 } (fixed point) + PinnedVector prof_count; // plain pixel count } - // Every per-ring array above is a device-to-host copy once per frame. A D2H copy into PAGEABLE + // The pinned arrays above are each a device-to-host copy once per frame. A D2H copy into PAGEABLE // memory blocks the host until it completes, whatever stream it was issued on - which would stall // Detect() between the plain pass and the clip passes, with the device then idle while the host - // enqueues them. Pinning the destinations makes the copies genuinely asynchronous, as the - // azimuthal-integration engine already does with its own. - CudaRegisteredVector host_sum_reg; - CudaRegisteredVector host_sum2_reg; - CudaRegisteredVector host_count_reg; - CudaRegisteredVector prof_sum_reg; - CudaRegisteredVector prof_sum2_reg; - CudaRegisteredVector prof_count_reg; + // enqueues them. In pinned memory the copies are genuinely asynchronous. AzimuthalIntegrationProfile last_profile; // filled every Run(), retrievable via GetProfile() diff --git a/rugnux/Rugnux.cpp b/rugnux/Rugnux.cpp index 4d5e408e2..0110bc8f1 100644 --- a/rugnux/Rugnux.cpp +++ b/rugnux/Rugnux.cpp @@ -421,6 +421,11 @@ ProcessResult Rugnux::Run(RugnuxObserver *observer) { if (result.consensus_cell) header_experiment.SetUnitCell(*result.consensus_cell); result.beam_centre_need_pxl = BeamCenterNeed_pxl(header_experiment); +#ifdef JFJOCH_USE_CUDA + // The engines went with their passes; the pinned buffers they kept for each other go with the run. + const size_t pinned_bytes = cuda_release_pinned_pool(); + Logger("Rugnux").Debug("Pinned host memory kept between engines, released: {:.1f} MB", pinned_bytes / 1e6); +#endif return result; } diff --git a/tests/CUDAMemHelpersTest.cpp b/tests/CUDAMemHelpersTest.cpp index dad8a2f3a..0e8bb2311 100644 --- a/tests/CUDAMemHelpersTest.cpp +++ b/tests/CUDAMemHelpersTest.cpp @@ -103,6 +103,30 @@ TEST_CASE("CudaDevicePtr_HandledPoolFailureLeavesNoError", "[CUDAMemHelpers]") { cudaMemPoolDestroy(capped); } +// Engines built one after another take the pinned buffers the previous one gave back - also a little +// larger ones, within the size class - while a frame-sized buffer is not kept at all. +TEST_CASE("CudaPinnedPool_ReusedByTheNextEngine", "[CUDAMemHelpers]") { + if (get_gpu_count() == 0) + SKIP("No CUDA GPU present"); + + cuda_release_pinned_pool(); + void *first = nullptr; + { + CudaHostPtr a(1000); + first = a.get(); + } + { + PinnedVector b(1020); + CHECK(static_cast(b.data()) == first); + } + { + CudaHostPtr large(CUDA_PINNED_POOL_MAX_BUFFER + 1); + large.get()[CUDA_PINNED_POOL_MAX_BUFFER] = 1; + } + CHECK(cuda_release_pinned_pool() == 4096); + CHECK(cuda_release_pinned_pool() == 0); +} + #endif // A merge that finds the card full waits for the GPU work running beside it: until the last one