From 90ba15871f2ee1a83a75b8258a9c09f7db676eed Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Wed, 14 Oct 2026 16:40:34 +0200 Subject: [PATCH] CudaFFTPlan: wait for the plan's NULL-stream setup before it is used, output unchanged A transform launched on a non-blocking stream straight after cufftPlanMany could read the plan before its setup, still queued on the NULL stream, had finished: Xid 31, a read at a fixed near-null address, "an illegal memory access" in the indexer threads. It showed at transform length 3072, in 3 of 4 runs of one rotation set where a pool built its FFT indexers at the start of integration (6 of 28 ci-tier sets, 9 of 25 null-model runs). It went away under CUDA_LAUNCH_BLOCKING=1 and with a cudaStreamSynchronize(NULL) after the plan (0 of 11 runs). The 3142 length the indexer used so far never hit it, but nothing ordered it either. Plans are built once per indexer, so the wait costs nothing measurable. The other cuFFT users (beam-centre FFT, model structure factors, rigid body) create raw plans and are not covered by this wrapper. Co-Authored-By: Claude Opus 5.5 (1M context) --- image_analysis/indexing/CUDAMemHelpers.h | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/image_analysis/indexing/CUDAMemHelpers.h b/image_analysis/indexing/CUDAMemHelpers.h index 0aa741556..6d2db7ad0 100644 --- a/image_analysis/indexing/CUDAMemHelpers.h +++ b/image_analysis/indexing/CUDAMemHelpers.h @@ -128,6 +128,13 @@ public: throw JFJochException(JFJochExceptionCategory::GPUCUDAError, GPUMemoryFailure("Creating the CUFFT plan (cufftPlanMany)", 0, true, ("cuFFT error " + std::to_string(r)).c_str())); + // The plan's setup can still be queued on the NULL stream when cufftPlanMany returns, and a + // non-blocking stream (CudaStream) does not wait for it: a transform launched on one straight + // away read the plan's tables before they were there - an illegal address in 3 of 4 runs where + // a rotation pool built its FFT indexers at transform length 3072. Waiting here orders it. + if (cudaStreamSynchronize(nullptr) != cudaSuccess) + throw JFJochException(JFJochExceptionCategory::GPUCUDAError, + "Creating the CUFFT plan: setup on the NULL stream failed"); } // Convenience overload for vector input