Stop allocating GPU and pinned memory nothing reads

Three resource fixes and two latent bugs, none of which changes a computed
number.

The preprocessed image has a host copy that only a CPU engine ever reads. On
the GPU path every engine reads the device buffer instead, and rugnux always
runs the fused adaptive finder, so that host copy is allocated, zeroed and
PAGE-LOCKED for nothing - 72 MB per worker, 3.5 GB over 48 of them, and a
cudaHostRegister each, which the driver serializes. It is now skipped by the
same condition that already decides whether the device copies the image back.
ImagePreprocessorBuffer keeps the pixel count separately so size() still
answers when the mirror was not allocated.

ROIIntegrationGPU asked device 0 for the SM count it sizes its grid from, while
workers are pinned round-robin across the GPUs - so on a multi-GPU node it
could size a grid from a card it never launches on. It asks the current device
now, like every other engine.

~CudaRegisteredVector called a function that throws out of a destructor, and
the move-assignment did the same from a noexcept function. Either would abort
the process rather than report the failure, and teardown - after a device
reset, or while another exception unwinds - is exactly where cudaHostUnregister
fails. Both now use an unchecked unregister, as every other destructor in that
header already does for its own teardown call. The throwing form stays for
rebind()/unregister(), which are called from live code.

Measured on a 16M-pixel rotation dataset: unchanged space group, merged
reflection count and merging statistics.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

image_analysis/MXAnalysisWithoutFPGA.cpp

@@ -59,7 +59,11 @@ MXAnalysisWithoutFPGA::MXAnalysisWithoutFPGA(const DiffractionExperiment &in_exp
 #ifdef JFJOCH_USE_CUDA
     } else {
         stream = std::make_shared<CudaStream>();
-        preprocessor_buffer = std::make_unique<ImagePreprocessorBufferGPU>(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<ImagePreprocessorBufferGPU>(
+                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

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<uint32_t> &npixel, std::vector<int32_t> &values) const {
     values.resize(npixel.size());

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

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) {

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;

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_;

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