Proper memcopy for localCoordToHostAsync

cms-sw · Feb 27, 2023 · 1cbb29c · 1cbb29c
1 parent e4a95d9
commit 1cbb29c
Showing 1 changed file with 4 additions and 2 deletions.
diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h
@@ -48,8 +48,10 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection<Track
   cms::cuda::host::unique_ptr<float[]> localCoordToHostAsync(cudaStream_t stream) const {
     auto ret = cms::cuda::make_host_unique<float[]>(4 * nHits(), stream);
     size_t rowSize = sizeof(float) * nHits();
-    cudaCheck(cudaMemcpyAsync(ret.get(), view().xLocal(), rowSize * 4, cudaMemcpyDefault, stream));
-
+
+    size_t srcPitch = ptrdiff_t(view().yLocal()) - ptrdiff_t(view().xLocal());
+    cudaCheck(cudaMemcpy2DAsync(ret.get(), rowSize, view().xLocal(), srcPitch, rowSize, 4, cudaMemcpyDeviceToHost, stream));
+
     return ret;
   }  //move to utilities