v1.0.0-rc.179 (#90)

* Rugnux: Stills processing needs only the unit cell and space group beyond what the file records - detector distance, beam centre, tilt, panel placement, saturation level, gain and polarisation plane are found from the images - and is up to twice as fast.
* Rugnux: Stills can be indexed without a GPU, with a CPU implementation of the fast-feedback indexer (macOS, computers without a GPU).
* Rugnux: Rotation and stills data are integrated in the detector's own units with one gain model, improving data from CCD detectors.
* Web frontend: The Rugnux live page shows the placed model with its electron density and anomalous difference map in a Mol* viewer.

Reviewed-on: #90
Co-authored-by: Filip Leonarski <filip.leonarski@psi.ch>

image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu

@@ -19,7 +19,8 @@ __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,
+// centre holds six values per tile: its centre x, y and the box x0, y0, x1, y1 it is clipped to.
+__global__ void gather_tiles_kernel(const int32_t *__restrict__ image, int width,
                                     const int32_t *__restrict__ centre, int half,
                                     int32_t *__restrict__ tiles, size_t count) {
     const int side = 2 * half + 1;
@@ -27,9 +28,10 @@ __global__ void gather_tiles_kernel(const int32_t *__restrict__ image, int width
     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)
+        const int32_t *c = centre + 6 * tile;
+        const int px = c[0] - half + p % side;
+        const int py = c[1] - half + p / side;
+        tiles[i] = (px < c[2] || py < c[3] || px >= c[4] || py >= c[5])
                    ? INT32_MIN : image[static_cast<size_t>(py) * width + px];
     }
 }
@@ -79,18 +81,24 @@ const int32_t *ImagePreprocessorBufferGPU::GatherTiles(const std::vector<int32_t
         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()));
+    if (6 * x.size() > tile_centre.size()) {
+        tile_centre.resize(std::max(6 * 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];
+        const PanelBox b = TileBox(x[i], y[i], width);
+        int32_t *c = tile_centre.data() + 6 * i;
+        c[0] = x[i];
+        c[1] = y[i];
+        c[2] = b.x0;
+        c[3] = b.y0;
+        c[4] = b.x1;
+        c[5] = b.y1;
     }
-    cuda_err(cudaMemcpyAsync(gpu_tile_centre.get(), tile_centre.data(), 2 * x.size() * sizeof(int32_t),
+    cuda_err(cudaMemcpyAsync(gpu_tile_centre.get(), tile_centre.data(), 6 * 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),
+    gather_tiles_kernel<<<blocks, 256, 0, gather_stream>>>(gpu_image.get(), 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),