From 2be86804228a484f8aca5ab00400a94eef85c32d Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Thu, 30 Jul 2026 23:05:11 +0200 Subject: [PATCH] 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) --- image_analysis/image_preprocessing/ImagePreprocessorGPU.cu | 5 ++++- image_analysis/indexing/CUDAMemHelpers.h | 5 ++++- 2 files changed, 8 insertions(+), 2 deletions(-) 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(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");