From 1cbb29c69bacfaefe0b6715d8d8c885b91400b5c Mon Sep 17 00:00:00 2001 From: AdrianoDee Date: Fri, 24 Feb 2023 09:00:53 +0100 Subject: [PATCH 1/2] Proper memcopy for localCoordToHostAsync --- .../TrackingRecHit/interface/TrackingRecHitSoADevice.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index c796725b2e8d0..244f214a57601 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 013581a7353f21e0d931beea7ed1398f1071c727 Mon Sep 17 00:00:00 2001 From: AdrianoDee Date: Mon, 27 Feb 2023 10:18:30 +0100 Subject: [PATCH 2/2] Code checks --- .../TrackingRecHit/interface/TrackingRecHitSoADevice.h | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index 244f214a57601..ab0043930558d 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h @@ -48,10 +48,11 @@ 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(); - + size_t srcPitch = ptrdiff_t(view().yLocal()) - ptrdiff_t(view().xLocal()); - cudaCheck(cudaMemcpy2DAsync(ret.get(), rowSize, view().xLocal(), srcPitch, rowSize, 4, cudaMemcpyDeviceToHost, stream)); - + cudaCheck( + cudaMemcpy2DAsync(ret.get(), rowSize, view().xLocal(), srcPitch, rowSize, 4, cudaMemcpyDeviceToHost, stream)); + return ret; } //move to utilities