Skip to content

Commit

Permalink
Merged devel_pfRecHitAlpakaES_integ_v2 from repository missirol with …
Browse files Browse the repository at this point in the history
…cms-merge-topic
  • Loading branch information
fllor committed Feb 23, 2023
2 parents cdc6edf + eec8f55 commit d24418b
Show file tree
Hide file tree
Showing 27 changed files with 728 additions and 26 deletions.
3 changes: 3 additions & 0 deletions RecoParticleFlow/PFClusterProducer/plugins/DeclsForKernels.h
Original file line number Diff line number Diff line change
Expand Up @@ -240,6 +240,7 @@ namespace PFClustering {
cms::cuda::device::unique_ptr<int> posL;
cms::cuda::device::unique_ptr<int> topH;
cms::cuda::device::unique_ptr<int> posH;
cms::cuda::device::unique_ptr<int> nTopoId;

cms::cuda::device::unique_ptr<float[]> pcrh_fracSum;
cms::cuda::device::unique_ptr<float4[]> pfc_prevPos4;
Expand All @@ -258,6 +259,8 @@ namespace PFClustering {

posH = cms::cuda::make_device_unique<int>(cudaStream);

nTopoId = cms::cuda::make_device_unique<int>(cudaStream);

pcrh_fracSum = cms::cuda::make_device_unique<float[]>(sizeof(float) * config.maxRH, cudaStream);
pfc_prevPos4 = cms::cuda::make_device_unique<float4[]>(sizeof(float4) * config.maxRH, cudaStream);
}
Expand Down
24 changes: 17 additions & 7 deletions RecoParticleFlow/PFClusterProducer/plugins/PFClusterCudaHCAL.cu
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,6 @@ using PFClustering::common::PFLayer;
constexpr const float PI_F = 3.141592654f;

// Number of neighbors considered for topo clustering
constexpr const int nNT = 8;

namespace PFClusterCudaHCAL {
//
Expand Down Expand Up @@ -1314,6 +1313,8 @@ namespace PFClusterCudaHCAL {
}

// Contraction in a single block

__device__ int nTopo;
__global__ void topoClusterContraction(size_t size,
int* pfrh_parent,
int* pfrh_isSeed,
Expand All @@ -1325,10 +1326,12 @@ namespace PFClusterCudaHCAL {
int* topoSeedList,
int* pcrhfracind,
float* pcrhfrac,
int* pcrhFracSize) {
__shared__ int notDone, totalSeedOffset, totalSeedFracOffset;
int* pcrhFracSize,
int* nTopoId) {
__shared__ int totalSeedOffset, totalSeedFracOffset;
if (threadIdx.x == 0) {
notDone = 0;
*nTopoId = 0;
nTopo = 0;
totalSeedOffset = 0;
totalSeedFracOffset = 0;
*pcrhFracSize = 0;
Expand Down Expand Up @@ -1372,6 +1375,7 @@ namespace PFClusterCudaHCAL {
// This is a valid topo ID
int offset = atomicAdd(&totalSeedOffset, topoSeedCount[topoId]);
topoSeedOffsets[topoId] = offset;
atomicAdd(&*nTopoId, 1);
}
}
__syncthreads();
Expand Down Expand Up @@ -1405,6 +1409,7 @@ namespace PFClusterCudaHCAL {
__syncthreads();

if (threadIdx.x == 0) {
nTopo = *nTopoId;
*pcrhFracSize = totalSeedFracOffset;
if (*pcrhFracSize > 200000) // DeclsForKernels.h maxPFCFracs
printf("At the end of topoClusterContraction, found large *pcrhFracSize = %d\n", *pcrhFracSize);
Expand Down Expand Up @@ -1663,7 +1668,7 @@ namespace PFClusterCudaHCAL {
nRH,
inputPFRecHits.pfrh_energy.get(),
inputPFRecHits.pfrh_x.get(),
inputPFRecHits.pfrh_y.get(),
inputPFRecHits.pfrh_y.get(),
inputPFRecHits.pfrh_z.get(),
outputGPU.pfrh_isSeed.get(),
outputGPU.pfrh_topoId.get(),
Expand Down Expand Up @@ -1732,11 +1737,16 @@ namespace PFClusterCudaHCAL {
outputGPU.topoSeedList.get(),
outputGPU.pcrh_fracInd.get(),
outputGPU.pcrh_frac.get(),
outputGPU.pcrhFracSize.get());
outputGPU.pcrhFracSize.get(),
scratchGPU.nTopoId.get());

dim3 grid((nRH + 31) / 32, (nRH + 31) / 32);
dim3 block(32, 32);

typeof(nTopo) h_nTopo;

cudaCheck(cudaMemcpyFromSymbolAsync(&h_nTopo, nTopo, sizeof(int), 0, cudaMemcpyDeviceToHost, cudaStream));

fillRhfIndex<<<grid, block, 0, cudaStream>>>(nRH,
outputGPU.pfrh_topoId.get(),
outputGPU.pfrh_isSeed.get(),
Expand All @@ -1746,7 +1756,7 @@ namespace PFClusterCudaHCAL {
scratchGPU.rhcount.get(),
outputGPU.pcrh_fracInd.get());

hcalFastCluster_selection<<<nRH, 256, 0, cudaStream>>>(pfClusParams.const_view(),
hcalFastCluster_selection<<<nRH, threadsPerBlock, 0, cudaStream>>>(pfClusParams.const_view(),
nRH,
inputPFRecHits.pfrh_x.get(),
inputPFRecHits.pfrh_y.get(),
Expand Down
4 changes: 4 additions & 0 deletions RecoParticleFlow/PFRecHitProducer/BuildFile.xml
Original file line number Diff line number Diff line change
@@ -1,4 +1,8 @@
<use name="FWCore/Framework"/>
<use name="FWCore/Utilities"/>
<use name="DataFormats/Portable"/>
<use name="DataFormats/SoATemplate"/>
<use name="Geometry/Records"/>
<use name="HeterogeneousCore/AlpakaCore"/>
<use name="HeterogeneousCore/AlpakaInterface"/>
<flags ALPAKA_BACKENDS="1"/>
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,8 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_JobConfigurationAlpakaRecord_h
#define RecoParticleFlow_PFRecHitProducer_interface_JobConfigurationAlpakaRecord_h

#include "FWCore/Framework/interface/EventSetupRecordImplementation.h"

class JobConfigurationAlpakaRecord : public edm::eventsetup::EventSetupRecordImplementation<JobConfigurationAlpakaRecord> {};

#endif // RecoParticleFlow_PFRecHitProducer_interface_JobConfigurationAlpakaRecord_h
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_AlpakaESTestData_h
#define RecoParticleFlow_PFRecHitProducer_interface_AlpakaESTestData_h

#include "DataFormats/Portable/interface/PortableHostCollection.h"
#include "HeterogeneousCore/AlpakaInterface/interface/CopyToDevice.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"

#include "RecoParticleFlow/PFRecHitProducer/interface/PFRecHitHBHEParamsAlpakaESDataSoA.h"

namespace reco {

using PFRecHitHBHEParamsAlpakaESDataHost = PortableHostCollection<PFRecHitHBHEParamsAlpakaESDataSoA>;

}

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHEParamsAlpakaESDataSoA_h
#define RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHEParamsAlpakaESDataSoA_h

#include "DataFormats/SoATemplate/interface/SoACommon.h"
#include "DataFormats/SoATemplate/interface/SoALayout.h"
#include "DataFormats/SoATemplate/interface/SoAView.h"

namespace reco {

GENERATE_SOA_LAYOUT(PFRecHitHBHEParamsAlpakaESDataSoALayout,
SOA_COLUMN(float, energyThresholds))

using PFRecHitHBHEParamsAlpakaESDataSoA = PFRecHitHBHEParamsAlpakaESDataSoALayout<>;

}

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESData_h
#define RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESData_h

#include "DataFormats/Portable/interface/PortableHostCollection.h"
#include "HeterogeneousCore/AlpakaInterface/interface/CopyToDevice.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"

#include "RecoParticleFlow/PFRecHitProducer/interface/PFRecHitHBHETopologyAlpakaESDataSoA.h"

namespace reco {

using PFRecHitHBHETopologyAlpakaESDataHost = PortableHostCollection<PFRecHitHBHETopologyAlpakaESDataSoA>;

}

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,27 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESDataSoA_h
#define RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESDataSoA_h

#include "DataFormats/SoATemplate/interface/SoACommon.h"
#include "DataFormats/SoATemplate/interface/SoALayout.h"
#include "DataFormats/SoATemplate/interface/SoAView.h"

namespace reco {

GENERATE_SOA_LAYOUT(PFRecHitHBHETopologyAlpakaESDataSoALayout,
SOA_COLUMN(float, positionX),
SOA_COLUMN(float, positionY),
SOA_COLUMN(float, positionZ),
SOA_COLUMN(int32_t, neighbour0),
SOA_COLUMN(int32_t, neighbour1),
SOA_COLUMN(int32_t, neighbour2),
SOA_COLUMN(int32_t, neighbour3),
SOA_COLUMN(int32_t, neighbour4),
SOA_COLUMN(int32_t, neighbour5),
SOA_COLUMN(int32_t, neighbour6),
SOA_COLUMN(int32_t, neighbour7))

using PFRecHitHBHETopologyAlpakaESDataSoA = PFRecHitHBHETopologyAlpakaESDataSoALayout<>;

}

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESRcd_h
#define RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESRcd_h

#include "FWCore/Framework/interface/EventSetupRecordImplementation.h"
#include "FWCore/Framework/interface/DependentRecordImplementation.h"
#include "Geometry/Records/interface/CaloGeometryRecord.h"
#include "Geometry/Records/interface/HcalRecNumberingRecord.h"

class PFRecHitHBHETopologyAlpakaESRcd : public edm::eventsetup::DependentRecordImplementation<PFRecHitHBHETopologyAlpakaESRcd, edm::mpl::Vector<HcalRecNumberingRecord, CaloGeometryRecord>> {};

#endif // RecoParticleFlow_PFRecHitProducer_interface_PFRecHitHBHETopologyAlpakaESRcd_h
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_alpaka_PFRecHitHBHEParamsAlpakaESData_h
#define RecoParticleFlow_PFRecHitProducer_interface_alpaka_PFRecHitHBHEParamsAlpakaESData_h

#include "DataFormats/Portable/interface/alpaka/PortableCollection.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"

#include "RecoParticleFlow/PFRecHitProducer/interface/PFRecHitHBHEParamsAlpakaESData.h"
#include "RecoParticleFlow/PFRecHitProducer/interface/PFRecHitHBHEParamsAlpakaESDataSoA.h"

namespace ALPAKA_ACCELERATOR_NAMESPACE {

using PFRecHitHBHEParamsAlpakaESDataHost = reco::PFRecHitHBHEParamsAlpakaESDataHost;
using PFRecHitHBHEParamsAlpakaESDataDevice = PortableCollection<reco::PFRecHitHBHEParamsAlpakaESDataSoA>;

} // namespace ALPAKA_ACCELERATOR_NAMESPACE

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
#ifndef RecoParticleFlow_PFRecHitProducer_interface_alpaka_PFRecHitHBHETopologyAlpakaESData_h
#define RecoParticleFlow_PFRecHitProducer_interface_alpaka_PFRecHitHBHETopologyAlpakaESData_h

#include "DataFormats/Portable/interface/alpaka/PortableCollection.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"

#include "RecoParticleFlow/PFRecHitProducer/interface/PFRecHitHBHETopologyAlpakaESData.h"
#include "RecoParticleFlow/PFRecHitProducer/interface/PFRecHitHBHETopologyAlpakaESDataSoA.h"

namespace ALPAKA_ACCELERATOR_NAMESPACE {

using PFRecHitHBHETopologyAlpakaESDataHost = reco::PFRecHitHBHETopologyAlpakaESDataHost;
using PFRecHitHBHETopologyAlpakaESDataDevice = PortableCollection<reco::PFRecHitHBHETopologyAlpakaESDataSoA>;

} // namespace ALPAKA_ACCELERATOR_NAMESPACE

#endif
42 changes: 23 additions & 19 deletions RecoParticleFlow/PFRecHitProducer/plugins/BuildFile.xml
Original file line number Diff line number Diff line change
@@ -1,26 +1,30 @@
<library name="RecoParticleFlowPFRecHitProducersPlugins" file="*.cc">
<use name="FWCore/Framework"/>
<use name="FWCore/ParameterSet"/>
<use name="FWCore/Utilities"/>
<use name="alpaka"/>
<use name="HeterogeneousCore/AlpakaTest"/>
<use name="DataFormats/Common"/>
<use name="DataFormats/ParticleFlowReco"/>
<use name="CommonTools/ParticleFlow"/>
<use name="DQMServices/Core"/>
<use name="FWCore/Framework"/>
<use name="FWCore/ParameterSet"/>
<use name="FWCore/Utilities"/>
<use name="alpaka"/>
<use name="HeterogeneousCore/AlpakaTest"/>
<use name="DataFormats/Common"/>
<use name="DataFormats/ParticleFlowReco"/>
<use name="CommonTools/ParticleFlow"/>
<use name="DQMServices/Core"/>
<flags EDM_PLUGIN="1"/>
</library>

<!-- alpaka-based portable plugins -->
<library name="RecoParticleFlowPFRecHitProducersPluginsPortable" file="alpaka/*.cc">
<use name="alpaka"/>
<use name="DataFormats/ParticleFlowReco"/>
<use name="FWCore/Framework"/>
<use name="FWCore/ParameterSet"/>
<use name="FWCore/Utilities"/>
<use name="HeterogeneousCore/AlpakaCore"/>
<use name="HeterogeneousCore/AlpakaInterface"/>
<use name="HeterogeneousCore/AlpakaTest"/>

<use name="alpaka"/>
<use name="DataFormats/ParticleFlowReco"/>
<use name="FWCore/Framework"/>
<use name="FWCore/MessageLogger"/>
<use name="FWCore/ParameterSet"/>
<use name="FWCore/Utilities"/>
<use name="Geometry/CaloGeometry"/>
<use name="Geometry/CaloTopology"/>
<use name="Geometry/Records"/>
<use name="HeterogeneousCore/AlpakaCore"/>
<use name="HeterogeneousCore/AlpakaInterface"/>
<use name="RecoParticleFlow/PFRecHitProducer"/>
<flags ALPAKA_BACKENDS="1"/>
<flags EDM_PLUGIN="1"/>
</library>
</library>
Original file line number Diff line number Diff line change
@@ -0,0 +1,61 @@
#include "FWCore/Framework/interface/EventSetupRecordIntervalFinder.h"
#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
#include "FWCore/ParameterSet/interface/ParameterSet.h"
#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
#include "FWCore/Utilities/interface/Exception.h"
#include "HeterogeneousCore/AlpakaCore/interface/alpaka/ESProducer.h"
#include "HeterogeneousCore/AlpakaCore/interface/alpaka/ModuleFactory.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"
#include "RecoParticleFlow/PFRecHitProducer/interface/JobConfigurationAlpakaRecord.h"
#include "RecoParticleFlow/PFRecHitProducer/interface/alpaka/PFRecHitHBHEParamsAlpakaESData.h"

namespace ALPAKA_ACCELERATOR_NAMESPACE {

class PFRecHitHBHEParamsESProducer : public ESProducer {
public:
PFRecHitHBHEParamsESProducer(edm::ParameterSet const& iConfig) :
energyThresholdsHB_(iConfig.getParameter<std::vector<double>>("energyThresholdsHB")),
energyThresholdsHE_(iConfig.getParameter<std::vector<double>>("energyThresholdsHE")) {

if (energyThresholdsHB_.size() != kMaxDepthHB) {
throw cms::Exception("InvalidConfiguration") << "\"energyThresholdsHB\" must be a cms.vdouble() of size " << kMaxDepthHB;
}

if (energyThresholdsHE_.size() != kMaxDepthHE) {
throw cms::Exception("InvalidConfiguration") << "\"energyThresholdsHE\" must be a cms.vdouble() of size " << kMaxDepthHE;
}

setWhatProduced(this);
}

static void fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
edm::ParameterSetDescription desc;
desc.add<std::string>("appendToDataLabel", "");
desc.add<std::vector<double>>("energyThresholdsHB", {0.1, 0.2, 0.3, 0.3});
desc.add<std::vector<double>>("energyThresholdsHE", {0.1, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2});
descriptions.addWithDefaultLabel(desc);
}

std::unique_ptr<PFRecHitHBHEParamsAlpakaESDataHost> produce(JobConfigurationAlpakaRecord const& iRecord) {
auto product = std::make_unique<PFRecHitHBHEParamsAlpakaESDataHost>(kMaxDepthHB + kMaxDepthHE, cms::alpakatools::host());
for (int idx = 0; idx < kMaxDepthHB; ++idx) {
product->view().energyThresholds()[idx] = energyThresholdsHB_[idx];
}
for (int idx = 0; idx < kMaxDepthHE; ++idx) {
product->view().energyThresholds()[idx+kMaxDepthHB] = energyThresholdsHE_[idx];
}
return product;
}

private:
constexpr static uint8_t kMaxDepthHB = 4;
constexpr static uint8_t kMaxDepthHE = 7;

std::vector<double> energyThresholdsHB_;
std::vector<double> energyThresholdsHE_;
};

} // namespace ALPAKA_ACCELERATOR_NAMESPACE

DEFINE_FWK_EVENTSETUP_ALPAKA_MODULE(PFRecHitHBHEParamsESProducer);
Loading

0 comments on commit d24418b

Please sign in to comment.