diff --git a/image_analysis/MXAnalysisWithoutFPGA.cpp b/image_analysis/MXAnalysisWithoutFPGA.cpp index 3cbf1ba5..63a55f58 100644 --- a/image_analysis/MXAnalysisWithoutFPGA.cpp +++ b/image_analysis/MXAnalysisWithoutFPGA.cpp @@ -59,7 +59,11 @@ MXAnalysisWithoutFPGA::MXAnalysisWithoutFPGA(const DiffractionExperiment &in_exp #ifdef JFJOCH_USE_CUDA } else { stream = std::make_shared(); - preprocessor_buffer = std::make_unique(experiment.GetPixelsNum()); + // The host copy of the preprocessed image is only read when a CPU engine wants it, which is + // the same condition that drives copy_image_to_host below. Skipping it also skips page-locking + // 4 bytes per pixel per worker. + preprocessor_buffer = std::make_unique( + experiment.GetPixelsNum(), /*host_mirror=*/!enable_fused_adaptive_gpu); // The preprocessed image only has to come back to the host if a CPU engine reads it. Every // engine built below runs on the GPU, except the CPU adaptive finder that is kept when the fused // GPU engine is off - so that is the one case that needs the copy. Every caller currently passes diff --git a/image_analysis/image_preprocessing/ImagePreprocessorBuffer.cpp b/image_analysis/image_preprocessing/ImagePreprocessorBuffer.cpp index 0a2728a3..46cae1d2 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorBuffer.cpp +++ b/image_analysis/image_preprocessing/ImagePreprocessorBuffer.cpp @@ -3,7 +3,11 @@ #include "ImagePreprocessorBuffer.h" -ImagePreprocessorBuffer::ImagePreprocessorBuffer(size_t npixels) : buffer(npixels) {} +ImagePreprocessorBuffer::ImagePreprocessorBuffer(size_t npixels, bool host_mirror) + : npixels(npixels), buffer(host_mirror ? npixels : 0) {} + +ImagePreprocessorBuffer::ImagePreprocessorBuffer(size_t npixels) + : ImagePreprocessorBuffer(npixels, true) {} void ImagePreprocessorBuffer::Gather(const std::vector &npixel, std::vector &values) const { values.resize(npixel.size()); diff --git a/image_analysis/image_preprocessing/ImagePreprocessorBuffer.h b/image_analysis/image_preprocessing/ImagePreprocessorBuffer.h index a0ef88c5..c5730515 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorBuffer.h +++ b/image_analysis/image_preprocessing/ImagePreprocessorBuffer.h @@ -9,7 +9,15 @@ class ImagePreprocessorBuffer { protected: + // The image size, kept apart from `buffer` so size() still answers when the host copy was not + // allocated at all (see the host_mirror constructor). + const size_t npixels; std::vector buffer; + + // host_mirror = false leaves `buffer` empty. On the GPU path the preprocessed image lives on the + // device and no CPU engine reads it back, so the four-bytes-per-pixel host copy would be + // allocated, zeroed and page-locked for nothing - 72 MB per worker on a 16M-pixel detector. + ImagePreprocessorBuffer(size_t npixels, bool host_mirror); public: explicit ImagePreprocessorBuffer(size_t npixels); virtual ~ImagePreprocessorBuffer() = default; @@ -23,7 +31,7 @@ public: const std::vector &getBuffer() const { return buffer; } int32_t &operator[](size_t i) { return buffer[i]; } const int32_t &operator[](size_t i) const { return buffer[i]; } - size_t size() const { return buffer.size(); } + size_t size() const { return npixels; } int32_t *data() { return buffer.data(); } const int32_t *data() const { return buffer.data(); } diff --git a/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu index ed02bd30..54265f5c 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu @@ -11,9 +11,10 @@ __global__ void gather_kernel(const int32_t *__restrict__ image, values[i] = image[npixel[i]]; } -ImagePreprocessorBufferGPU::ImagePreprocessorBufferGPU(size_t npixel) - : ImagePreprocessorBuffer(npixel), +ImagePreprocessorBufferGPU::ImagePreprocessorBufferGPU(size_t npixel, bool host_mirror) + : ImagePreprocessorBuffer(npixel, host_mirror), gpu_image(npixel), + // A no-op when the mirror was not allocated: CudaRegisteredVector skips an empty vector. buffer_reg(buffer), gpu_gather_index(MAX_GATHER), gpu_gather_value(MAX_GATHER) { diff --git a/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.h b/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.h index d03a94aa..dd223be1 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.h +++ b/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.h @@ -23,7 +23,9 @@ class ImagePreprocessorBufferGPU : public ImagePreprocessorBuffer { CudaStream gather_stream; public: - explicit ImagePreprocessorBufferGPU(size_t npixel); + // host_mirror = false skips the host copy of the preprocessed image entirely (and with it the + // page-locking): pass it when every engine reading this buffer runs on the device. + explicit ImagePreprocessorBufferGPU(size_t npixel, bool host_mirror = true); int32_t *getGPUBuffer() override; const int32_t *getGPUBuffer() const override; diff --git a/image_analysis/indexing/CUDAMemHelpers.h b/image_analysis/indexing/CUDAMemHelpers.h index 3af9fbd0..3f975e96 100644 --- a/image_analysis/indexing/CUDAMemHelpers.h +++ b/image_analysis/indexing/CUDAMemHelpers.h @@ -195,6 +195,13 @@ class CudaRegisteredVector { if (err != cudaSuccess) throw JFJochException(JFJochExceptionCategory::GPUCUDAError, "cudaHostUnregister failed"); } + // Unchecked, for the destructor and the move-assignment: both are noexcept, so throwing out of + // them aborts the process instead of reporting the failure - and teardown is exactly where + // cudaHostUnregister fails (after a device reset, or while another exception is unwinding). + // The other destructors in this header ignore their teardown status for the same reason. + static void unregisterPtrNoThrow(void* ptr) noexcept { + cudaHostUnregister(ptr); + } public: // Non-owning wrapper. Does NOT provide accessors to the vector. @@ -211,7 +218,7 @@ public: ~CudaRegisteredVector() { if (registered_ && vec_ && !vec_->empty()) { - unregisterPtr(vec_->data()); + unregisterPtrNoThrow(vec_->data()); } } @@ -226,7 +233,7 @@ public: if (this != &other) { // Clean current registration if (registered_ && vec_ && !vec_->empty()) { - unregisterPtr(vec_->data()); + unregisterPtrNoThrow(vec_->data()); } vec_ = other.vec_; registered_ = other.registered_; diff --git a/image_analysis/roi/ROIIntegrationGPU.cu b/image_analysis/roi/ROIIntegrationGPU.cu index 68fe1cbb..db215670 100644 --- a/image_analysis/roi/ROIIntegrationGPU.cu +++ b/image_analysis/roi/ROIIntegrationGPU.cu @@ -107,8 +107,12 @@ ROIIntegrationGPU::ROIIntegrationGPU(const DiffractionExperiment &experiment, st host_max(roi_count), max_init(roi_count, INT_MIN) { + // The current device, not device 0: workers are pinned round-robin across the GPUs, so device 0's + // SM count can belong to a different card than the one these kernels launch on. + int device = 0; + cuda_err(cudaGetDevice(&device)); cudaDeviceProp prop{}; - cuda_err(cudaGetDeviceProperties(&prop, 0)); + cuda_err(cudaGetDeviceProperties(&prop, device)); threads = 128; blocks = 4 * prop.multiProcessorCount; shared_needed = roi_count * (5 * sizeof(unsigned long long) + sizeof(int));