Skip to content

Commit 32f4925

Browse files
cmsbuildcmsbuild
authored
Merge pull request cms-sw#34664 from stahlleiton/CMSSW_12_0_X_Patatrack_HIon_v0
Adapted EcalRawToDigi and SiPixelRawToCluster GPU code for HIon
2 parents 8c5d4c0 + 49c4c70 commit 32f4925

File tree

5 files changed

+24
-16
lines changed

5 files changed

+24
-16
lines changed

‎EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@ class EcalRawToDigiGPU : public edm::stream::EDProducer<edm::ExternalWork> {
3333

3434
cms::cuda::ContextState cudaState_;
3535

36+
const uint32_t maxFedSize_;
3637
std::vector<int> fedsToUnpack_;
3738

3839
ecal::raw::ConfigurationParameters config_;
@@ -44,6 +45,7 @@ void EcalRawToDigiGPU::fillDescriptions(edm::ConfigurationDescriptions& confDesc
4445
edm::ParameterSetDescription desc;
4546

4647
desc.add<edm::InputTag>("InputLabel", edm::InputTag("rawDataCollector"));
48+
desc.add<uint32_t>("maxFedSize", ecal::raw::nbytes_per_fed_max);
4749
std::vector<int> feds(54);
4850
for (uint32_t i = 0; i < 54; ++i)
4951
feds[i] = i + 601;
@@ -62,6 +64,7 @@ EcalRawToDigiGPU::EcalRawToDigiGPU(const edm::ParameterSet& ps)
6264
digisEBToken_{produces<OutputProduct>(ps.getParameter<std::string>("digisLabelEB"))},
6365
digisEEToken_{produces<OutputProduct>(ps.getParameter<std::string>("digisLabelEE"))},
6466
eMappingToken_{esConsumes<ecal::raw::ElectronicsMappingGPU, EcalMappingElectronicsRcd>()},
67+
maxFedSize_{ps.getParameter<uint32_t>("maxFedSize")},
6568
fedsToUnpack_{ps.getParameter<std::vector<int>>("FEDs")} {
6669
config_.maxChannelsEB = ps.getParameter<uint32_t>("maxChannelsEB");
6770
config_.maxChannelsEE = ps.getParameter<uint32_t>("maxChannelsEE");
@@ -91,15 +94,15 @@ void EcalRawToDigiGPU::acquire(edm::Event const& event,
9194

9295
// input cpu data
9396
ecal::raw::InputDataCPU inputCPU = {
94-
cms::cuda::make_host_unique<unsigned char[]>(ecal::raw::nfeds_max * ecal::raw::nbytes_per_fed_max, ctx.stream()),
97+
cms::cuda::make_host_unique<unsigned char[]>(ecal::raw::nfeds_max * maxFedSize_, ctx.stream()),
9598
cms::cuda::make_host_unique<uint32_t[]>(ecal::raw::nfeds_max, ctx.stream()),
9699
cms::cuda::make_host_unique<int[]>(ecal::raw::nfeds_max, ctx.stream())};
97100

98101
// input data gpu
99-
ecal::raw::InputDataGPU inputGPU = {cms::cuda::make_device_unique<unsigned char[]>(
100-
ecal::raw::nfeds_max * ecal::raw::nbytes_per_fed_max, ctx.stream()),
101-
cms::cuda::make_device_unique<uint32_t[]>(ecal::raw::nfeds_max, ctx.stream()),
102-
cms::cuda::make_device_unique<int[]>(ecal::raw::nfeds_max, ctx.stream())};
102+
ecal::raw::InputDataGPU inputGPU = {
103+
cms::cuda::make_device_unique<unsigned char[]>(ecal::raw::nfeds_max * maxFedSize_, ctx.stream()),
104+
cms::cuda::make_device_unique<uint32_t[]>(ecal::raw::nfeds_max, ctx.stream()),
105+
cms::cuda::make_device_unique<int[]>(ecal::raw::nfeds_max, ctx.stream())};
103106

104107
// output cpu
105108
outputCPU_ = {cms::cuda::make_host_unique<uint32_t[]>(2, ctx.stream())};

‎RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -78,6 +78,7 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
7878
const bool isRun2_;
7979
const bool includeErrors_;
8080
const bool useQuality_;
81+
const uint32_t maxFedWords_;
8182
const SiPixelClusterThresholds clusterThresholds_;
8283
};
8384

@@ -92,6 +93,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi
9293
isRun2_(iConfig.getParameter<bool>("isRun2")),
9394
includeErrors_(iConfig.getParameter<bool>("IncludeErrors")),
9495
useQuality_(iConfig.getParameter<bool>("UseQualityInfo")),
96+
maxFedWords_(iConfig.getParameter<uint32_t>("MaxFEDWords")),
9597
clusterThresholds_{iConfig.getParameter<int32_t>("clusterThreshold_layer1"),
9698
iConfig.getParameter<int32_t>("clusterThreshold_otherLayers")} {
9799
if (includeErrors_) {
@@ -105,7 +107,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi
105107

106108
edm::Service<CUDAService> cs;
107109
if (cs->enabled()) {
108-
wordFedAppender_ = std::make_unique<pixelgpudetails::SiPixelRawToClusterGPUKernel::WordFedAppender>();
110+
wordFedAppender_ = std::make_unique<pixelgpudetails::SiPixelRawToClusterGPUKernel::WordFedAppender>(maxFedWords_);
109111
}
110112
}
111113

@@ -114,6 +116,7 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d
114116
desc.add<bool>("isRun2", true);
115117
desc.add<bool>("IncludeErrors", true);
116118
desc.add<bool>("UseQualityInfo", false);
119+
desc.add<uint32_t>("MaxFEDWords", pixelgpudetails::MAX_FED * pixelgpudetails::MAX_WORD);
117120
desc.add<int32_t>("clusterThreshold_layer1", kSiPixelClusterThresholdsDefaultPhase1.layer1);
118121
desc.add<int32_t>("clusterThreshold_otherLayers", kSiPixelClusterThresholdsDefaultPhase1.otherLayers);
119122
desc.add<edm::InputTag>("InputLabel", edm::InputTag("rawDataCollector"));
@@ -245,6 +248,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent,
245248
std::move(errors_),
246249
wordCounterGPU,
247250
fedCounter,
251+
maxFedWords_,
248252
useQuality_,
249253
includeErrors_,
250254
edm::MessageDrop::instance()->debugEnabled,

‎RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu‎

Lines changed: 7 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -35,12 +35,9 @@
3535

3636
namespace pixelgpudetails {
3737

38-
// number of words for all the FEDs
39-
constexpr uint32_t MAX_FED_WORDS = pixelgpudetails::MAX_FED * pixelgpudetails::MAX_WORD;
40-
41-
SiPixelRawToClusterGPUKernel::WordFedAppender::WordFedAppender() {
42-
word_ = cms::cuda::make_host_noncached_unique<unsigned int[]>(MAX_FED_WORDS, cudaHostAllocWriteCombined);
43-
fedId_ = cms::cuda::make_host_noncached_unique<unsigned char[]>(MAX_FED_WORDS, cudaHostAllocWriteCombined);
38+
SiPixelRawToClusterGPUKernel::WordFedAppender::WordFedAppender(uint32_t maxFedWords) {
39+
word_ = cms::cuda::make_host_noncached_unique<unsigned int[]>(maxFedWords, cudaHostAllocWriteCombined);
40+
fedId_ = cms::cuda::make_host_noncached_unique<unsigned char[]>(maxFedWords, cudaHostAllocWriteCombined);
4441
}
4542

4643
void SiPixelRawToClusterGPUKernel::WordFedAppender::initializeWordFed(int fedId,
@@ -505,19 +502,20 @@ namespace pixelgpudetails {
505502
SiPixelFormatterErrors &&errors,
506503
const uint32_t wordCounter,
507504
const uint32_t fedCounter,
505+
const uint32_t maxFedWords,
508506
bool useQualityInfo,
509507
bool includeErrors,
510508
bool debug,
511509
cudaStream_t stream) {
512510
nDigis = wordCounter;
513511

514512
#ifdef GPU_DEBUG
515-
std::cout << "decoding " << wordCounter << " digis. Max is " << pixelgpudetails::MAX_FED_WORDS << std::endl;
513+
std::cout << "decoding " << wordCounter << " digis. Max is " << maxFedWords << std::endl;
516514
#endif
517515

518-
digis_d = SiPixelDigisCUDA(pixelgpudetails::MAX_FED_WORDS, stream);
516+
digis_d = SiPixelDigisCUDA(maxFedWords, stream);
519517
if (includeErrors) {
520-
digiErrors_d = SiPixelDigiErrorsCUDA(pixelgpudetails::MAX_FED_WORDS, std::move(errors), stream);
518+
digiErrors_d = SiPixelDigiErrorsCUDA(maxFedWords, std::move(errors), stream);
521519
}
522520
clusters_d = SiPixelClustersCUDA(gpuClustering::maxNumModules, stream);
523521

‎RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -120,6 +120,7 @@ namespace pixelgpudetails {
120120
class WordFedAppender {
121121
public:
122122
WordFedAppender();
123+
WordFedAppender(uint32_t maxFedWords);
123124
~WordFedAppender() = default;
124125

125126
void initializeWordFed(int fedId, unsigned int wordCounterGPU, const cms_uint32_t* src, unsigned int length);
@@ -149,6 +150,7 @@ namespace pixelgpudetails {
149150
SiPixelFormatterErrors&& errors,
150151
const uint32_t wordCounter,
151152
const uint32_t fedCounter,
153+
const uint32_t maxFedWords,
152154
bool useQualityInfo,
153155
bool includeErrors,
154156
bool debug,

‎RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -75,7 +75,8 @@ namespace gpuClustering {
7575
}
7676

7777
//init hist (ymax=416 < 512 : 9bits)
78-
constexpr uint32_t maxPixInModule = 4000;
78+
//6000 max pixels required for HI operations with no measurable impact on pp performance
79+
constexpr uint32_t maxPixInModule = 6000;
7980
constexpr auto nbins = phase1PixelTopology::numColsInModule + 2; //2+2;
8081
using Hist = cms::cuda::HistoContainer<uint16_t, nbins, maxPixInModule, 9, uint16_t>;
8182
__shared__ Hist hist;

0 commit comments

Comments
 (0)