diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu
index 2de35c7d..de4401de 100644
--- a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu
+++ b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu
@@ -424,8 +424,12 @@ BraggIntegrationEngineGPU::BraggIntegrationEngineGPU(const DiffractionExperiment
     const int max_Gf = 2 * max_Rf + 1;
     fit_shared_bytes = static_cast<size_t>(max_Gf) * max_Gf * sizeof(float);
 
+    // The current device, not device 0: workers are pinned round-robin across GPUs, so device 0's
+    // shared-memory size can belong to a different card than the one these kernels launch on.
+    int device = 0;
+    cuda_err(cudaGetDevice(&device));
     cudaDeviceProp prop{};
-    cuda_err(cudaGetDeviceProperties(&prop, 0));
+    cuda_err(cudaGetDeviceProperties(&prop, device));
     if (fit_shared_bytes > prop.sharedMemPerBlock)
         throw JFJochException(JFJochExceptionCategory::GPUCUDAError,
                               "BraggIntegrationEngineGPU: profile grid exceeds shared memory (r2 too large)");
diff --git a/image_analysis/spot_finding/SpotExtractorGPU.cu b/image_analysis/spot_finding/SpotExtractorGPU.cu
index 5c8d7c99..ec40d6cb 100644
--- a/image_analysis/spot_finding/SpotExtractorGPU.cu
+++ b/image_analysis/spot_finding/SpotExtractorGPU.cu
@@ -296,7 +296,11 @@ SpotExtractorGPU::SpotExtractorGPU(int32_t in_width, int32_t in_height, std::sha
     const size_t npixel = static_cast<size_t>(in_width) * in_height;
     if (npixel % 32 != 0)
         mask.back() = ~((1u << (npixel % 32)) - 1u);
-    cuda_err(cudaMemcpy(gpu_res_mask, mask.data(), nwords * sizeof(uint32_t), cudaMemcpyHostToDevice));
+    // On this engine's stream, then synchronised - the default stream is non-blocking, so a NULL-stream
+    // copy is not ordered against the kernels that read this mask.
+    cuda_err(cudaMemcpyAsync(gpu_res_mask, mask.data(), nwords * sizeof(uint32_t),
+                             cudaMemcpyHostToDevice, *stream));
+    cuda_err(cudaStreamSynchronize(*stream));
 }
 
 void SpotExtractorGPU::SetResolutionMask(const std::vector<uint32_t> &packed_mask) {