diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu
index 1a1f36cb..02bd5c44 100644
--- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu
+++ b/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);
diff --git a/image_analysis/indexing/CUDAMemHelpers.h b/image_analysis/indexing/CUDAMemHelpers.h
index 8eb25348..ce97bf42 100644
--- a/image_analysis/indexing/CUDAMemHelpers.h
+++ b/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");