rugnux pre-scan: spot width on per-spot squares, engines built during the read, tiers looked ahead, measurements memoised

The spot-width stage of the pre-scan copied every preprocessed frame back to the host (72 MB a
frame on a 16M detector, page-locked host mirrors per worker) for measurements that read a few
pixels about each spot. Now:

- MeasureSpotsOnTiles: the flux curves and the footprint read (2*24+1)^2 squares about each spot,
  gathered on the card (ImagePreprocessorBufferGPU::GatherTiles) - no copy_image_to_host, no host
  mirror. The flux curves are bounded at R_BKG_OUT = 20 px. The footprint keeps its arithmetic; a
  per-iteration check that its window is inside the square, and a spot whose window walks off it is
  measured again from scratch on an 80 px square, then on the whole frame. Where the squares would
  be more pixels than the frame, the frame is read. Same pixels in the same order, so the same bits.
  Harness over 34 rotation sets (~1500 frames): 0 differences; misses at 24 on a few spots of most
  frames, at 54 on up to 80% of frames of a few sets, at 80 on none.
- The azimuthal mapping and the per-worker spot engines are built while the frames are read.
- The tiers no longer wait for each other: workers run on through the visit order, each tier's
  answer is drawn from exactly its own frames, and frames measured past the tier the answer stops
  at are thrown away - a frame is measured the same whichever tier it is read for.
- PreScanMemo keeps the spot measurement, the defective pixels and the background gain under the
  pre-scan key plus the spot-finding and azimuthal settings, so a copy of the run that pre-scans
  again on the same inputs reads no frames.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3

image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu

@@ -3,6 +3,9 @@
 
 #include "ImagePreprocessorBufferGPU.h"
 
+#include <algorithm>
+#include <climits>
+
 static inline void cuda_err(cudaError_t val) {
     if (val != cudaSuccess)
         throw JFJochException(JFJochExceptionCategory::GPUCUDAError, cudaGetErrorString(val));
@@ -16,6 +19,21 @@ __global__ void gather_kernel(const int32_t *__restrict__ image,
         values[i] = image[npixel[i]];
 }
 
+__global__ void gather_tiles_kernel(const int32_t *__restrict__ image, int width, int height,
+                                    const int32_t *__restrict__ centre, int half,
+                                    int32_t *__restrict__ tiles, size_t count) {
+    const int side = 2 * half + 1;
+    const size_t area = static_cast<size_t>(side) * side;
+    for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < count; i += static_cast<size_t>(blockDim.x) * gridDim.x) {
+        const size_t tile = i / area;
+        const int p = static_cast<int>(i - tile * area);
+        const int px = centre[2 * tile] - half + p % side;
+        const int py = centre[2 * tile + 1] - half + p / side;
+        tiles[i] = (px < 0 || py < 0 || px >= width || py >= height)
+                   ? INT32_MIN : image[static_cast<size_t>(py) * width + px];
+    }
+}
+
 ImagePreprocessorBufferGPU::ImagePreprocessorBufferGPU(size_t npixel, bool host_mirror)
     : ImagePreprocessorBuffer(npixel, host_mirror),
       gpu_image(npixel),
@@ -49,3 +67,42 @@ void ImagePreprocessorBufferGPU::Gather(const std::vector<uint32_t> &npixel, std
                     cudaMemcpyDeviceToHost, gather_stream);
     cudaStreamSynchronize(gather_stream);
 }
+
+const int32_t *ImagePreprocessorBufferGPU::GatherTiles(const std::vector<int32_t> &x, const std::vector<int32_t> &y,
+                                                       int width, int half) {
+    const size_t side = 2 * half + 1;
+    const size_t count = x.size() * side * side;
+    if (count == 0)
+        return nullptr;
+    if (count > tile_capacity) {
+        tile_capacity = std::max(count, 2 * tile_capacity);
+        gpu_tiles = CudaDevicePtr<int32_t>(tile_capacity);
+        host_tiles.resize(tile_capacity);
+    }
+    if (2 * x.size() > tile_centre.size()) {
+        tile_centre.resize(std::max(2 * x.size(), 2 * tile_centre.size()));
+        gpu_tile_centre = CudaDevicePtr<int32_t>(tile_centre.size());
+    }
+    for (size_t i = 0; i < x.size(); i++) {
+        tile_centre[2 * i] = x[i];
+        tile_centre[2 * i + 1] = y[i];
+    }
+    cuda_err(cudaMemcpyAsync(gpu_tile_centre.get(), tile_centre.data(), 2 * x.size() * sizeof(int32_t),
+                             cudaMemcpyHostToDevice, gather_stream));
+    const int blocks = static_cast<int>(std::min<size_t>((count + 255) / 256, 4096));
+    gather_tiles_kernel<<<blocks, 256, 0, gather_stream>>>(gpu_image.get(), width, static_cast<int>(npixels / width),
+                                                           gpu_tile_centre.get(), half, gpu_tiles.get(), count);
+    cuda_err(cudaGetLastError());
+    cuda_err(cudaMemcpyAsync(host_tiles.data(), gpu_tiles.get(), count * sizeof(int32_t),
+                             cudaMemcpyDeviceToHost, gather_stream));
+    cuda_err(cudaStreamSynchronize(gather_stream));
+    return host_tiles.data();
+}
+
+const int32_t *ImagePreprocessorBufferGPU::HostImage() {
+    buffer.resize(npixels);
+    cuda_err(cudaMemcpyAsync(buffer.data(), gpu_image.get(), npixels * sizeof(int32_t),
+                             cudaMemcpyDeviceToHost, gather_stream));
+    cuda_err(cudaStreamSynchronize(gather_stream));
+    return buffer.data();
+}