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<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));
     }
 }
 
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> &coord, size_t nspots) {