CUDA: let worker streams run concurrently

Every per-thread stream was created with cudaStreamDefault, and the 20 MB raw
image upload went to the legacy NULL stream. A NULL-stream operation implicitly
synchronises with every blocking stream in the process, so with one engine per
worker thread no two workers' GPU work could ever overlap - the whole GPU
pipeline ran serially however many threads were asked for.

Create the streams non-blocking and put the upload on the engine's own stream.
Measured on 2000 serial stills, interleaved, medians of three: 24.6 -> 19.2 s at
-N 32 (-22%), 32.7 -> 21.0 s at -N 16 (-36%), CPU utilisation 436-570% -> 723-859%.
Output bit-identical - same observations, uniques, completeness, R-meas, CC1/2,
error model and cell. The stream is synchronised at the end of the same function,
so the ordering the code relies on is unchanged.

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

image_analysis/image_preprocessing/ImagePreprocessorGPU.cu

@@ -141,7 +141,10 @@ ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed
     if (sat_value > saturation_limit)
         sat_value = static_cast<T>(saturation_limit);
 
-    cudaMemcpy(gpu_decompressed_image, input, npixels * sizeof(T), cudaMemcpyHostToDevice);
+    // On this engine's own stream, not the NULL stream: a NULL-stream copy implicitly synchronises with
+    // every blocking stream in the process, which serialised all workers behind whichever one was
+    // uploading. The stream is synchronised at the end of this function, so the ordering is unchanged.
+    cudaMemcpyAsync(gpu_decompressed_image, input, npixels * sizeof(T), cudaMemcpyHostToDevice, *stream);
 
     cpu_stats[0] = ImageStatistics{.max_value = INT64_MIN, .min_value = INT64_MAX};
     cudaMemcpyAsync(gpu_stats, cpu_stats.data(), sizeof(ImageStatistics), cudaMemcpyHostToDevice, *stream);

image_analysis/indexing/CUDAMemHelpers.h

@@ -12,7 +12,10 @@
 class CudaStream {
     cudaStream_t stream_ = nullptr;
 public:
-    CudaStream(unsigned int flags = cudaStreamDefault) {
+    // Non-blocking by default: a stream created with cudaStreamDefault synchronises against the legacy
+    // NULL stream, so any NULL-stream operation anywhere in the process serialises every worker's GPU
+    // work against every other's. With one engine per worker thread that costs most of the parallelism.
+    CudaStream(unsigned int flags = cudaStreamNonBlocking) {
         if (cudaStreamCreateWithFlags(&stream_, flags) != cudaSuccess)
             throw JFJochException(JFJochExceptionCategory::GPUCUDAError,
                                   "Failed to create CUDA stream");