CMS 3D CMS Logo

List of all members | Public Member Functions
pixelClustering::ClusterChargeCut< TrackerTraits > Struct Template Reference

#include <ClusterChargeCut.h>

Public Member Functions

template<typename TAcc >
ALPAKA_FN_ACC void operator() (TAcc const &acc, SiPixelDigisSoAView digi_view, SiPixelClustersSoAView clus_view, SiPixelClusterThresholds clusterThresholds, const uint32_t numElements) const
 

Detailed Description

template<typename TrackerTraits>
struct pixelClustering::ClusterChargeCut< TrackerTraits >

Definition at line 20 of file ClusterChargeCut.h.

Member Function Documentation

◆ operator()()

template<typename TrackerTraits >
template<typename TAcc >
ALPAKA_FN_ACC void pixelClustering::ClusterChargeCut< TrackerTraits >::operator() ( TAcc const &  acc,
SiPixelDigisSoAView  digi_view,
SiPixelClustersSoAView  clus_view,
SiPixelClusterThresholds  clusterThresholds,
const uint32_t  numElements 
) const
inline

Definition at line 22 of file ClusterChargeCut.h.

References gpuClustering::adc, ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), cms::cudacompat::atomicAdd(), cms::alpakatools::blockPrefixScan(), ALCARECOTkAlJpsiMuMu_cff::charge, DMR_cfg::chargeCut, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), gpuClustering::endModule, SiPixelClusterThresholds::getThresholdForLayerOnCondition(), caHitNtupletGeneratorKernels::good, mps_fire::i, cms::alpakatools::independent_group_elements(), cms::alpakatools::independent_groups(), pixelClustering::invalidModuleId, phase1PixelTopology::layerStart, pixelClustering::maxNumClustersPerModules, pixelClustering::maxNumModules, SiStripPI::min, callgraph::module, gpuClustering::moduleId, gpuClustering::newclusId, phase1PixelTopology::numberOfModules, gpuClustering::numElements, hltrates_dqm_sourceclient-live_cfg::offset, convertSQLiteXML::ok, cms::alpakatools::once_per_block(), gpuClustering::startBPIX2, and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ws.

27  {
29 
30  auto& charge = alpaka::declareSharedVar<int32_t[maxNumClustersPerModules], __COUNTER__>(acc);
31  auto& ok = alpaka::declareSharedVar<uint8_t[maxNumClustersPerModules], __COUNTER__>(acc);
32  auto& newclusId = alpaka::declareSharedVar<uint16_t[maxNumClustersPerModules], __COUNTER__>(acc);
33 
35 
38 
39  auto endModule = clus_view[0].moduleStart();
41  auto firstPixel = clus_view[1 + module].moduleStart();
42  auto thisModuleId = digi_view[firstPixel].moduleId();
43  while (thisModuleId == invalidModuleId and firstPixel < numElements) {
44  // skip invalid or duplicate pixels
45  ++firstPixel;
46  thisModuleId = digi_view[firstPixel].moduleId();
47  }
48  if (firstPixel >= numElements) {
49  // reached the end of the input while skipping the invalid pixels, nothing left to do
50  break;
51  }
52  if (thisModuleId != clus_view[module].moduleId()) {
53  // reached the end of the module while skipping the invalid pixels, skip this module
54  continue;
55  }
57 
58  uint32_t nclus = clus_view[thisModuleId].clusInModule();
59  if (nclus == 0)
60  return;
61 
63  printf("Warning: too many clusters in module %u in block %u: %u > %d\n",
64  thisModuleId,
65  module,
66  nclus,
68 
69  if (nclus > maxNumClustersPerModules) {
70  // remove excess FIXME find a way to cut charge first....
71  for (auto i : cms::alpakatools::independent_group_elements(acc, firstPixel, numElements)) {
72  if (digi_view[i].moduleId() == invalidModuleId)
73  continue; // not valid
74  if (digi_view[i].moduleId() != thisModuleId)
75  break; // end of module
76  if (digi_view[i].clus() >= maxNumClustersPerModules) {
77  digi_view[i].moduleId() = invalidModuleId;
78  digi_view[i].clus() = invalidModuleId;
79  }
80  }
82  }
83 
84 #ifdef GPU_DEBUG
85  if (thisModuleId % 100 == 1)
87  printf("start cluster charge cut for module %d in block %d\n", thisModuleId, module);
88 #endif
89 
91  for (auto i : cms::alpakatools::independent_group_elements(acc, nclus)) {
92  charge[i] = 0;
93  }
94  alpaka::syncBlockThreads(acc);
95 
96  for (auto i : cms::alpakatools::independent_group_elements(acc, firstPixel, numElements)) {
97  if (digi_view[i].moduleId() == invalidModuleId)
98  continue; // not valid
99  if (digi_view[i].moduleId() != thisModuleId)
100  break; // end of module
101  alpaka::atomicAdd(acc,
102  &charge[digi_view[i].clus()],
103  static_cast<int32_t>(digi_view[i].adc()),
104  alpaka::hierarchy::Threads{});
105  }
106  alpaka::syncBlockThreads(acc);
107 
108  auto chargeCut = clusterThresholds.getThresholdForLayerOnCondition(thisModuleId < startBPIX2);
109 
110  bool good = true;
111  for (auto i : cms::alpakatools::independent_group_elements(acc, nclus)) {
112  newclusId[i] = ok[i] = (charge[i] >= chargeCut) ? 1 : 0;
113  if (0 == ok[i])
114  good = false;
115  }
116 
117  // if all clusters are above threshold, do nothing
118  if (alpaka::syncBlockThreadsPredicate<alpaka::BlockAnd>(acc, good))
119  continue;
120 
121  // renumber
122  auto& ws = alpaka::declareSharedVar<uint16_t[32], __COUNTER__>(acc);
123  // FIXME this value should come from cms::alpakatools::blockPrefixScan itself
124  constexpr uint32_t maxThreads = 1024;
125  auto minClust = std::min(nclus, maxThreads);
126 
128  if constexpr (maxNumClustersPerModules > maxThreads) //only if needed
129  {
130  for (uint32_t offset = maxThreads; offset < nclus; offset += maxThreads) {
132  for (uint32_t i : cms::alpakatools::independent_group_elements(acc, offset, nclus)) {
133  uint32_t prevBlockEnd = (i / maxThreads) * maxThreads - 1;
134  newclusId[i] += newclusId[prevBlockEnd];
135  }
136  alpaka::syncBlockThreads(acc);
137  }
138  }
139  ALPAKA_ASSERT_ACC(nclus >= newclusId[nclus - 1]);
140 
141  clus_view[thisModuleId].clusInModule() = newclusId[nclus - 1];
142 
143  // reassign id
144  for (auto i : cms::alpakatools::independent_group_elements(acc, firstPixel, numElements)) {
145  if (digi_view[i].moduleId() == invalidModuleId)
146  continue; // not valid
147  if (digi_view[i].moduleId() != thisModuleId)
148  break; // end of module
149  if (0 == ok[digi_view[i].clus()])
150  digi_view[i].moduleId() = digi_view[i].clus() = invalidModuleId;
151  else
152  digi_view[i].clus() = newclusId[digi_view[i].clus()] - 1;
153  }
154 
155  // done
156  alpaka::syncBlockThreads(acc);
157  }
158  }
ALPAKA_FN_ACC constexpr bool once_per_block(TAcc const &acc)
ALPAKA_FN_ACC auto independent_group_elements(TAcc const &acc, TArgs... args)
ALPAKA_FN_ACC auto independent_groups(TAcc const &acc, TArgs... args)
chargeCut
Definition: DMR_cfg.py:159
constexpr uint16_t numberOfModules
constexpr int32_t getThresholdForLayerOnCondition(bool isLayer1) const noexcept
constexpr int startBPIX2
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ int32_t *__restrict__ uint32_t numElements
__shared__ uint16_t newclusId[maxNumClustersPerModules]
constexpr int32_t maxNumClustersPerModules
ALPAKA_FN_ACC ALPAKA_FN_INLINE void blockPrefixScan(const TAcc &acc, T const *ci, T *co, int32_t size, T *ws=nullptr)
Definition: prefixScan.h:47
static constexpr uint32_t layerStart[numberOfLayers+1]
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ moduleId
constexpr uint16_t maxNumModules
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61
constexpr uint16_t invalidModuleId
uint16_t *__restrict__ uint16_t const *__restrict__ adc