Skip to content

Commit

Permalink
Apply code formatting
Browse files Browse the repository at this point in the history
  • Loading branch information
fwyzard committed Jul 12, 2020
1 parent b6e7d1b commit 2a223ba
Show file tree
Hide file tree
Showing 34 changed files with 358 additions and 441 deletions.
10 changes: 5 additions & 5 deletions CUDADataFormats/CaloCommon/interface/Common.h
Original file line number Diff line number Diff line change
Expand Up @@ -26,9 +26,9 @@ namespace calo {
uint32_t size;
};

template<>
template <>
struct AddSize<tags::DevPtr> {
uint32_t size;
uint32_t size;
};

struct ViewStoragePolicy {
Expand All @@ -43,9 +43,9 @@ namespace calo {
struct DevStoragePolicy {
using TagType = tags::DevPtr;

template<typename T>
template <typename T>
struct StorageSelector {
using type = cms::cuda::device::unique_ptr<T[]>;
using type = cms::cuda::device::unique_ptr<T[]>;
};
};

Expand All @@ -58,7 +58,7 @@ namespace calo {
using type = std::vector<T, Allocator<T>>;
};
};

template <typename T>
using CUDAHostAllocatorAlias = cms::cuda::HostAllocator<T>;

Expand Down
2 changes: 1 addition & 1 deletion CUDADataFormats/EcalDigi/interface/DigisCollection.h
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@

namespace ecal {

template<typename StoragePolicy>
template <typename StoragePolicy>
struct DigisCollection : public ::calo::common::AddSize<typename StoragePolicy::TagType> {
DigisCollection() = default;
DigisCollection(DigisCollection const &) = default;
Expand Down
6 changes: 4 additions & 2 deletions CUDADataFormats/EcalRecHitSoA/interface/EcalRecHit.h
Original file line number Diff line number Diff line change
Expand Up @@ -23,8 +23,10 @@ namespace ecal {
typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type time;
// should we remove the following, since already included in "extra" ?
typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type chi2;
typename StoragePolicy::template StorageSelector<uint32_t>::type extra; // packed uint32_t for timeError, chi2, energyError
typename StoragePolicy::template StorageSelector<uint32_t>::type flagBits; // store rechit condition (see Flags enum) in a bit-wise way
typename StoragePolicy::template StorageSelector<uint32_t>::type
extra; // packed uint32_t for timeError, chi2, energyError
typename StoragePolicy::template StorageSelector<uint32_t>::type
flagBits; // store rechit condition (see Flags enum) in a bit-wise way

typename StoragePolicy::template StorageSelector<uint32_t>::type did;

Expand Down
15 changes: 6 additions & 9 deletions EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h
Original file line number Diff line number Diff line change
Expand Up @@ -35,19 +35,16 @@ namespace ecal {
};

struct OutputDataGPU {
DigisCollection<::calo::common::DevStoragePolicy> digisEB, digisEE;
DigisCollection<::calo::common::DevStoragePolicy> digisEB, digisEE;

// FIXME: we should separate max channels parameter for eb and ee
// FIXME: replace hardcoded values
void allocate(ConfigurationParameters const &config, cudaStream_t cudaStream) {
digisEB.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannels, cudaStream);
digisEE.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannels, cudaStream);

digisEB.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannels, cudaStream);
digisEE.ids = cms::cuda::make_device_unique<uint32_t[]>( config.maxChannels, cudaStream);
digisEB.data = cms::cuda::make_device_unique<uint16_t[]>(config.maxChannels, cudaStream);
digisEE.data = cms::cuda::make_device_unique<uint16_t[]>(config.maxChannels, cudaStream);

digisEB.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannels, cudaStream);
digisEE.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannels, cudaStream);
}
};

Expand Down
6 changes: 2 additions & 4 deletions EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc
Original file line number Diff line number Diff line change
Expand Up @@ -81,10 +81,8 @@ void EcalCPUDigisProducer::fillDescriptions(edm::ConfigurationDescriptions& conf

EcalCPUDigisProducer::EcalCPUDigisProducer(const edm::ParameterSet& ps)
: // input digi collections in GPU-friendly format
digisInEBToken_{
consumes<InputProduct>(ps.getParameter<edm::InputTag>("digisInLabelEB"))},
digisInEEToken_{
consumes<InputProduct>(ps.getParameter<edm::InputTag>("digisInLabelEE"))},
digisInEBToken_{consumes<InputProduct>(ps.getParameter<edm::InputTag>("digisInLabelEB"))},
digisInEEToken_{consumes<InputProduct>(ps.getParameter<edm::InputTag>("digisInLabelEE"))},
// output digi collections in legacy format
digisOutEBToken_{produces<EBDigiCollection>(ps.getParameter<std::string>("digisOutLabelEB"))},
digisOutEEToken_{produces<EEDigiCollection>(ps.getParameter<std::string>("digisOutLabelEE"))},
Expand Down
34 changes: 10 additions & 24 deletions EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc
Original file line number Diff line number Diff line change
Expand Up @@ -64,8 +64,7 @@ EcalRawToDigiGPU::EcalRawToDigiGPU(const edm::ParameterSet& ps)
config_.maxChannels = ps.getParameter<uint32_t>("maxChannels");
}

EcalRawToDigiGPU::~EcalRawToDigiGPU() {
}
EcalRawToDigiGPU::~EcalRawToDigiGPU() {}

void EcalRawToDigiGPU::acquire(edm::Event const& event,
edm::EventSetup const& setup,
Expand All @@ -86,35 +85,22 @@ void EcalRawToDigiGPU::acquire(edm::Event const& event,
event.getByToken(rawDataToken_, rawDataHandle);

// scratch
ecal::raw::ScratchDataGPU scratchGPU = {
cms::cuda::make_device_unique<uint32_t[]>(2, ctx.stream())
};
ecal::raw::ScratchDataGPU scratchGPU = {cms::cuda::make_device_unique<uint32_t[]>(2, ctx.stream())};

// input cpu data
ecal::raw::InputDataCPU inputCPU = {
cms::cuda::make_host_unique<unsigned char[]>(
ecal::raw::nfeds_max * ecal::raw::nbytes_per_fed_max, ctx.stream()),
cms::cuda::make_host_unique<uint32_t[]>(
ecal::raw::nfeds_max, ctx.stream()),
cms::cuda::make_host_unique<int[]>(
ecal::raw::nfeds_max, ctx.stream())
};
cms::cuda::make_host_unique<unsigned char[]>(ecal::raw::nfeds_max * ecal::raw::nbytes_per_fed_max, ctx.stream()),
cms::cuda::make_host_unique<uint32_t[]>(ecal::raw::nfeds_max, ctx.stream()),
cms::cuda::make_host_unique<int[]>(ecal::raw::nfeds_max, ctx.stream())};

// input data gpu
ecal::raw::InputDataGPU inputGPU = {
cms::cuda::make_device_unique<unsigned char[]>(
ecal::raw::nfeds_max * ecal::raw::nbytes_per_fed_max, ctx.stream()),
cms::cuda::make_device_unique<uint32_t[]>(
ecal::raw::nfeds_max, ctx.stream()),
cms::cuda::make_device_unique<int[]>(
ecal::raw::nfeds_max, ctx.stream())
};
ecal::raw::InputDataGPU inputGPU = {cms::cuda::make_device_unique<unsigned char[]>(
ecal::raw::nfeds_max * ecal::raw::nbytes_per_fed_max, ctx.stream()),
cms::cuda::make_device_unique<uint32_t[]>(ecal::raw::nfeds_max, ctx.stream()),
cms::cuda::make_device_unique<int[]>(ecal::raw::nfeds_max, ctx.stream())};

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

// output gpu
outputGPU_.allocate(config_, ctx.stream());
Expand Down
7 changes: 5 additions & 2 deletions EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu
Original file line number Diff line number Diff line change
Expand Up @@ -299,8 +299,11 @@ namespace ecal {
uint32_t const nfedsWithData,
uint32_t const nbytesTotal) {
// transfer
cudaCheck(cudaMemcpyAsync(
inputGPU.data.get(), inputCPU.data.get(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
cudaCheck(cudaMemcpyAsync(inputGPU.data.get(),
inputCPU.data.get(),
nbytesTotal * sizeof(unsigned char),
cudaMemcpyHostToDevice,
cudaStream));
cudaCheck(cudaMemcpyAsync(inputGPU.offsets.get(),
inputCPU.offsets.get(),
nfedsWithData * sizeof(uint32_t),
Expand Down
32 changes: 7 additions & 25 deletions EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
Original file line number Diff line number Diff line change
Expand Up @@ -57,45 +57,27 @@ namespace hcal {

void allocate(ConfigurationParameters const &config, cudaStream_t cudaStream) {
digisF01HE.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannelsF01HE*compute_stride<Flavor01>(config.nsamplesF01HE),
cudaStream
);
config.maxChannelsF01HE * compute_stride<Flavor01>(config.nsamplesF01HE), cudaStream);
//cudaCheck(
// cudaMalloc((void **)&digisF01HE.data,
// config.maxChannelsF01HE * sizeof(uint16_t) * compute_stride<Flavor01>(config.nsamplesF01HE)));
digisF01HE.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannelsF01HE,
cudaStream
);
digisF01HE.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannelsF01HE, cudaStream);
//cudaCheck(cudaMalloc((void **)&digisF01HE.ids, sizeof(uint32_t) * config.maxChannelsF01HE));

digisF5HB.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannelsF5HB * compute_stride<Flavor5>(config.nsamplesF5HB),
cudaStream
);
config.maxChannelsF5HB * compute_stride<Flavor5>(config.nsamplesF5HB), cudaStream);
//cudaCheck(cudaMalloc((void **)&digisF5HB.data,
// config.maxChannelsF5HB * sizeof(uint16_t) * compute_stride<Flavor5>(config.nsamplesF5HB)));
digisF5HB.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannelsF5HB,
cudaStream
);
digisF5HB.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannelsF5HB, cudaStream);
//cudaCheck(cudaMalloc((void **)&digisF5HB.ids, sizeof(uint32_t) * config.maxChannelsF5HB));
digisF5HB.npresamples = cms::cuda::make_device_unique<uint8_t[]>(
config.maxChannelsF5HB,
cudaStream
);
digisF5HB.npresamples = cms::cuda::make_device_unique<uint8_t[]>(config.maxChannelsF5HB, cudaStream);
//cudaCheck(cudaMalloc((void **)&digisF5HB.npresamples, sizeof(uint8_t) * config.maxChannelsF5HB));

digisF3HB.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannelsF3HB * compute_stride<Flavor3>(config.nsamplesF3HB),
cudaStream
);
config.maxChannelsF3HB * compute_stride<Flavor3>(config.nsamplesF3HB), cudaStream);
//cudaCheck(cudaMalloc((void **)&digisF3HB.data,
// config.maxChannelsF3HB * sizeof(uint16_t) * compute_stride<Flavor3>(config.nsamplesF3HB)));
digisF3HB.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannelsF3HB,
cudaStream
);
digisF3HB.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannelsF3HB, cudaStream);
//cudaCheck(cudaMalloc((void **)&digisF3HB.ids, config.maxChannelsF3HB * sizeof(uint32_t)));
}
};
Expand Down
10 changes: 7 additions & 3 deletions EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
Original file line number Diff line number Diff line change
Expand Up @@ -567,14 +567,18 @@ namespace hcal {
uint32_t const nfedsWithData,
uint32_t const nbytesTotal) {
// transfer
cudaCheck(cudaMemcpyAsync(
inputGPU.data.get(), inputCPU.data.get(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
cudaCheck(cudaMemcpyAsync(inputGPU.data.get(),
inputCPU.data.get(),
nbytesTotal * sizeof(unsigned char),
cudaMemcpyHostToDevice,
cudaStream));
cudaCheck(cudaMemcpyAsync(inputGPU.offsets.get(),
inputCPU.offsets.get(),
nfedsWithData * sizeof(uint32_t),
cudaMemcpyHostToDevice,
cudaStream));
cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters.get(), 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
cudaCheck(
cudaMemsetAsync(scratchGPU.pChannelsCounters.get(), 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
cudaCheck(cudaMemcpyAsync(
inputGPU.feds.get(), inputCPU.feds.get(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));

Expand Down
37 changes: 8 additions & 29 deletions EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
Original file line number Diff line number Diff line change
Expand Up @@ -125,40 +125,19 @@ void HcalDigisProducerGPU::acquire(edm::Event const& event,

// flavor 0/1 get devie blobs
df01_.data = cms::cuda::make_device_unique<uint16_t[]>(
config_.maxChannelsF01HE * hcal::compute_stride<hcal::Flavor01>(
config_.nsamplesF01HE),
ctx.stream()
);
df01_.ids = cms::cuda::make_device_unique<uint32_t[]>(
config_.maxChannelsF01HE,
ctx.stream()
);
config_.maxChannelsF01HE * hcal::compute_stride<hcal::Flavor01>(config_.nsamplesF01HE), ctx.stream());
df01_.ids = cms::cuda::make_device_unique<uint32_t[]>(config_.maxChannelsF01HE, ctx.stream());

// flavor3 get device blobs
df3_.data = cms::cuda::make_device_unique<uint16_t[]>(
config_.maxChannelsF3HB * hcal::compute_stride<hcal::Flavor3>(
config_.nsamplesF3HB),
ctx.stream()
);
df3_.ids = cms::cuda::make_device_unique<uint32_t[]>(
config_.maxChannelsF3HB,
ctx.stream()
);
df3_.data = cms::cuda::make_device_unique<uint16_t[]>(
config_.maxChannelsF3HB * hcal::compute_stride<hcal::Flavor3>(config_.nsamplesF3HB), ctx.stream());
df3_.ids = cms::cuda::make_device_unique<uint32_t[]>(config_.maxChannelsF3HB, ctx.stream());

// flavor5 get device blobs
df5_.data = cms::cuda::make_device_unique<uint16_t[]>(
config_.maxChannelsF5HB * hcal::compute_stride<hcal::Flavor5>(
config_.nsamplesF5HB),
ctx.stream()
);
df5_.ids = cms::cuda::make_device_unique<uint32_t[]>(
config_.maxChannelsF5HB,
ctx.stream()
);
df5_.npresamples = cms::cuda::make_device_unique<uint8_t[]>(
config_.maxChannelsF5HB,
ctx.stream()
);
config_.maxChannelsF5HB * hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB), ctx.stream());
df5_.ids = cms::cuda::make_device_unique<uint32_t[]>(config_.maxChannelsF5HB, ctx.stream());
df5_.npresamples = cms::cuda::make_device_unique<uint8_t[]>(config_.maxChannelsF5HB, ctx.stream());

for (auto const& hbhe : *hbheDigis) {
auto const id = hbhe.id().rawId();
Expand Down
46 changes: 12 additions & 34 deletions EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
Original file line number Diff line number Diff line change
Expand Up @@ -83,8 +83,7 @@ HcalRawToDigiGPU::HcalRawToDigiGPU(const edm::ParameterSet& ps)
config_.nsamplesF3HB = ps.getParameter<uint32_t>("nsamplesF3HB");
}

HcalRawToDigiGPU::~HcalRawToDigiGPU() {
}
HcalRawToDigiGPU::~HcalRawToDigiGPU() {}

void HcalRawToDigiGPU::acquire(edm::Event const& event,
edm::EventSetup const& setup,
Expand All @@ -106,44 +105,23 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,

// scratch
hcal::raw::ScratchDataGPU scratchGPU = {
cms::cuda::make_device_unique<uint32_t[]>(
hcal::raw::numOutputCollections,
ctx.stream())
};
cms::cuda::make_device_unique<uint32_t[]>(hcal::raw::numOutputCollections, ctx.stream())};

// input cpu data
hcal::raw::InputDataCPU inputCPU = {
cms::cuda::make_host_unique<unsigned char[]>(
hcal::raw::utca_nfeds_max * hcal::raw::nbytes_per_fed_max,
ctx.stream()),
cms::cuda::make_host_unique<uint32_t[]>(
hcal::raw::utca_nfeds_max,
ctx.stream()),
cms::cuda::make_host_unique<int[]>(
hcal::raw::utca_nfeds_max,
ctx.stream())
};
hcal::raw::InputDataCPU inputCPU = {cms::cuda::make_host_unique<unsigned char[]>(
hcal::raw::utca_nfeds_max * hcal::raw::nbytes_per_fed_max, ctx.stream()),
cms::cuda::make_host_unique<uint32_t[]>(hcal::raw::utca_nfeds_max, ctx.stream()),
cms::cuda::make_host_unique<int[]>(hcal::raw::utca_nfeds_max, ctx.stream())};

// input data gpu
hcal::raw::InputDataGPU inputGPU = {
cms::cuda::make_device_unique<unsigned char[]>(
hcal::raw::utca_nfeds_max * hcal::raw::nbytes_per_fed_max,
ctx.stream()),
cms::cuda::make_device_unique<uint32_t[]>(
hcal::raw::utca_nfeds_max,
ctx.stream()),
cms::cuda::make_device_unique<int[]>(
hcal::raw::utca_nfeds_max,
ctx.stream())
};
cms::cuda::make_device_unique<unsigned char[]>(hcal::raw::utca_nfeds_max * hcal::raw::nbytes_per_fed_max,
ctx.stream()),
cms::cuda::make_device_unique<uint32_t[]>(hcal::raw::utca_nfeds_max, ctx.stream()),
cms::cuda::make_device_unique<int[]>(hcal::raw::utca_nfeds_max, ctx.stream())};

// output cpu
outputCPU_ = {
cms::cuda::make_host_unique<uint32_t[]>(
hcal::raw::numOutputCollections,
ctx.stream()
)
};
outputCPU_ = {cms::cuda::make_host_unique<uint32_t[]>(hcal::raw::numOutputCollections, ctx.stream())};

// output gpu
outputGPU_.allocate(config_, ctx.stream());
Expand Down Expand Up @@ -214,7 +192,7 @@ void HcalRawToDigiGPU::produce(edm::Event& event, edm::EventSetup const& setup)
ctx.emplace(event, digisF01HEToken_, std::move(outputGPU_.digisF01HE));
ctx.emplace(event, digisF5HBToken_, std::move(outputGPU_.digisF5HB));
ctx.emplace(event, digisF3HBToken_, std::move(outputGPU_.digisF3HB));

// reset ptrs that are carried as members
outputCPU_.nchannels.reset();
}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -35,8 +35,10 @@ int main(int argc, char *argv[]) {

gStyle->SetOptStat("ourme");

edm::Wrapper<ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>> *wgpuEB = nullptr;
edm::Wrapper<ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>> *wgpuEE = nullptr;
edm::Wrapper<ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>> *wgpuEB =
nullptr;
edm::Wrapper<ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>> *wgpuEE =
nullptr;
edm::Wrapper<EBUncalibratedRecHitCollection> *wcpuEB = nullptr;
edm::Wrapper<EEUncalibratedRecHitCollection> *wcpuEE = nullptr;

Expand Down Expand Up @@ -152,10 +154,14 @@ int main(int argc, char *argv[]) {
std::cout << "validating file " << fileName << std::endl;
TFile rf{fileName.c_str()};
TTree *rt = (TTree *)rf.Get("Events");
rt->SetBranchAddress("calocommonCUDAHostAllocatorAliascalocommonVecStoragePolicyecalUncalibratedRecHit_ecalCPUUncalibRecHitProducer_EcalUncalibRecHitsEB_RECO.",
&wgpuEB);
rt->SetBranchAddress("calocommonCUDAHostAllocatorAliascalocommonVecStoragePolicyecalUncalibratedRecHit_ecalCPUUncalibRecHitProducer_EcalUncalibRecHitsEE_RECO.",
&wgpuEE);
rt->SetBranchAddress(
"calocommonCUDAHostAllocatorAliascalocommonVecStoragePolicyecalUncalibratedRecHit_ecalCPUUncalibRecHitProducer_"
"EcalUncalibRecHitsEB_RECO.",
&wgpuEB);
rt->SetBranchAddress(
"calocommonCUDAHostAllocatorAliascalocommonVecStoragePolicyecalUncalibratedRecHit_ecalCPUUncalibRecHitProducer_"
"EcalUncalibRecHitsEE_RECO.",
&wgpuEE);
rt->SetBranchAddress("EcalUncalibratedRecHitsSorted_ecalMultiFitUncalibRecHit_EcalUncalibRecHitsEB_RECO.", &wcpuEB);
rt->SetBranchAddress("EcalUncalibratedRecHitsSorted_ecalMultiFitUncalibRecHit_EcalUncalibRecHitsEE_RECO.", &wcpuEE);

Expand Down
Loading

0 comments on commit 2a223ba

Please sign in to comment.