-
Notifications
You must be signed in to change notification settings - Fork 4.3k
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
ECAL Phase 2 weights method amplitude reconstruction on GPU #37695
Merged
cmsbuild
merged 23 commits into
cms-sw:master
from
ChrisSandever:ChrisSPhase2EcalRecoGPU
Aug 4, 2022
Merged
Changes from all commits
Commits
Show all changes
23 commits
Select commit
Hold shift + click to select a range
7979baa
Added amplitudeError
ChrisSandever 2cf6995
Added GPU weights reconstruction module for Phase2EcalRecoGPU
ChrisSandever c9ab238
Modified EcalCPUUncalibRecHitProducer to be Phase 1 & 2 compatible
ChrisSandever 491e688
Modified EcalUncalibRecHitConvertGPU2CPUFormat to be Phase 1 & 2 comp…
ChrisSandever 5fd429d
Added configuration for testing Phase2EcalRecoGPU
ChrisSandever e35ae4d
Modified ecalUncalibRecHitPhase2_cff to use a switch producer
ChrisSandever b6eb69a
Added test configuration for Phase2EcalRecoGPU
ChrisSandever 724935f
Added isPhase1 for amplitudeError issue
ChrisSandever d8c9212
re-added testEcalUncalibRecHitPhase2WeightsProducerGPU_harvesting.py
ChrisSandever e9f49d8
Changed isPhase1 to not(isPhase2)
ChrisSandever 5a3e720
Added DeclsForKernelsPhase2.h
ChrisSandever 67d3466
Changed converter to use switch description
ChrisSandever f9eddde
Added Phase2 CPU to GPU converter
ChrisSandever b59b8c1
Changed converter to use switch description
ChrisSandever 487051f
Code checks
ChrisSandever c51e579
Changed Ph2 to Phase2
ChrisSandever fcd98f8
Changed Ph2 to Phase2
ChrisSandever 1dcac8f
Changed Ph2 to Phase2
ChrisSandever 722ef69
Changed weights_ to use cms::cuda::HostAllocator
ChrisSandever b1ba084
Chanhed Ph2 to Phase2 and removes EE objects
ChrisSandever e48b02d
Code checks
ChrisSandever be0268c
Made else statement in converter
ChrisSandever 0820f19
Reverted kernel changes and added comments
ChrisSandever File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
28 changes: 28 additions & 0 deletions
28
RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernelsPhase2.h
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,28 @@ | ||
| #ifndef RecoLocalCalo_EcalRecProducers_plugins_DeclsForKernelsPhase2_h | ||
| #define RecoLocalCalo_EcalRecProducers_plugins_DeclsForKernelsPhase2_h | ||
|
|
||
| #include "CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h" | ||
| #include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" | ||
|
|
||
| namespace ecal { | ||
| namespace weights { | ||
|
|
||
| struct EventOutputDataGPU { | ||
| UncalibratedRecHit<::calo::common::DevStoragePolicy> recHits; | ||
|
|
||
| void allocate(uint32_t digi_size, cudaStream_t cudaStream) { | ||
| auto const size = digi_size; | ||
| recHits.amplitudesAll = | ||
| cms::cuda::make_device_unique<reco::ComputationScalarType[]>(size * EcalDataFrame::MAXSAMPLES, cudaStream); | ||
| recHits.amplitude = cms::cuda::make_device_unique<reco::StorageScalarType[]>(size, cudaStream); | ||
| recHits.amplitudeError = cms::cuda::make_device_unique<reco::StorageScalarType[]>(size, cudaStream); | ||
| recHits.chi2 = cms::cuda::make_device_unique<reco::StorageScalarType[]>(size, cudaStream); | ||
| recHits.pedestal = cms::cuda::make_device_unique<reco::StorageScalarType[]>(size, cudaStream); | ||
| recHits.did = cms::cuda::make_device_unique<uint32_t[]>(size, cudaStream); | ||
| recHits.flags = cms::cuda::make_device_unique<uint32_t[]>(size, cudaStream); | ||
| } | ||
| }; | ||
| } //namespace weights | ||
| } //namespace ecal | ||
|
|
||
| #endif // RecoLocalCalo_EcalRecProducers_plugins_DeclsForKernelsPhase2_h |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
103 changes: 103 additions & 0 deletions
103
RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,103 @@ | ||
| #include "FWCore/Framework/interface/stream/EDProducer.h" | ||
| #include "FWCore/Framework/interface/Event.h" | ||
| #include "FWCore/Framework/interface/EventSetup.h" | ||
| #include "FWCore/Framework/interface/MakerMacros.h" | ||
|
|
||
| #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" | ||
| #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" | ||
| #include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" | ||
|
|
||
| #include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" | ||
| #include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" | ||
|
|
||
| #include "DeclsForKernelsPhase2.h" | ||
|
|
||
| class EcalPhase2DigiToGPUProducer : public edm::stream::EDProducer<edm::ExternalWork> { | ||
| public: | ||
| explicit EcalPhase2DigiToGPUProducer(const edm::ParameterSet& ps); | ||
| ~EcalPhase2DigiToGPUProducer() override = default; | ||
| static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); | ||
|
|
||
| void acquire(edm::Event const& event, edm::EventSetup const& setup, edm::WaitingTaskWithArenaHolder holder) override; | ||
| void produce(edm::Event& evt, edm::EventSetup const& setup) override; | ||
|
|
||
| private: | ||
| const edm::EDGetTokenT<EBDigiCollectionPh2> digiCollectionToken_; | ||
| const edm::EDPutTokenT<cms::cuda::Product<ecal::DigisCollection<calo::common::DevStoragePolicy>>> | ||
| digisCollectionToken_; | ||
| uint32_t size_; | ||
|
|
||
| ecal::DigisCollection<::calo::common::DevStoragePolicy> digis_; | ||
|
|
||
| cms::cuda::ContextState cudaState_; | ||
| }; | ||
|
|
||
| void EcalPhase2DigiToGPUProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { | ||
| edm::ParameterSetDescription desc; | ||
|
|
||
| desc.add<edm::InputTag>("BarrelDigis", edm::InputTag("simEcalUnsuppressedDigis", "")); | ||
| desc.add<std::string>("digisLabelEB", "ebDigis"); | ||
|
|
||
| descriptions.addWithDefaultLabel(desc); | ||
| } | ||
|
|
||
| EcalPhase2DigiToGPUProducer::EcalPhase2DigiToGPUProducer(const edm::ParameterSet& ps) | ||
| : digiCollectionToken_(consumes<EBDigiCollectionPh2>(ps.getParameter<edm::InputTag>("BarrelDigis"))), | ||
| digisCollectionToken_(produces<cms::cuda::Product<ecal::DigisCollection<calo::common::DevStoragePolicy>>>( | ||
| ps.getParameter<std::string>("digisLabelEB"))) {} | ||
|
|
||
| void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, | ||
| edm::EventSetup const& setup, | ||
| edm::WaitingTaskWithArenaHolder holder) { | ||
| cms::cuda::ScopedContextAcquire ctx{event.streamID(), std::move(holder), cudaState_}; | ||
|
|
||
| //input data from event | ||
| const auto& pdigis = event.get(digiCollectionToken_); | ||
|
|
||
| size_ = pdigis.size(); | ||
|
|
||
| digis_.size = size_; | ||
| //allocate device pointers for output | ||
| digis_.ids = cms::cuda::make_device_unique<uint32_t[]>(size_, ctx.stream()); | ||
| digis_.data = cms::cuda::make_device_unique<uint16_t[]>(size_ * EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); | ||
|
|
||
| //allocate host pointers for holding product data and id vectors | ||
| auto idstmp = cms::cuda::make_host_unique<uint32_t[]>(size_, ctx.stream()); | ||
| auto datatmp = cms::cuda::make_host_unique<uint16_t[]>(size_ * EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); | ||
|
|
||
| //iterate over digis | ||
| uint32_t i = 0; | ||
| for (const auto& pdigi : pdigis) { | ||
| const int nSamples = pdigi.size(); | ||
| //assign id to output vector | ||
| idstmp.get()[i] = pdigi.id(); | ||
| //iterate over sample in digi | ||
| for (int sample = 0; sample < nSamples; ++sample) { | ||
| //get samples from input digi | ||
| EcalLiteDTUSample thisSample = pdigi[sample]; | ||
| //assign adc data to output | ||
| datatmp.get()[i * nSamples + sample] = thisSample.raw(); | ||
| } | ||
| ++i; | ||
| } | ||
|
|
||
| //copy output vectors into member variable device pointers for the output struct | ||
|
|
||
| cudaCheck( | ||
| cudaMemcpyAsync(digis_.ids.get(), idstmp.get(), size_ * sizeof(uint32_t), cudaMemcpyHostToDevice, ctx.stream())); | ||
| cudaCheck(cudaMemcpyAsync(digis_.data.get(), | ||
| datatmp.get(), | ||
| size_ * EcalDataFrame_Ph2::MAXSAMPLES * sizeof(uint16_t), | ||
| cudaMemcpyHostToDevice, | ||
| ctx.stream())); | ||
| } | ||
|
|
||
| void EcalPhase2DigiToGPUProducer::produce(edm::Event& event, edm::EventSetup const& setup) { | ||
| //get cuda context state for producer | ||
| cms::cuda::ScopedContextProduce ctx{cudaState_}; | ||
|
|
||
| //emplace output in the context | ||
| ctx.emplace(event, digisCollectionToken_, std::move(digis_)); | ||
| } | ||
|
|
||
| DEFINE_FWK_MODULE(EcalPhase2DigiToGPUProducer); | ||
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
In fact, I think this could also be done asynchronously, without the
acquire()/produce()split.@makortel is the expert here, but could you try to make the change and see if it works ?
edm::stream::EDProducer<>acquire()to the beginning ofproduce()size_anddigis_local variablescudaState_variable?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The possibility of dropping the
ExternalWorkwas something that has been mentioned already. However, sinceholderwas used to constructctxwe thought that it does not work in the end sinceholderis not available inproduce(). But usingctxwith typecms::cuda::ScopedContextProducecould make that possible.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
For what it's worth, I agree.
Right, without
ExternalWorktheScopedContextProduceshould be used.