Add Alpaka Implementation of PFClusterProducer

DataFormats/ParticleFlowReco/interface/PFClusterHostCollection.h

@@ -0,0 +1,13 @@
+#ifndef DataFormats_ParticleFlowReco_interface_PFClusterHostCollection_h
+#define DataFormats_ParticleFlowReco_interface_PFClusterHostCollection_h
+
+#include "DataFormats/ParticleFlowReco/interface/PFClusterSoA.h"
+#include "DataFormats/Portable/interface/PortableHostCollection.h"
+
+namespace reco {
+
+  using PFClusterHostCollection = PortableHostCollection<PFClusterSoA>;
+
+}  // namespace reco
+
+#endif  // DataFormats_ParticleFlowReco_interface_PFClusterHostCollection_h

DataFormats/ParticleFlowReco/interface/PFClusterSoA.h

@@ -0,0 +1,29 @@
+#ifndef DataFormats_ParticleFlowReco_interface_PFClusterSoA_h
+#define DataFormats_ParticleFlowReco_interface_PFClusterSoA_h
+
+#include "DataFormats/SoATemplate/interface/SoACommon.h"
+#include "DataFormats/SoATemplate/interface/SoALayout.h"
+#include "DataFormats/SoATemplate/interface/SoAView.h"
+
+namespace reco {
+
+  GENERATE_SOA_LAYOUT(PFClusterSoALayout,
+                      SOA_COLUMN(int, depth),
+                      SOA_COLUMN(int, seedRHIdx),
+                      SOA_COLUMN(int, topoId),
+                      SOA_COLUMN(int, rhfracSize),
+                      SOA_COLUMN(int, rhfracOffset),
+                      SOA_COLUMN(float, energy),
+                      SOA_COLUMN(float, x),
+                      SOA_COLUMN(float, y),
+                      SOA_COLUMN(float, z),
+                      SOA_COLUMN(int, topoRHCount),
+                      SOA_SCALAR(int, nTopos),
+                      SOA_SCALAR(int, nSeeds),
+                      SOA_SCALAR(int, nRHFracs),
+                      SOA_SCALAR(int, size)  // nRH
+  )
+  using PFClusterSoA = PFClusterSoALayout<>;
+}  // namespace reco
+
+#endif  // DataFormats_ParticleFlowReco_interface_PFClusterSoA_h

DataFormats/ParticleFlowReco/interface/PFRecHitFractionHostCollection.h

@@ -0,0 +1,11 @@
+#ifndef DataFormats_ParticleFlowReco_interface_PFRecHitFractionHostCollection_h
+#define DataFormats_ParticleFlowReco_interface_PFRecHitFractionHostCollection_h
+
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionSoA.h"
+#include "DataFormats/Portable/interface/PortableHostCollection.h"
+
+namespace reco {
+  using PFRecHitFractionHostCollection = PortableHostCollection<PFRecHitFractionSoA>;
+}
+
+#endif  // DataFormats_ParticleFlowReco_interface_PFRecHitFractionHostCollection_h

DataFormats/ParticleFlowReco/interface/PFRecHitFractionSoA.h

@@ -0,0 +1,18 @@
+#ifndef DataFormats_ParticleFlowReco_interface_PFRecHitFractionSoA_h
+#define DataFormats_ParticleFlowReco_interface_PFRecHitFractionSoA_h
+
+#include "DataFormats/SoATemplate/interface/SoACommon.h"
+#include "DataFormats/SoATemplate/interface/SoALayout.h"
+#include "DataFormats/SoATemplate/interface/SoAView.h"
+
+namespace reco {
+
+  GENERATE_SOA_LAYOUT(PFRecHitFractionSoALayout,
+                      SOA_COLUMN(float, frac),
+                      SOA_COLUMN(int, pfrhIdx),
+                      SOA_COLUMN(int, pfcIdx))
+
+  using PFRecHitFractionSoA = PFRecHitFractionSoALayout<>;
+}  // namespace reco
+
+#endif  // DataFormats_ParticleFlowReco_interface_PFRecHitFractionSoA_h

DataFormats/ParticleFlowReco/interface/alpaka/CaloRecHitDeviceCollection.h

@@ -12,4 +12,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::reco {
 
 }  // namespace ALPAKA_ACCELERATOR_NAMESPACE::reco
 
+// check that the portable device collection for the host device is the same as the portable host collection
+ASSERT_DEVICE_MATCHES_HOST_COLLECTION(reco::CaloRecHitDeviceCollection, reco::CaloRecHitHostCollection);
+
 #endif  // DataFormats_ParticleFlowReco_interface_alpaka_CaloRecHitDeviceCollection_h

DataFormats/ParticleFlowReco/interface/alpaka/PFClusterDeviceCollection.h

@@ -0,0 +1,18 @@
+#ifndef DataFormats_ParticleFlowReco_interface_alpaka_PFClusterDeviceCollection_h
+#define DataFormats_ParticleFlowReco_interface_alpaka_PFClusterDeviceCollection_h
+
+#include "DataFormats/Portable/interface/alpaka/PortableCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/PFClusterSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFClusterHostCollection.h"
+
+namespace ALPAKA_ACCELERATOR_NAMESPACE::reco {
+
+  using ::reco::PFClusterHostCollection;
+
+  using PFClusterDeviceCollection = PortableCollection<::reco::PFClusterSoA>;
+}  // namespace ALPAKA_ACCELERATOR_NAMESPACE::reco
+
+// check that the portable device collection for the host device is the same as the portable host collection
+ASSERT_DEVICE_MATCHES_HOST_COLLECTION(reco::PFClusterDeviceCollection, reco::PFClusterHostCollection);
+
+#endif  // DataFormats_ParticleFlowReco_interface_alpaka_PFClusterDeviceCollection_h

DataFormats/ParticleFlowReco/interface/alpaka/PFRecHitDeviceCollection.h

@@ -13,4 +13,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::reco {
 
 }  // namespace ALPAKA_ACCELERATOR_NAMESPACE::reco
 
+// check that the portable device collection for the host device is the same as the portable host collection
+ASSERT_DEVICE_MATCHES_HOST_COLLECTION(reco::PFRecHitDeviceCollection, reco::PFRecHitHostCollection);
+
 #endif  // DataFormats_ParticleFlowReco_interface_alpaka_PFRecHitDeviceCollection_h

DataFormats/ParticleFlowReco/interface/alpaka/PFRecHitFractionDeviceCollection.h

@@ -0,0 +1,18 @@
+#ifndef DataFormats_ParticleFlowReco_interface_alpaka_PFRecHitFractionDeviceCollection_h
+#define DataFormats_ParticleFlowReco_interface_alpaka_PFRecHitFractionDeviceCollection_h
+
+#include "DataFormats/Portable/interface/alpaka/PortableCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionHostCollection.h"
+
+namespace ALPAKA_ACCELERATOR_NAMESPACE::reco {
+
+  using ::reco::PFRecHitFractionHostCollection;
+
+  using PFRecHitFractionDeviceCollection = PortableCollection<::reco::PFRecHitFractionSoA>;
+}  // namespace ALPAKA_ACCELERATOR_NAMESPACE::reco
+
+// check that the portable device collection for the host device is the same as the portable host collection
+ASSERT_DEVICE_MATCHES_HOST_COLLECTION(reco::PFRecHitFractionDeviceCollection, reco::PFRecHitFractionHostCollection);
+
+#endif  // DataFormats_ParticleFlowReco_interface_alpaka_PFRecHitFractionDeviceCollection_h

DataFormats/ParticleFlowReco/src/alpaka/classes_cuda.h

@@ -2,5 +2,9 @@
 #include "DataFormats/Common/interface/Wrapper.h"
 #include "DataFormats/ParticleFlowReco/interface/CaloRecHitSoA.h"
 #include "DataFormats/ParticleFlowReco/interface/PFRecHitSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFClusterSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionSoA.h"
 #include "DataFormats/ParticleFlowReco/interface/alpaka/CaloRecHitDeviceCollection.h"
 #include "DataFormats/ParticleFlowReco/interface/alpaka/PFRecHitDeviceCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/alpaka/PFClusterDeviceCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/alpaka/PFRecHitFractionDeviceCollection.h"

DataFormats/ParticleFlowReco/src/alpaka/classes_cuda_def.xml

@@ -6,5 +6,13 @@
   <class name="alpaka_cuda_async::reco::PFRecHitDeviceCollection" persistent="false"/>
   <class name="edm::DeviceProduct<alpaka_cuda_async::reco::PFRecHitDeviceCollection>" persistent="false"/>
   <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::reco::PFRecHitDeviceCollection>>" persistent="false"/>
+
+  <class name="alpaka_cuda_async::reco::PFClusterDeviceCollection" persistent="false"/>
+  <class name="edm::DeviceProduct<alpaka_cuda_async::reco::PFClusterDeviceCollection>" persistent="false"/>
+  <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::reco::PFClusterDeviceCollection>>" persistent="false"/>
+
+  <class name="alpaka_cuda_async::reco::PFRecHitFractionDeviceCollection" persistent="false"/>
+  <class name="edm::DeviceProduct<alpaka_cuda_async::reco::PFRecHitFractionDeviceCollection>" persistent="false"/>
+  <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::reco::PFRecHitFractionDeviceCollection>>" persistent="false"/>
 </lcgdict>
 

DataFormats/ParticleFlowReco/src/alpaka/classes_rocm.h

@@ -2,5 +2,9 @@
 #include "DataFormats/Common/interface/Wrapper.h"
 #include "DataFormats/ParticleFlowReco/interface/CaloRecHitSoA.h"
 #include "DataFormats/ParticleFlowReco/interface/PFRecHitSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFClusterSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionSoA.h"
 #include "DataFormats/ParticleFlowReco/interface/alpaka/CaloRecHitDeviceCollection.h"
 #include "DataFormats/ParticleFlowReco/interface/alpaka/PFRecHitDeviceCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/alpaka/PFClusterDeviceCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/alpaka/PFRecHitFractionDeviceCollection.h"

DataFormats/ParticleFlowReco/src/alpaka/classes_rocm_def.xml

@@ -6,4 +6,12 @@
   <class name="alpaka_rocm_async::reco::PFRecHitDeviceCollection" persistent="false"/>
   <class name="edm::DeviceProduct<alpaka_rocm_async::reco::PFRecHitDeviceCollection>" persistent="false"/>
   <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::reco::PFRecHitDeviceCollection>>" persistent="false"/>
+
+  <class name="alpaka_rocm_async::reco::PFClusterDeviceCollection" persistent="false"/>
+  <class name="edm::DeviceProduct<alpaka_rocm_async::reco::PFClusterDeviceCollection>" persistent="false"/>
+  <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::reco::PFClusterDeviceCollection>>" persistent="false"/>
+
+  <class name="alpaka_rocm_async::reco::PFRecHitFractionDeviceCollection" persistent="false"/>
+  <class name="edm::DeviceProduct<alpaka_rocm_async::reco::PFRecHitFractionDeviceCollection>" persistent="false"/>
+  <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::reco::PFRecHitFractionDeviceCollection>>" persistent="false"/>
 </lcgdict>

DataFormats/ParticleFlowReco/src/classes_serial.h

@@ -3,3 +3,8 @@
 #include "DataFormats/ParticleFlowReco/interface/CaloRecHitSoA.h"
 #include "DataFormats/ParticleFlowReco/interface/PFRecHitHostCollection.h"
 #include "DataFormats/ParticleFlowReco/interface/PFRecHitSoA.h"
+
+#include "DataFormats/ParticleFlowReco/interface/PFClusterSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFClusterHostCollection.h"
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionSoA.h"
+#include "DataFormats/ParticleFlowReco/interface/PFRecHitFractionHostCollection.h"

DataFormats/ParticleFlowReco/src/classes_serial_def.xml

@@ -30,4 +30,36 @@
   ]]>
   </read>
   <class name="edm::Wrapper<reco::PFRecHitHostCollection>" splitLevel="0"/>
+
+  <class name="reco::PFClusterSoA"/>
+  <class name="reco::PFClusterSoA::View"/>
+  <class name="reco::PFClusterHostCollection"/>
+  <read
+    sourceClass="reco::PFClusterHostCollection"
+    targetClass="reco::PFClusterHostCollection"
+    version="[1-]"
+    source="reco::PFClusterSoA layout_;"
+    target="buffer_,layout_,view_"
+    embed="false">
+  <![CDATA[
+    reco::PFClusterHostCollection::ROOTReadStreamer(newObj, onfile.layout_);
+  ]]>
+  </read>
+  <class name="edm::Wrapper<reco::PFClusterHostCollection>" splitLevel="0"/>
+
+  <class name="reco::PFRecHitFractionSoA"/>
+  <class name="reco::PFRecHitFractionSoA::View"/>
+  <class name="reco::PFRecHitFractionHostCollection"/>
+  <read
+    sourceClass="reco::PFRecHitFractionHostCollection"
+    targetClass="reco::PFRecHitFractionHostCollection"
+    version="[1-]"
+    source="reco::PFRecHitFractionSoA layout_;"
+    target="buffer_,layout_,view_"
+    embed="false">
+  <![CDATA[
+    reco::PFRecHitFractionHostCollection::ROOTReadStreamer(newObj, onfile.layout_);
+  ]]>
+  </read>
+  <class name="edm::Wrapper<reco::PFRecHitFractionHostCollection>" splitLevel="0"/>
 </lcgdict>

HeterogeneousCore/AlpakaInterface/interface/atomicMaxF.h

@@ -0,0 +1,33 @@
+#ifndef HeterogeneousCore_AlpakaCore_interface_atomicMaxF_h
+#define HeterogeneousCore_AlpakaCore_interface_atomicMaxF_h
+#include <alpaka/alpaka.hpp>
+
+#include "FWCore/Utilities/interface/bit_cast.h"
+#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
+
+#if defined(__CUDA_ARCH__) or defined(__HIP_DEVICE_COMPILE__)
+template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
+static __device__ __forceinline__ float atomicMaxF(const TAcc& acc, float* address, float val) {
+  int ret = __float_as_int(*address);
+  while (val > __int_as_float(ret)) {
+    int old = ret;
+    if ((ret = atomicCAS((int*)address, old, __float_as_int(val))) == old)
+      break;
+  }
+  return __int_as_float(ret);
+}
+#else
+template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
+ALPAKA_FN_ACC ALPAKA_FN_INLINE static float atomicMaxF(const TAcc& acc, float* address, float val) {
+  // CPU implementation uses edm::bit_cast
+  int ret = edm::bit_cast<int>(*address);
+  while (val > edm::bit_cast<float>(ret)) {
+    int old = ret;
+    if ((ret = alpaka::atomicCas(acc, (int*)address, old, edm::bit_cast<int>(val))) == old)
+      break;
+  }
+  return edm::bit_cast<float>(ret);
+}
+#endif  // __CUDA_ARCH__ or __HIP_DEVICE_COMPILE__
+
+#endif  // HeterogeneousCore_AlpakaCore_interface_atomicMaxF_h

RecoParticleFlow/PFClusterProducer/BuildFile.xml

@@ -18,6 +18,15 @@
 <use name="vdt_headers"/>
 <use name="rootmath"/>
 <use name="root"/>
+<use name="eigen"/>
+<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"/>
 <export>
   <lib name="1"/>
 </export>

RecoParticleFlow/PFClusterProducer/interface/PFClusterParamsHostCollection.h

@@ -0,0 +1,14 @@
+#ifndef RecoParticleFlow_PFClusterProducer_interface_PFClusterParamsHostCollection_h
+#define RecoParticleFlow_PFClusterProducer_interface_PFClusterParamsHostCollection_h
+
+#include "DataFormats/Portable/interface/PortableHostCollection.h"
+
+#include "RecoParticleFlow/PFClusterProducer/interface/PFClusterParamsSoA.h"
+
+namespace reco {
+
+  using PFClusterParamsHostCollection = PortableHostCollection<PFClusterParamsSoA>;
+
+}
+
+#endif

RecoParticleFlow/PFClusterProducer/interface/PFClusterParamsSoA.h

@@ -0,0 +1,49 @@
+#ifndef RecoParticleFlow_PFClusterProducer_interface_PFRecHitHBHEParamsSoA_h
+#define RecoParticleFlow_PFClusterProducer_interface_PFRecHitHBHEParamsSoA_h
+
+#include "DataFormats/SoATemplate/interface/SoACommon.h"
+#include "DataFormats/SoATemplate/interface/SoALayout.h"
+#include "DataFormats/SoATemplate/interface/SoAView.h"
+
+namespace reco {
+
+  GENERATE_SOA_LAYOUT(PFClusterParamsSoALayout,
+                      SOA_SCALAR(int32_t, nNeigh),
+                      SOA_SCALAR(float, seedPt2ThresholdEB),
+                      SOA_SCALAR(float, seedPt2ThresholdEE),
+                      SOA_COLUMN(float, seedEThresholdEB_vec),
+                      SOA_COLUMN(float, seedEThresholdEE_vec),
+                      SOA_COLUMN(float, topoEThresholdEB_vec),
+                      SOA_COLUMN(float, topoEThresholdEE_vec),
+                      SOA_SCALAR(float, showerSigma2),
+                      SOA_SCALAR(float, minFracToKeep),
+                      SOA_SCALAR(float, minFracTot),
+                      SOA_SCALAR(uint32_t, maxIterations),
+                      SOA_SCALAR(bool, excludeOtherSeeds),
+                      SOA_SCALAR(float, stoppingTolerance),
+                      SOA_SCALAR(float, minFracInCalc),
+                      SOA_SCALAR(float, minAllowedNormalization),
+                      SOA_COLUMN(float, recHitEnergyNormInvEB_vec),
+                      SOA_COLUMN(float, recHitEnergyNormInvEE_vec),
+                      SOA_SCALAR(float, barrelTimeResConsts_corrTermLowE),
+                      SOA_SCALAR(float, barrelTimeResConsts_threshLowE),
+                      SOA_SCALAR(float, barrelTimeResConsts_noiseTerm),
+                      SOA_SCALAR(float, barrelTimeResConsts_constantTermLowE2),
+                      SOA_SCALAR(float, barrelTimeResConsts_noiseTermLowE),
+                      SOA_SCALAR(float, barrelTimeResConsts_threshHighE),
+                      SOA_SCALAR(float, barrelTimeResConsts_constantTerm2),
+                      SOA_SCALAR(float, barrelTimeResConsts_resHighE2),
+                      SOA_SCALAR(float, endcapTimeResConsts_corrTermLowE),
+                      SOA_SCALAR(float, endcapTimeResConsts_threshLowE),
+                      SOA_SCALAR(float, endcapTimeResConsts_noiseTerm),
+                      SOA_SCALAR(float, endcapTimeResConsts_constantTermLowE2),
+                      SOA_SCALAR(float, endcapTimeResConsts_noiseTermLowE),
+                      SOA_SCALAR(float, endcapTimeResConsts_threshHighE),
+                      SOA_SCALAR(float, endcapTimeResConsts_constantTerm2),
+                      SOA_SCALAR(float, endcapTimeResConsts_resHighE2))
+
+  using PFClusterParamsSoA = PFClusterParamsSoALayout<>;
+
+}  // namespace reco
+
+#endif

RecoParticleFlow/PFClusterProducer/interface/PFClusteringEdgeVarsSoA.h

@@ -0,0 +1,17 @@
+#ifndef RecoParticleFlow_PFRecHitProducer_interface_PFClusteringEdgeVarsSoA_h
+#define RecoParticleFlow_PFRecHitProducer_interface_PFClusteringEdgeVarsSoA_h
+
+#include "DataFormats/SoATemplate/interface/SoACommon.h"
+#include "DataFormats/SoATemplate/interface/SoALayout.h"
+#include "DataFormats/SoATemplate/interface/SoAView.h"
+
+namespace reco {
+
+  GENERATE_SOA_LAYOUT(PFClusteringEdgeVarsSoALayout,
+                      SOA_COLUMN(int, pfrh_edgeIdx),   // needs nRH + 1 allocation
+                      SOA_COLUMN(int, pfrh_edgeList))  // needs nRH + maxNeighbors allocation
+
+  using PFClusteringEdgeVarsSoA = PFClusteringEdgeVarsSoALayout<>;
+}  // namespace reco
+
+#endif

RecoParticleFlow/PFClusterProducer/interface/PFClusteringVarsSoA.h

@@ -0,0 +1,31 @@
+#ifndef RecoParticleFlow_PFClusterProducer_interface_PFClusteringVarsSoA_h
+#define RecoParticleFlow_PFClusterProducer_interface_PFClusteringVarsSoA_h
+
+#include "DataFormats/SoATemplate/interface/SoACommon.h"
+#include "DataFormats/SoATemplate/interface/SoALayout.h"
+#include "DataFormats/SoATemplate/interface/SoAView.h"
+
+namespace reco {
+
+  GENERATE_SOA_LAYOUT(PFClusteringVarsSoALayout,
+                      SOA_COLUMN(int, pfrh_topoId),
+                      SOA_COLUMN(int, pfrh_isSeed),
+                      SOA_COLUMN(int, pfrh_passTopoThresh),
+                      SOA_COLUMN(int, topoSeedCount),
+                      SOA_COLUMN(int, topoRHCount),
+                      SOA_COLUMN(int, seedFracOffsets),
+                      SOA_COLUMN(int, topoSeedOffsets),
+                      SOA_COLUMN(int, topoSeedList),
+                      SOA_SCALAR(int, pcrhFracSize),
+                      SOA_COLUMN(int, rhCount),
+                      SOA_SCALAR(int, nEdges),
+                      SOA_COLUMN(int, rhcount),
+                      SOA_SCALAR(int, nTopos),
+                      SOA_COLUMN(int, topoIds),
+                      SOA_SCALAR(int, nRHFracs),
+                      SOA_COLUMN(int, rhIdxToSeedIdx))
+
+  using PFClusteringVarsSoA = PFClusteringVarsSoALayout<>;
+}  // namespace reco
+
+#endif