SiPixelRecHits/plugins/SiPixelRecHitSoAFromCUDA.cc

0001 #include <cuda_runtime.h>
0002
0003 #include <fmt/printf.h>
0004
0005 #include "CUDADataFormats/Common/interface/HostProduct.h"
0006 #include "CUDADataFormats/Common/interface/Product.h"
0007 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h"
0008 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h"
0009 #include "DataFormats/Common/interface/DetSetVectorNew.h"
0010 #include "DataFormats/Common/interface/Handle.h"
0011 #include "DataFormats/SiPixelCluster/interface/SiPixelCluster.h"
0012 #include "DataFormats/TrackerRecHit2D/interface/SiPixelRecHitCollection.h"
0013 #include "FWCore/Framework/interface/Event.h"
0014 #include "FWCore/Framework/interface/EventSetup.h"
0015 #include "FWCore/Framework/interface/MakerMacros.h"
0016 #include "FWCore/Framework/interface/stream/EDProducer.h"
0017 #include "FWCore/MessageLogger/interface/MessageLogger.h"
0018 #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
0019 #include "FWCore/ParameterSet/interface/ParameterSet.h"
0020 #include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
0021 #include "FWCore/Utilities/interface/InputTag.h"
0022 #include "Geometry/CommonDetUnit/interface/PixelGeomDetUnit.h"
0023 #include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
0024 #include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
0025 #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
0026 #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
0027
0028 template <typename TrackerTraits>
0029 class SiPixelRecHitSoAFromCUDAT : public edm::stream::EDProducer<edm::ExternalWork> {
0030 public:
0031   explicit SiPixelRecHitSoAFromCUDAT(const edm::ParameterSet& iConfig);
0032   ~SiPixelRecHitSoAFromCUDAT() override = default;
0033
0034   static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
0035   using HMSstorage = HostProduct<uint32_t[]>;
0036   using HitsOnHost = TrackingRecHitSoAHost<TrackerTraits>;
0037   using HitsOnDevice = TrackingRecHitSoADevice<TrackerTraits>;
0038
0039 private:
0040   void acquire(edm::Event const& iEvent,
0041                edm::EventSetup const& iSetup,
0042                edm::WaitingTaskWithArenaHolder waitingTaskHolder) override;
0043   void produce(edm::Event& iEvent, edm::EventSetup const& iSetup) override;
0044
0045   const edm::EDGetTokenT<cms::cuda::Product<HitsOnDevice>> hitsTokenGPU_;  // CUDA hits
0046   const edm::EDPutTokenT<HitsOnHost> hitsPutTokenCPU_;
0047   const edm::EDPutTokenT<HMSstorage> hostPutToken_;
0048
0049   uint32_t nHits_;
0050   HitsOnHost hits_h_;
0051 };
0052
0053 template <typename TrackerTraits>
0054 SiPixelRecHitSoAFromCUDAT<TrackerTraits>::SiPixelRecHitSoAFromCUDAT(const edm::ParameterSet& iConfig)
0055     : hitsTokenGPU_(consumes(iConfig.getParameter<edm::InputTag>("pixelRecHitSrc"))),
0056       hitsPutTokenCPU_(produces<HitsOnHost>()),
0057       hostPutToken_(produces<HMSstorage>()) {}
0058
0059 template <typename TrackerTraits>
0060 void SiPixelRecHitSoAFromCUDAT<TrackerTraits>::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
0061   edm::ParameterSetDescription desc;
0062   desc.add<edm::InputTag>("pixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA"));
0063
0064   descriptions.addWithDefaultLabel(desc);
0065 }
0066
0067 template <typename TrackerTraits>
0068 void SiPixelRecHitSoAFromCUDAT<TrackerTraits>::acquire(edm::Event const& iEvent,
0069                                                        edm::EventSetup const& iSetup,
0070                                                        edm::WaitingTaskWithArenaHolder waitingTaskHolder) {
0071   cms::cuda::Product<HitsOnDevice> const& inputDataWrapped = iEvent.get(hitsTokenGPU_);
0072   cms::cuda::ScopedContextAcquire ctx{inputDataWrapped, std::move(waitingTaskHolder)};
0073   auto const& inputData = ctx.get(inputDataWrapped);
0074
0075   nHits_ = inputData.nHits();
0076   hits_h_ = HitsOnHost(nHits_, ctx.stream());
0077   cudaCheck(cudaMemcpyAsync(hits_h_.buffer().get(),
0078                             inputData.const_buffer().get(),
0079                             inputData.bufferSize(),
0080                             cudaMemcpyDeviceToHost,
0081                             ctx.stream()));  // Copy data from Device to Host
0082   LogDebug("SiPixelRecHitSoAFromCUDA") << "copying to cpu SoA" << inputData.nHits() << " Hits";
0083 }
0084
0085 template <typename TrackerTraits>
0086 void SiPixelRecHitSoAFromCUDAT<TrackerTraits>::produce(edm::Event& iEvent, edm::EventSetup const& es) {
0087   auto hmsp = std::make_unique<uint32_t[]>(TrackerTraits::numberOfModules + 1);
0088
0089   if (nHits_ > 0)
0090     std::copy(hits_h_.view().hitsModuleStart().begin(), hits_h_.view().hitsModuleStart().end(), hmsp.get());
0091
0092   iEvent.emplace(hostPutToken_, std::move(hmsp));
0093   iEvent.emplace(hitsPutTokenCPU_, std::move(hits_h_));
0094 }
0095
0096 using SiPixelRecHitSoAFromCUDAPhase1 = SiPixelRecHitSoAFromCUDAT<pixelTopology::Phase1>;
0097 DEFINE_FWK_MODULE(SiPixelRecHitSoAFromCUDAPhase1);
0098
0099 using SiPixelRecHitSoAFromCUDAPhase2 = SiPixelRecHitSoAFromCUDAT<pixelTopology::Phase2>;
0100 DEFINE_FWK_MODULE(SiPixelRecHitSoAFromCUDAPhase2);
0101
0102 using SiPixelRecHitSoAFromCUDAHIonPhase1 = SiPixelRecHitSoAFromCUDAT<pixelTopology::HIonPhase1>;
0103 DEFINE_FWK_MODULE(SiPixelRecHitSoAFromCUDAHIonPhase1);