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) <noreply@anthropic.com>
leonarski_fandClaude Opus 5.5
90ba15871f