From cb97832fb34ed68d537d508ae0bb48c97858e58e Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Sun, 27 Sep 2026 17:11:38 +0200 Subject: [PATCH] BraggIntegrationEngineGPU, FFTIndexerGPU: upload on the engine's stream, not the NULL stream The engines' kernels run on their own non-blocking stream, which is not ordered after the legacy NULL stream. The counter reset (cudaMemset, asynchronous to the host) and the radial-background kernel upload in the Bragg engine constructor, and the direction-grid upload in FFTIndexerGPU (a pageable cudaMemcpy returns once staged, before the DMA completes) were issued on the NULL stream, so the first kernel reading them was not ordered after them. Now cudaMemsetAsync/cudaMemcpyAsync on the engine's stream. In the Bragg constructor that is this->stream: the parameter of the same name has been moved from. compute-sanitizer --track-stream-ordered-races all flagged the Bragg ones (16 + 16 reports on rugnux -e 150 myob; 1 in the [CUDAMemHelpers] tests); 0 after. Outputs byte-identical on myob, cytc, lyso. Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_01D1G8gJVAy6gp1K5Dz3NE5C --- .../bragg_integration/BraggIntegrationEngineGPU.cu | 8 +++++--- image_analysis/indexing/FFTIndexerGPU.cu | 9 +++++---- 2 files changed, 10 insertions(+), 7 deletions(-) diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu index 8aea69444..84cd2305d 100644 --- a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu +++ b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu @@ -797,7 +797,9 @@ BraggIntegrationEngineGPU::BraggIntegrationEngineGPU(const DiffractionExperiment d_invd2(2), d_counts(COUNT_SLOTS) { threads = 128; - cuda_err(cudaMemset(d_counts, 0, sizeof(unsigned long long) * COUNT_SLOTS)); + // On the engine's stream (the member: the parameter was moved from), where the kernels counting + // into it run - the NULL stream is not ordered before those. + cuda_err(cudaMemsetAsync(d_counts, 0, sizeof(unsigned long long) * COUNT_SLOTS, *this->stream)); // Fit profile grid: R for empirical / box, up to 3R (radially elongated) for the Gaussian. const int max_Rf = empirical ? R : 3 * R; @@ -841,8 +843,8 @@ BraggIntegrationEngineGPU::BraggIntegrationEngineGPU(const DiffractionExperiment d_rad_sum = CudaDevicePtr(n_rad); d_rad_cnt = CudaDevicePtr(n_rad); d_k_diff = CudaDevicePtr(k_diff.size()); - cuda_err(cudaMemcpy(d_k_diff, k_diff.data(), sizeof(float) * k_diff.size(), - cudaMemcpyHostToDevice)); + cuda_err(cudaMemcpyAsync(d_k_diff, k_diff.data(), sizeof(float) * k_diff.size(), + cudaMemcpyHostToDevice, *this->stream)); } } diff --git a/image_analysis/indexing/FFTIndexerGPU.cu b/image_analysis/indexing/FFTIndexerGPU.cu index a769801f0..7672fc1a5 100644 --- a/image_analysis/indexing/FFTIndexerGPU.cu +++ b/image_analysis/indexing/FFTIndexerGPU.cu @@ -194,10 +194,11 @@ void FFTIndexerGPU::DirectionsChanged() { } // Checked like every other CUDA call here: a silently failed upload leaves the device holding the // PREVIOUS grid, so the directions the host reads the results against are not the ones the kernel - // used - wrong rows reported as good ones, and nothing says so. - cuda_err(cudaMemcpy(d_dir_x, dir_x.data(), nDirections * sizeof(float), cudaMemcpyHostToDevice)); - cuda_err(cudaMemcpy(d_dir_y, dir_y.data(), nDirections * sizeof(float), cudaMemcpyHostToDevice)); - cuda_err(cudaMemcpy(d_dir_z, dir_z.data(), nDirections * sizeof(float), cudaMemcpyHostToDevice)); + // used - wrong rows reported as good ones, and nothing says so. On the engine's stream, which the + // kernels reading them run on: a cudaMemcpy on the NULL stream is not ordered before those. + cuda_err(cudaMemcpyAsync(d_dir_x, dir_x.data(), nDirections * sizeof(float), cudaMemcpyHostToDevice, stream)); + cuda_err(cudaMemcpyAsync(d_dir_y, dir_y.data(), nDirections * sizeof(float), cudaMemcpyHostToDevice, stream)); + cuda_err(cudaMemcpyAsync(d_dir_z, dir_z.data(), nDirections * sizeof(float), cudaMemcpyHostToDevice, stream)); } void FFTIndexerGPU::ExecuteFFT(const std::vector &coord, size_t nspots) {