mx/Jungfraujoch
3
2
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases 81 Wiki Activity

BraggIntegrationEngineGPU, FFTIndexerGPU: upload on the engine's stream, not the NULL stream

Browse Source 
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) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01D1G8gJVAy6gp1K5Dz3NE5C
This commit is contained in:
leonarski_f
2026-09-27 17:11:38 +02:00
co-authored by Claude Opus 5.5
parent ca1baffef4
commit cb97832fb3
2 changed files with 10 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu
+5 -3
View File
@@ -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<unsigned long long>(n_rad);
d_rad_cnt = CudaDevicePtr<int>(n_rad);
d_k_diff = CudaDevicePtr<float>(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));
}
}
image_analysis/indexing/FFTIndexerGPU.cu
+5 -4
View File
@@ -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> &coord, size_t nspots) {