Merge pull request #40870 from AdrianoDee/fix_localCoordToHostAsync

[13_0_X] Fix `cudaMemcpyAsync` for `localCoordToHostAsync`
cms-sw · Mar 1, 2023 · 89e3634 · 89e3634
2 parents 885f88d + 013581a
commit 89e3634
Showing 1 changed file with 4 additions and 1 deletion.
diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h
@@ -48,7 +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