From 0e49a367337570dd2e867ee929420c4e29b288ea Mon Sep 17 00:00:00 2001 From: Vincenzo Innocente Date: Thu, 5 May 2022 12:06:49 +0200 Subject: [PATCH] compiles --- .../interface/TrackingRecHit2DHeterogeneous.h | 54 ++++++++++--------- .../src/TrackingRecHit2DHeterogeneous.cc | 18 +++---- 2 files changed, 37 insertions(+), 35 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h index d2d33d7b4e43d..5ba91165d506c 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h @@ -32,7 +32,7 @@ class TrackingRecHit2DHeterogeneous { }; template - using unique_ptr = typename memoryPool::unique_ptr; + using buffer = typename memoryPool::buffer; using PhiBinner = TrackingRecHit2DSOAView::PhiBinner; @@ -69,12 +69,12 @@ class TrackingRecHit2DHeterogeneous { auto phiBinnerStorage() { return m_phiBinnerStorage; } auto iphi() { return m_iphi; } - unique_ptr localCoordToHostAsync(cudaStream_t stream) const; + buffer localCoordToHostAsync(cudaStream_t stream) const; - unique_ptr hitsModuleStartToHostAsync(cudaStream_t stream) const; + buffer hitsModuleStartToHostAsync(cudaStream_t stream) const; - unique_ptr store16ToHostAsync(cudaStream_t stream) const; - unique_ptr store32ToHostAsync(cudaStream_t stream) const; + buffer store16ToHostAsync(cudaStream_t stream) const; + buffer store32ToHostAsync(cudaStream_t stream) const; // needs specialization for Host void copyFromGPU(TrackingRecHit2DHeterogeneous const* input, cudaStream_t stream); @@ -84,13 +84,13 @@ class TrackingRecHit2DHeterogeneous { static constexpr uint32_t n32 = 10; // number of elements in m_store32 static_assert(sizeof(uint32_t) == sizeof(float)); // just stating the obvious static_assert(n32 == static_cast(Storage32::kLayers)); - unique_ptr m_store16; //! - unique_ptr m_store32; //! + buffer m_store16; //! + buffer m_store32; //! - unique_ptr m_PhiBinnerStore; //! - unique_ptr m_AverageGeometryStore; //! + buffer m_PhiBinnerStore; //! + buffer m_AverageGeometryStore; //! - unique_ptr m_view; //! + buffer m_view; //! uint32_t m_nHits; int32_t m_offsetBPIX2; @@ -120,15 +120,16 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous( cudaStream_t stream, TrackingRecHit2DHeterogeneous const* input) : m_nHits(nHits), m_offsetBPIX2(offsetBPIX2), m_hitsModuleStart(hitsModuleStart) { + using namespace memoryPool::cuda; - auto view = make_unique(1,stream,memoryPool::onCPU==where ? memoryPool::onCPU : memoryPool::onHost); + auto view = make_buffer(1,stream,memoryPool::onCPU==where ? memoryPool::onCPU : memoryPool::onHost); m_nMaxModules = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules; view->m_nHits = nHits; view->m_nMaxModules = m_nMaxModules; - m_view = make_unique(1,stream,where); // leave it on host and pass it by value? - m_AverageGeometryStore = make_unique(1,stream,where); + m_view = make_buffer(1,stream,where); // leave it on host and pass it by value? + m_AverageGeometryStore = make_buffer(1,stream,where); view->m_averageGeometry = m_AverageGeometryStore.get(); view->m_cpeParams = cpeParams; view->m_hitsModuleStart = hitsModuleStart; @@ -136,7 +137,7 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous( // if empy do not bother if (0 == nHits) { if constexpr (memoryPool::onDevice == where) { - cudaCheck(cudaMemcpyAsync(m_view.get(), view.get(), sizeof(TrackingRecHit2DSOAView), cudaMemcpyHostToDevice, stream)); + cudaCheck(memoryPool::cuda::copy(m_view, view, sizeof(TrackingRecHit2DSOAView), stream)); } else { m_view.reset(view.release()); // NOLINT: std::move() breaks CUDA version } @@ -158,9 +159,9 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous( auto nL = isPhase2 ? phase2PixelTopology::numberOfLayers : phase1PixelTopology::numberOfLayers; - m_store16 = make_unique(nHits * n16, stream, where); - m_store32 = make_unique(nHits * n32 + nL + 1, stream, where); - m_PhiBinnerStore = make_unique(1,stream,where); + m_store16 = make_buffer(nHits * n16, stream, where); + m_store32 = make_buffer(nHits * n32 + nL + 1, stream, where); + m_PhiBinnerStore = make_buffer(1,stream,where); } static_assert(sizeof(TrackingRecHit2DSOAView::hindex_type) == sizeof(float)); @@ -199,7 +200,7 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous( // transfer view if constexpr (memoryPool::onDevice == where) { - cudaCheck(cudaMemcpyAsync(m_view.get(), view.get(), sizeof(TrackingRecHit2DSOAView), cudaMemcpyHostToDevice, stream)); + cudaCheck(memoryPool::cuda::copy(m_view, view, sizeof(TrackingRecHit2DSOAView), stream)); } else { m_view.reset(view.release()); // NOLINT: std::move() breaks CUDA version } @@ -211,25 +212,26 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous( float* store32, uint16_t* store16, uint32_t* modules, int nHits, cudaStream_t stream) : m_nHits(nHits), m_hitsModuleStart(modules) { - auto view = make_unique(1,stream,memoryPool::onCPU==where? memoryPool::onCPU : memoryPool::onHost); + using namespace memoryPool::cuda; + auto view = make_buffer(1,stream,memoryPool::onCPU==where? memoryPool::onCPU : memoryPool::onHost); - m_view = make_unique(stream); + m_view = make_buffer(1,stream,where); view->m_nHits = nHits; if (0 == nHits) { if constexpr (memoryPool::onDevice == where) { - cudaCheck(cudaMemcpyAsync(m_view.get(), view.get(), sizeof(TrackingRecHit2DSOAView), cudaMemcpyHostToDevice, stream)); + cudaCheck(memoryPool::cuda::copy(m_view, view, sizeof(TrackingRecHit2DSOAView), stream)); } else { m_view = std::move(view); } return; } - m_store16 = make_unique(nHits * n16, stream,where); - m_store32 = make_unique(nHits * n32, stream,where); - m_PhiBinnerStore = make_unique(1,stream,where); - m_AverageGeometryStore = make_unique(1,stream,where); + m_store16 = make_buffer(nHits * n16, stream,where); + m_store32 = make_buffer(nHits * n32, stream,where); + m_PhiBinnerStore = make_buffer(1,stream,where); + m_AverageGeometryStore = make_buffer(1,stream,where); view->m_averageGeometry = m_AverageGeometryStore.get(); view->m_hitsModuleStart = m_hitsModuleStart; @@ -270,7 +272,7 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous( // transfer view if constexpr (memoryPool::onDevice == where) { - cudaCheck(cudaMemcpyAsync(m_view.get(), view.get(), sizeof(TrackingRecHit2DSOAView), cudaMemcpyHostToDevice, stream)); + cudaCheck(memoryPool::cuda::copy(m_view, view, sizeof(TrackingRecHit2DSOAView), stream)); } else { m_view = std::move(view); } diff --git a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc index e9bef7c4eaf89..70f17036c8146 100644 --- a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc +++ b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc @@ -5,32 +5,32 @@ using namespace memoryPool::cuda; template <> -memoryPool::unique_ptr TrackingRecHit2DGPU::localCoordToHostAsync(cudaStream_t stream) const { - auto ret = make_unique(5 * nHits(), stream,onHost); +memoryPool::buffer TrackingRecHit2DGPU::localCoordToHostAsync(cudaStream_t stream) const { + auto ret = make_buffer(5 * nHits(), stream,memoryPool::onHost); cudaCheck( cudaMemcpyAsync(ret.get(), m_store32.get(), 5 * nHits(), cudaMemcpyDefault,stream)); return ret; } template <> -memoryPool::unique_ptr TrackingRecHit2DGPU::store32ToHostAsync(cudaStream_t stream) const { - auto ret = make_unique(static_cast(n32) * nHits(), stream,onHost)); +memoryPool::buffer TrackingRecHit2DGPU::store32ToHostAsync(cudaStream_t stream) const { + auto ret = make_buffer(static_cast(n32) * nHits(), stream,memoryPool::onHost); cudaCheck( cudaMemcpyAsync(ret.get(), m_store32.get(), static_cast(n32) * nHits(), cudaMemcpyDefault,stream)); return ret; } template <> -memoryPool::unique_ptr TrackingRecHit2DGPU::store16ToHostAsync(cudaStream_t stream) const { - auto ret = make_unique(static_cast(n16) * nHits(), stream,onHost)); +memoryPool::buffer TrackingRecHit2DGPU::store16ToHostAsync(cudaStream_t stream) const { + auto ret = make_buffer(static_cast(n16) * nHits(), stream,memoryPool::onHost); cudaCheck( - cudaMemcpyAsync(ret.get(), m_store16, static_cast(n16) * nHits(), cudaMemcpyDefault,stream)); + cudaMemcpyAsync(ret.get(), m_store16.get(), static_cast(n16) * nHits(), cudaMemcpyDefault,stream)); return ret; } template <> -memoryPool::unique_ptr TrackingRecHit2DGPU::hitsModuleStartToHostAsync(cudaStream_t stream) const { - auto ret = make_unique(nMaxModules() + 1, stream,onHost)); +memoryPool::buffer TrackingRecHit2DGPU::hitsModuleStartToHostAsync(cudaStream_t stream) const { + auto ret = make_buffer(nMaxModules() + 1, stream,memoryPool::onHost); cudaCheck( cudaMemcpyAsync(ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (nMaxModules() + 1), cudaMemcpyDefault, stream)); return ret;