From 3eafe4f6d9578d2ab0f75870798d84f40cb3437d Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Tue, 23 Apr 2019 08:18:43 -0500 Subject: [PATCH 1/4] Produce a copy of the beamspot data on the GPU Implement the beamspot host-to-device transfer in its own EDProducer, making use of beginStream()-allocated write-combined memory for the transfer. --- CUDADataFormats/BeamSpot/BuildFile.xml | 8 ++ .../BeamSpot/interface/BeamSpotCUDA.h | 32 +++++++ CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc | 9 ++ CUDADataFormats/BeamSpot/src/classes.h | 8 ++ CUDADataFormats/BeamSpot/src/classes_def.xml | 4 + .../python/Reconstruction_cff.py | 6 +- .../plugins/BeamSpotToCUDA.cc | 83 +++++++++++++++++++ .../BeamSpotProducer/plugins/BuildFile.xml | 21 +++-- .../BeamSpotProducer/python/BeamSpot_cff.py | 7 ++ 9 files changed, 167 insertions(+), 11 deletions(-) create mode 100644 CUDADataFormats/BeamSpot/BuildFile.xml create mode 100644 CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h create mode 100644 CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc create mode 100644 CUDADataFormats/BeamSpot/src/classes.h create mode 100644 CUDADataFormats/BeamSpot/src/classes_def.xml create mode 100644 RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc diff --git a/CUDADataFormats/BeamSpot/BuildFile.xml b/CUDADataFormats/BeamSpot/BuildFile.xml new file mode 100644 index 0000000000000..75f3d15738429 --- /dev/null +++ b/CUDADataFormats/BeamSpot/BuildFile.xml @@ -0,0 +1,8 @@ + + + + + + + + diff --git a/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h b/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h new file mode 100644 index 0000000000000..800634d2f5270 --- /dev/null +++ b/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h @@ -0,0 +1,32 @@ +#ifndef CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h +#define CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h + +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" + +#include + +class BeamSpotCUDA { +public: + // alignas(128) doesn't really make sense as there is only one + // beamspot per event? + struct Data { + float x, y, z; // position + // TODO: add covariance matrix + + float sigmaZ; + float beamWidthX, beamWidthY; + float dxdz, dydz; + float emittanceX, emittanceY; + float betaStar; + }; + + BeamSpotCUDA() = default; + BeamSpotCUDA(Data const* data_h, cudaStream_t stream); + + Data const* data() const { return data_d_.get(); } + +private: + cms::cuda::device::unique_ptr data_d_; +}; + +#endif diff --git a/CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc b/CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc new file mode 100644 index 0000000000000..575fcf63b8eaa --- /dev/null +++ b/CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc @@ -0,0 +1,9 @@ +#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" + +#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" + +BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cudaStream_t stream) { + data_d_ = cms::cuda::make_device_unique(stream); + cudaCheck(cudaMemcpyAsync(data_d_.get(), data_h, sizeof(Data), cudaMemcpyHostToDevice, stream)); +} diff --git a/CUDADataFormats/BeamSpot/src/classes.h b/CUDADataFormats/BeamSpot/src/classes.h new file mode 100644 index 0000000000000..f79c8c9452c63 --- /dev/null +++ b/CUDADataFormats/BeamSpot/src/classes.h @@ -0,0 +1,8 @@ +#ifndef CUDADataFormats_BeamSpot_classes_h +#define CUDADataFormats_BeamSpot_classes_h + +#include "CUDADataFormats/Common/interface/Product.h" +#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" +#include "DataFormats/Common/interface/Wrapper.h" + +#endif diff --git a/CUDADataFormats/BeamSpot/src/classes_def.xml b/CUDADataFormats/BeamSpot/src/classes_def.xml new file mode 100644 index 0000000000000..198edeebe7c73 --- /dev/null +++ b/CUDADataFormats/BeamSpot/src/classes_def.xml @@ -0,0 +1,4 @@ + + + + diff --git a/Configuration/StandardSequences/python/Reconstruction_cff.py b/Configuration/StandardSequences/python/Reconstruction_cff.py index edf85287421b3..24776960bc877 100644 --- a/Configuration/StandardSequences/python/Reconstruction_cff.py +++ b/Configuration/StandardSequences/python/Reconstruction_cff.py @@ -103,7 +103,7 @@ from RecoLocalCalo.Castor.Castor_cff import * from RecoLocalCalo.Configuration.hcalGlobalReco_cff import * -globalreco_trackingTask = cms.Task(offlineBeamSpot, +globalreco_trackingTask = cms.Task(offlineBeamSpotTask, MeasurementTrackerEventPreSplitting, # unclear where to put this siPixelClusterShapeCachePreSplitting, # unclear where to put this standalonemuontrackingTask, @@ -117,7 +117,7 @@ ########################################## # offlineBeamSpot is reconstructed before mixing in fastSim ########################################## -_fastSim_globalreco_trackingTask = globalreco_trackingTask.copyAndExclude([offlineBeamSpot,MeasurementTrackerEventPreSplitting,siPixelClusterShapeCachePreSplitting]) +_fastSim_globalreco_trackingTask = globalreco_trackingTask.copyAndExclude([offlineBeamSpotTask,MeasurementTrackerEventPreSplitting,siPixelClusterShapeCachePreSplitting]) fastSim.toReplaceWith(globalreco_trackingTask,_fastSim_globalreco_trackingTask) _phase2_timing_layer_globalreco_trackingTask = globalreco_trackingTask.copy() @@ -212,7 +212,7 @@ reconstruction_trackingOnly = cms.Sequence(reconstruction_trackingOnlyTask) reconstruction_pixelTrackingOnlyTask = cms.Task( pixeltrackerlocalrecoTask, - offlineBeamSpot, + offlineBeamSpotTask, siPixelClusterShapeCachePreSplitting, recopixelvertexingTask ) diff --git a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc new file mode 100644 index 0000000000000..d9f8490ba54e0 --- /dev/null +++ b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc @@ -0,0 +1,83 @@ +#include "CUDADataFormats/Common/interface/Product.h" +#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" +#include "DataFormats/BeamSpot/interface/BeamSpot.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/global/EDProducer.h" +#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" +#include "FWCore/ParameterSet/interface/ParameterSetDescription.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/ServiceRegistry/interface/Service.h" +#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" +#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h" +#include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" + +#include + +namespace { + class BSHost { + public: + BSHost() : bs{cms::cuda::make_host_noncached_unique(cudaHostAllocWriteCombined)} {} + BeamSpotCUDA::Data* get() { return bs.get(); } + + private: + cms::cuda::host::noncached::unique_ptr bs; + }; +} // namespace + +class BeamSpotToCUDA : public edm::global::EDProducer> { +public: + explicit BeamSpotToCUDA(const edm::ParameterSet& iConfig); + ~BeamSpotToCUDA() override = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + + std::unique_ptr beginStream(edm::StreamID) const override { + edm::Service cs; + if (cs->enabled()) { + return std::make_unique(); + } else { + return nullptr; + } + } + void produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const override; + +private: + edm::EDGetTokenT bsGetToken_; + edm::EDPutTokenT> bsPutToken_; +}; + +BeamSpotToCUDA::BeamSpotToCUDA(const edm::ParameterSet& iConfig) + : bsGetToken_{consumes(iConfig.getParameter("src"))}, + bsPutToken_{produces>()} {} + +void BeamSpotToCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("src", edm::InputTag("offlineBeamSpot")); + descriptions.add("offlineBeamSpotCUDA", desc); +} + +void BeamSpotToCUDA::produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const { + cms::cuda::ScopedContextProduce ctx{streamID}; + + const reco::BeamSpot& bs = iEvent.get(bsGetToken_); + + BeamSpotCUDA::Data* bsHost = streamCache(streamID)->get(); + + bsHost->x = bs.x0(); + bsHost->y = bs.y0(); + bsHost->z = bs.z0(); + + bsHost->sigmaZ = bs.sigmaZ(); + bsHost->beamWidthX = bs.BeamWidthX(); + bsHost->beamWidthY = bs.BeamWidthY(); + bsHost->dxdz = bs.dxdz(); + bsHost->dydz = bs.dydz(); + bsHost->emittanceX = bs.emittanceX(); + bsHost->emittanceY = bs.emittanceY(); + bsHost->betaStar = bs.betaStar(); + + ctx.emplace(iEvent, bsPutToken_, bsHost, ctx.stream()); +} + +DEFINE_FWK_MODULE(BeamSpotToCUDA); diff --git a/RecoVertex/BeamSpotProducer/plugins/BuildFile.xml b/RecoVertex/BeamSpotProducer/plugins/BuildFile.xml index c7498da5ae33f..97d6d07a27631 100644 --- a/RecoVertex/BeamSpotProducer/plugins/BuildFile.xml +++ b/RecoVertex/BeamSpotProducer/plugins/BuildFile.xml @@ -1,13 +1,12 @@ + + + + + - - - - - - - + @@ -40,4 +39,10 @@ - + + + + + + + diff --git a/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py b/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py index deb62255199e5..a6264a322c808 100644 --- a/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py +++ b/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py @@ -1,4 +1,11 @@ import FWCore.ParameterSet.Config as cms from RecoVertex.BeamSpotProducer.BeamSpot_cfi import * +from RecoVertex.BeamSpotProducer.offlineBeamSpotCUDA_cfi import offlineBeamSpotCUDA +offlineBeamSpotTask = cms.Task(offlineBeamSpot) + +from Configuration.ProcessModifiers.gpu_cff import gpu +_offlineBeamSpotTask_gpu = offlineBeamSpotTask.copy() +_offlineBeamSpotTask_gpu.add(offlineBeamSpotCUDA) +gpu.toReplaceWith(offlineBeamSpotTask, _offlineBeamSpotTask_gpu) From a946344e90076148dff072825874d6f8b11111d8 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Fri, 7 Aug 2020 20:13:26 +0200 Subject: [PATCH 2/4] Update the BeamSpotCUDA class Make the BeamSpotCUDA movable and explicitly non-copiable (as was already the case due to the device::unique_ptr data member). Remove the cudaMemcpyAsync from the BeamSpotCUDA data format, and move it to the BeamSpotToCUDA producer. --- .../BeamSpot/interface/BeamSpotCUDA.h | 26 +++++++--- CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc | 9 ---- CUDADataFormats/BeamSpot/src/classes.h | 2 +- .../plugins/BeamSpotToCUDA.cc | 47 +++++++++++++------ 4 files changed, 52 insertions(+), 32 deletions(-) delete mode 100644 CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc diff --git a/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h b/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h index 800634d2f5270..d23e5be4a87e4 100644 --- a/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h +++ b/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h @@ -1,15 +1,14 @@ #ifndef CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h #define CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h -#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" - #include +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" + class BeamSpotCUDA { public: - // alignas(128) doesn't really make sense as there is only one - // beamspot per event? - struct Data { + // align to the CUDA L1 cache line size + struct alignas(128) Data { float x, y, z; // position // TODO: add covariance matrix @@ -20,13 +19,26 @@ class BeamSpotCUDA { float betaStar; }; + // default constructor, required by cms::cuda::Product BeamSpotCUDA() = default; - BeamSpotCUDA(Data const* data_h, cudaStream_t stream); + // constructor that allocates cached device memory on the given CUDA stream + BeamSpotCUDA(cudaStream_t stream) { data_d_ = cms::cuda::make_device_unique(stream); } + + // movable, non-copiable + BeamSpotCUDA(BeamSpotCUDA const&) = delete; + BeamSpotCUDA(BeamSpotCUDA&&) = default; + BeamSpotCUDA& operator=(BeamSpotCUDA const&) = delete; + BeamSpotCUDA& operator=(BeamSpotCUDA&&) = default; + + Data* data() { return data_d_.get(); } Data const* data() const { return data_d_.get(); } + cms::cuda::device::unique_ptr& ptr() { return data_d_; } + cms::cuda::device::unique_ptr const& ptr() const { return data_d_; } + private: cms::cuda::device::unique_ptr data_d_; }; -#endif +#endif // CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h diff --git a/CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc b/CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc deleted file mode 100644 index 575fcf63b8eaa..0000000000000 --- a/CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc +++ /dev/null @@ -1,9 +0,0 @@ -#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" - -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" - -BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cudaStream_t stream) { - data_d_ = cms::cuda::make_device_unique(stream); - cudaCheck(cudaMemcpyAsync(data_d_.get(), data_h, sizeof(Data), cudaMemcpyHostToDevice, stream)); -} diff --git a/CUDADataFormats/BeamSpot/src/classes.h b/CUDADataFormats/BeamSpot/src/classes.h index f79c8c9452c63..5aebe536f8a01 100644 --- a/CUDADataFormats/BeamSpot/src/classes.h +++ b/CUDADataFormats/BeamSpot/src/classes.h @@ -5,4 +5,4 @@ #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" #include "DataFormats/Common/interface/Wrapper.h" -#endif +#endif // CUDADataFormats_BeamSpot_classes_h diff --git a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc index d9f8490ba54e0..9c6f57b8132b5 100644 --- a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc +++ b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc @@ -1,41 +1,55 @@ -#include "CUDADataFormats/Common/interface/Product.h" +#include + #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" +#include "CUDADataFormats/Common/interface/Product.h" #include "DataFormats/BeamSpot/interface/BeamSpot.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "FWCore/Framework/interface/global/EDProducer.h" #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" -#include "FWCore/ParameterSet/interface/ParameterSetDescription.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/ParameterSet/interface/ParameterSetDescription.h" #include "FWCore/ServiceRegistry/interface/Service.h" #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" -#include - namespace { - class BSHost { + + class BeamSpotHost { public: - BSHost() : bs{cms::cuda::make_host_noncached_unique(cudaHostAllocWriteCombined)} {} - BeamSpotCUDA::Data* get() { return bs.get(); } + BeamSpotHost() : data_h_{cms::cuda::make_host_noncached_unique(cudaHostAllocWriteCombined)} {} + + BeamSpotHost(BeamSpotHost const&) = delete; + BeamSpotHost(BeamSpotHost&&) = default; + + BeamSpotHost& operator=(BeamSpotHost const&) = delete; + BeamSpotHost& operator=(BeamSpotHost&&) = default; + + BeamSpotCUDA::Data* data() { return data_h_.get(); } + BeamSpotCUDA::Data const* data() const { return data_h_.get(); } + + cms::cuda::host::noncached::unique_ptr& ptr() { return data_h_; } + cms::cuda::host::noncached::unique_ptr const& ptr() const { return data_h_; } private: - cms::cuda::host::noncached::unique_ptr bs; + cms::cuda::host::noncached::unique_ptr data_h_; }; + } // namespace -class BeamSpotToCUDA : public edm::global::EDProducer> { +class BeamSpotToCUDA : public edm::global::EDProducer> { public: explicit BeamSpotToCUDA(const edm::ParameterSet& iConfig); ~BeamSpotToCUDA() override = default; static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); - std::unique_ptr beginStream(edm::StreamID) const override { + std::unique_ptr beginStream(edm::StreamID) const override { edm::Service cs; if (cs->enabled()) { - return std::make_unique(); + return std::make_unique(); } else { return nullptr; } @@ -43,8 +57,8 @@ class BeamSpotToCUDA : public edm::global::EDProducer> void produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const override; private: - edm::EDGetTokenT bsGetToken_; - edm::EDPutTokenT> bsPutToken_; + const edm::EDGetTokenT bsGetToken_; + const edm::EDPutTokenT> bsPutToken_; }; BeamSpotToCUDA::BeamSpotToCUDA(const edm::ParameterSet& iConfig) @@ -62,7 +76,7 @@ void BeamSpotToCUDA::produce(edm::StreamID streamID, edm::Event& iEvent, const e const reco::BeamSpot& bs = iEvent.get(bsGetToken_); - BeamSpotCUDA::Data* bsHost = streamCache(streamID)->get(); + auto& bsHost = streamCache(streamID)->ptr(); bsHost->x = bs.x0(); bsHost->y = bs.y0(); @@ -77,7 +91,10 @@ void BeamSpotToCUDA::produce(edm::StreamID streamID, edm::Event& iEvent, const e bsHost->emittanceY = bs.emittanceY(); bsHost->betaStar = bs.betaStar(); - ctx.emplace(iEvent, bsPutToken_, bsHost, ctx.stream()); + BeamSpotCUDA bsDevice(ctx.stream()); + cms::cuda::copyAsync(bsDevice.ptr(), bsHost, ctx.stream()); + + ctx.emplace(iEvent, bsPutToken_, std::move(bsDevice)); } DEFINE_FWK_MODULE(BeamSpotToCUDA); From 51f18ad645011763db0eccde292a389abde0b7af Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Fri, 21 Aug 2020 09:22:34 +0200 Subject: [PATCH 3/4] Rename the default module name to offlineBeamSpotToCUDA --- RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc | 2 +- RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc index 9c6f57b8132b5..3ca01c4e827b2 100644 --- a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc +++ b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc @@ -68,7 +68,7 @@ BeamSpotToCUDA::BeamSpotToCUDA(const edm::ParameterSet& iConfig) void BeamSpotToCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; desc.add("src", edm::InputTag("offlineBeamSpot")); - descriptions.add("offlineBeamSpotCUDA", desc); + descriptions.add("offlineBeamSpotToCUDA", desc); } void BeamSpotToCUDA::produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const { diff --git a/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py b/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py index a6264a322c808..7cc651af22106 100644 --- a/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py +++ b/RecoVertex/BeamSpotProducer/python/BeamSpot_cff.py @@ -1,11 +1,11 @@ import FWCore.ParameterSet.Config as cms from RecoVertex.BeamSpotProducer.BeamSpot_cfi import * -from RecoVertex.BeamSpotProducer.offlineBeamSpotCUDA_cfi import offlineBeamSpotCUDA +from RecoVertex.BeamSpotProducer.offlineBeamSpotToCUDA_cfi import offlineBeamSpotToCUDA offlineBeamSpotTask = cms.Task(offlineBeamSpot) from Configuration.ProcessModifiers.gpu_cff import gpu _offlineBeamSpotTask_gpu = offlineBeamSpotTask.copy() -_offlineBeamSpotTask_gpu.add(offlineBeamSpotCUDA) +_offlineBeamSpotTask_gpu.add(offlineBeamSpotToCUDA) gpu.toReplaceWith(offlineBeamSpotTask, _offlineBeamSpotTask_gpu) From d2dfd25f882cc4a1ab25c2bcfbb35363b2964ef9 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Fri, 21 Aug 2020 09:23:22 +0200 Subject: [PATCH 4/4] Move the underlying data type BeamSpotCUDA::Data to a separate transient struct BeamSpotPOD --- .../BeamSpot/interface/BeamSpotCUDA.h | 25 ++++++------------- DataFormats/BeamSpot/interface/BeamSpotPOD.h | 20 +++++++++++++++ DataFormats/BeamSpot/src/classes.h | 9 ++----- DataFormats/BeamSpot/src/classes_def.xml | 3 +++ .../plugins/BeamSpotToCUDA.cc | 13 +++++----- 5 files changed, 39 insertions(+), 31 deletions(-) create mode 100644 DataFormats/BeamSpot/interface/BeamSpotPOD.h diff --git a/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h b/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h index d23e5be4a87e4..7b04fac67b9f1 100644 --- a/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h +++ b/CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h @@ -3,27 +3,16 @@ #include +#include "DataFormats/BeamSpot/interface/BeamSpotPOD.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" class BeamSpotCUDA { public: - // align to the CUDA L1 cache line size - struct alignas(128) Data { - float x, y, z; // position - // TODO: add covariance matrix - - float sigmaZ; - float beamWidthX, beamWidthY; - float dxdz, dydz; - float emittanceX, emittanceY; - float betaStar; - }; - // default constructor, required by cms::cuda::Product BeamSpotCUDA() = default; // constructor that allocates cached device memory on the given CUDA stream - BeamSpotCUDA(cudaStream_t stream) { data_d_ = cms::cuda::make_device_unique(stream); } + BeamSpotCUDA(cudaStream_t stream) { data_d_ = cms::cuda::make_device_unique(stream); } // movable, non-copiable BeamSpotCUDA(BeamSpotCUDA const&) = delete; @@ -31,14 +20,14 @@ class BeamSpotCUDA { BeamSpotCUDA& operator=(BeamSpotCUDA const&) = delete; BeamSpotCUDA& operator=(BeamSpotCUDA&&) = default; - Data* data() { return data_d_.get(); } - Data const* data() const { return data_d_.get(); } + BeamSpotPOD* data() { return data_d_.get(); } + BeamSpotPOD const* data() const { return data_d_.get(); } - cms::cuda::device::unique_ptr& ptr() { return data_d_; } - cms::cuda::device::unique_ptr const& ptr() const { return data_d_; } + cms::cuda::device::unique_ptr& ptr() { return data_d_; } + cms::cuda::device::unique_ptr const& ptr() const { return data_d_; } private: - cms::cuda::device::unique_ptr data_d_; + cms::cuda::device::unique_ptr data_d_; }; #endif // CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h diff --git a/DataFormats/BeamSpot/interface/BeamSpotPOD.h b/DataFormats/BeamSpot/interface/BeamSpotPOD.h new file mode 100644 index 0000000000000..5920db6a18930 --- /dev/null +++ b/DataFormats/BeamSpot/interface/BeamSpotPOD.h @@ -0,0 +1,20 @@ +#ifndef DataFormats_BeamSpot_interface_BeamSpotPOD_h +#define DataFormats_BeamSpot_interface_BeamSpotPOD_h + +// This struct is a transient-only, simplified representation of the beamspot +// data used as the underlying type for data transfers and operations in +// heterogeneous code (e.g. in CUDA code). + +// The covariance matrix is not used in that code, so is left out here. + +// align to the CUDA L1 cache line size +struct alignas(128) BeamSpotPOD { + float x, y, z; // position + float sigmaZ; + float beamWidthX, beamWidthY; + float dxdz, dydz; + float emittanceX, emittanceY; + float betaStar; +}; + +#endif // DataFormats_BeamSpot_interface_BeamSpotPOD_h diff --git a/DataFormats/BeamSpot/src/classes.h b/DataFormats/BeamSpot/src/classes.h index efe580a4cc8f5..2365c90fdc257 100644 --- a/DataFormats/BeamSpot/src/classes.h +++ b/DataFormats/BeamSpot/src/classes.h @@ -1,8 +1,3 @@ - -#include "DataFormats/Common/interface/Wrapper.h" #include "DataFormats/BeamSpot/interface/BeamSpot.h" -#include "Math/Cartesian3D.h" -#include "DataFormats/Math/interface/Vector3D.h" -#include "Math/Polar3D.h" -#include "Math/CylindricalEta3D.h" -#include "DataFormats/Math/interface/Vector.h" +#include "DataFormats/BeamSpot/interface/BeamSpotPOD.h" +#include "DataFormats/Common/interface/Wrapper.h" diff --git a/DataFormats/BeamSpot/src/classes_def.xml b/DataFormats/BeamSpot/src/classes_def.xml index 9c99903ae8bd4..af3d4e4627c58 100644 --- a/DataFormats/BeamSpot/src/classes_def.xml +++ b/DataFormats/BeamSpot/src/classes_def.xml @@ -4,4 +4,7 @@ + + + diff --git a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc index 3ca01c4e827b2..8b0de1c739076 100644 --- a/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc +++ b/RecoVertex/BeamSpotProducer/plugins/BeamSpotToCUDA.cc @@ -3,6 +3,7 @@ #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" #include "CUDADataFormats/Common/interface/Product.h" #include "DataFormats/BeamSpot/interface/BeamSpot.h" +#include "DataFormats/BeamSpot/interface/BeamSpotPOD.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "FWCore/Framework/interface/global/EDProducer.h" @@ -19,7 +20,7 @@ namespace { class BeamSpotHost { public: - BeamSpotHost() : data_h_{cms::cuda::make_host_noncached_unique(cudaHostAllocWriteCombined)} {} + BeamSpotHost() : data_h_{cms::cuda::make_host_noncached_unique(cudaHostAllocWriteCombined)} {} BeamSpotHost(BeamSpotHost const&) = delete; BeamSpotHost(BeamSpotHost&&) = default; @@ -27,14 +28,14 @@ namespace { BeamSpotHost& operator=(BeamSpotHost const&) = delete; BeamSpotHost& operator=(BeamSpotHost&&) = default; - BeamSpotCUDA::Data* data() { return data_h_.get(); } - BeamSpotCUDA::Data const* data() const { return data_h_.get(); } + BeamSpotPOD* data() { return data_h_.get(); } + BeamSpotPOD const* data() const { return data_h_.get(); } - cms::cuda::host::noncached::unique_ptr& ptr() { return data_h_; } - cms::cuda::host::noncached::unique_ptr const& ptr() const { return data_h_; } + cms::cuda::host::noncached::unique_ptr& ptr() { return data_h_; } + cms::cuda::host::noncached::unique_ptr const& ptr() const { return data_h_; } private: - cms::cuda::host::noncached::unique_ptr data_h_; + cms::cuda::host::noncached::unique_ptr data_h_; }; } // namespace