From 875c6a180aaa0ef42a3b1403243850dbec0ea941 Mon Sep 17 00:00:00 2001 From: AdrianoDee Date: Fri, 24 Feb 2023 09:11:16 +0100 Subject: [PATCH 1/2] Proper memcopy for localCoordToHostAsync --- .../TrackingRecHit/interface/TrackingRecHitSoADevice.h | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index c796725b2e8d0..eaedb3973bc9c 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h @@ -48,8 +48,10 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection localCoordToHostAsync(cudaStream_t stream) const { auto ret = cms::cuda::make_host_unique(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 From 9636d8383ba289e9752762f63c6b7c8273397474 Mon Sep 17 00:00:00 2001 From: AdrianoDee Date: Mon, 27 Feb 2023 10:11:37 +0100 Subject: [PATCH 2/2] Code checks --- .../TrackingRecHit/interface/TrackingRecHitSoADevice.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index eaedb3973bc9c..ab0043930558d 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h @@ -50,8 +50,9 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection