CMS 3D CMS Logo

Classes | Typedefs | Functions | Variables
gpuPixelDoublets Namespace Reference

Classes

struct  CellCutsT
 

Typedefs

template<typename TrackerTraits >
using CellNeighbors = caStructures::CellNeighborsT< TrackerTraits >
 
template<typename TrackerTraits >
using CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
 
template<typename TrackerTraits >
using CellTracks = caStructures::CellTracksT< TrackerTraits >
 
template<typename TrackerTraits >
using CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
 
template<typename TrackerTraits >
using HitsConstView = typename GPUCACellT< TrackerTraits >::HitsConstView
 
template<typename TrackerTraits >
using OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
 
using PhiBinner = typename TrackingRecHitSoA< TrackerTraits >::PhiBinner
 

Functions

template<typename TrackerTraits >
 __attribute__ ((always_inline)) void fishbone(HitsConstView< TrackerTraits > hh
 
template<typename TrackerTraits >
__device__ __attribute__ ((always_inline)) void doubletsFromHisto(uint32_t nPairs
 
 __syncthreads ()
 
 assert (offsets)
 
 for (int i=first;i< nHits - isOuterHitOfCell.offset;i+=gridDim.x *blockDim.x) isOuterHitOfCell.container[i].reset()
 
 for (int idy=firstY, nt=nHits - offset;idy< nt;idy+=gridDim.y *blockDim.y)
 
 if (0==first)
 
 if (threadIdx.y==0 &&threadIdx.x==0)
 

Variables

uint32_t cc [maxCellsPerHit]
 
int CellNeighborsVector< TrackerTraits > * cellNeighbors
 
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > * cellNeighborsContainer
 
GPUCACellT< TrackerTraits > * cells
 
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > * cellTracks
 
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > CellTracks< TrackerTraits > * cellTracksContainer
 
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const int32_t bool checkTrack
 
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int CellCutsT< TrackerTraits > *const cuts
 
uint16_t d [maxCellsPerHit]
 
const bool doPtCut = cuts.doPtCut_
 
const bool doZ0Cut = cuts.doZ0Cut_
 
int first = blockIdx.x * blockDim.x + threadIdx.x
 
auto firstX = threadIdx.x
 
auto firstY = threadIdx.y + blockIdx.y * blockDim.y
 
constexpr auto getDoubletsFromHistoMaxBlockSize = 64
 
constexpr auto getDoubletsFromHistoMinBlocksPerMP = 16
 
const float hardPtCut = cuts.ptCut_
 
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > hh
 
auto idy = blockIdx.y * blockDim.y + threadIdx.y
 
__shared__ uint32_t innerLayerCumulativeSize [TrackerTraits::nPairs]
 
auto const isOuterHitOfCell = isOuterHitOfCellWrap.container
 
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const isOuterHitOfCellWrap
 
uint8_t l [maxCellsPerHit]
 
auto layerSize = [=](uint8_t li) { return offsets[li + 1] - offsets[li]; }
 
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int maxNumOfDoublets
 
const float minRadius = hardPtCut * 87.78f
 
const float minRadius2T4 = 4.f * minRadius * minRadius
 
float n [maxCellsPerHit]
 
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int nActualPairs
 
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ nCells
 
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const int32_t nHits
 
__shared__ uint32_t ntot
 
int32_t offset = isOuterHitOfCellWrap.offset
 
uint32_t const *__restrict__ offsets = hh.hitsLayerStart().data()
 
uint32_t pairLayerId = 0
 
auto const &__restrict__ phiBinner = hh.phiBinner()
 
auto stride = blockDim.x
 
float x [maxCellsPerHit]
 
float y [maxCellsPerHit]
 
float z [maxCellsPerHit]
 
const float z0cut = cuts.z0Cut_
 

Typedef Documentation

◆ CellNeighbors

template<typename TrackerTraits >
using gpuPixelDoublets::CellNeighbors = typedef caStructures::CellNeighborsT<TrackerTraits>

Definition at line 20 of file gpuFishbone.h.

◆ CellNeighborsVector

template<typename TrackerTraits >
using gpuPixelDoublets::CellNeighborsVector = typedef caStructures::CellNeighborsVectorT<TrackerTraits>

Definition at line 24 of file gpuFishbone.h.

◆ CellTracks

template<typename TrackerTraits >
using gpuPixelDoublets::CellTracks = typedef caStructures::CellTracksT<TrackerTraits>

Definition at line 22 of file gpuFishbone.h.

◆ CellTracksVector

template<typename TrackerTraits >
using gpuPixelDoublets::CellTracksVector = typedef caStructures::CellTracksVectorT<TrackerTraits>

Definition at line 26 of file gpuFishbone.h.

◆ HitsConstView

template<typename TrackerTraits >
using gpuPixelDoublets::HitsConstView = typedef typename GPUCACellT<TrackerTraits>::HitsConstView

Definition at line 30 of file gpuFishbone.h.

◆ OuterHitOfCell

template<typename TrackerTraits >
using gpuPixelDoublets::OuterHitOfCell = typedef caStructures::OuterHitOfCellT<TrackerTraits>

Definition at line 28 of file gpuFishbone.h.

◆ PhiBinner

using gpuPixelDoublets::PhiBinner = typedef typename TrackingRecHitSoA<TrackerTraits>::PhiBinner

Definition at line 144 of file gpuPixelDoubletsAlgos.h.

Function Documentation

◆ __attribute__() [1/2]

template<typename TrackerTraits >
gpuPixelDoublets::__attribute__ ( (always_inline)  )
inline

◆ __attribute__() [2/2]

template<typename TrackerTraits >
__device__ gpuPixelDoublets::__attribute__ ( (always_inline)  )
inline

◆ __syncthreads()

gpuPixelDoublets::__syncthreads ( )
inline

Definition at line 132 of file cudaCompat.h.

132 {}

◆ assert()

gpuPixelDoublets::assert ( offsets  )

◆ for() [1/2]

gpuPixelDoublets::for ( )

Definition at line 174 of file gpuPixelDoubletsAlgos.h.

References cms::cudacompat::__ldg(), funct::abs(), assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::atomicSub(), newFWLiteAna::bin, cellNeighbors, cells, cellTracks, cuts, HLT_2023v12_cff::doClusterCut, doPtCut, doZ0Cut, MillePedeFileConverter_cfg::e, first, full, hh, mps_fire::i, SurfaceOrientation::inner, innerLayerCumulativeSize, isOuterHitOfCell, dqmiolumiharvest::j, dqmdumpme::k, GetRecoTauVFromDQM_MC_cff::kk, pixelTopology::last_barrel_layer, TrajectorySeedProducer_cfi::layerPairs, gpuClustering::maxNumModules, maxNumOfDoublets, phase1PixelTopology::maxr, phase1PixelTopology::maxz, SiStripPI::min, minRadius2T4, phase1PixelTopology::minz, LaserClient_cfi::nbins, nCells, TrackingDataMCValidation_Standalone_cff::nmin, phase1PixelTopology::nPairs, offsets, SurfaceOrientation::outer, AlCaHLTBitMon_ParallelJobs::p, pairLayerId, phiBinner, submitPVValidationJobs::ptcut, short2phi(), stride, compareTotals::tot, and z0cut.

174  {
176  ;
177  --pairLayerId; // move to lower_bound ??
178 
182 
184  uint8_t outer = TrackerTraits::layerPairs[2 * pairLayerId + 1];
185  assert(outer > inner);
186 
187  auto hoff = PhiBinner::histOff(outer);
188  auto i = (0 == pairLayerId) ? j : j - innerLayerCumulativeSize[pairLayerId - 1];
189  i += offsets[inner];
190 
191  assert(i >= offsets[inner]);
192  assert(i < offsets[inner + 1]);
193 
194  // found hit corresponding to our cuda thread, now do the job
195 
196  if (hh[i].detectorIndex() > gpuClustering::maxNumModules)
197  continue; // invalid
198 
199  /* maybe clever, not effective when zoCut is on
200  auto bpos = (mi%8)/4; // if barrel is 1 for z>0
201  auto fpos = (outer>3) & (outer<7);
202  if ( ((inner<3) & (outer>3)) && bpos!=fpos) continue;
203  */
204 
205  auto mez = hh[i].zGlobal();
206 
208  continue;
209 
210  if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts.clusterCut(hh, i))
211  continue;
212 
213  auto mep = hh[i].iphi();
214  auto mer = hh[i].rGlobal();
215 
216  auto ptcut = [&](int j, int16_t idphi) {
217  auto r2t4 = minRadius2T4;
218  auto ri = mer;
219  auto ro = hh[j].rGlobal();
220  auto dphi = short2phi(idphi);
221  return dphi * dphi * (r2t4 - ri * ro) > (ro - ri) * (ro - ri);
222  };
223  auto z0cutoff = [&](int j) {
224  auto zo = hh[j].zGlobal();
225  auto ro = hh[j].rGlobal();
226  auto dr = ro - mer;
227  return dr > TrackerTraits::maxr[pairLayerId] || dr < 0 || std::abs((mez * ro - mer * zo)) > z0cut * dr;
228  };
229 
230  auto iphicut = cuts.phiCuts[pairLayerId];
231 
232  auto kl = PhiBinner::bin(int16_t(mep - iphicut));
233  auto kh = PhiBinner::bin(int16_t(mep + iphicut));
234  auto incr = [](auto& k) { return k = (k + 1) % PhiBinner::nbins(); };
235 
236 #ifdef GPU_DEBUG
237  int tot = 0;
238  int nmin = 0;
239  int tooMany = 0;
240 #endif
241 
242  auto khh = kh;
243  incr(khh);
244  for (auto kk = kl; kk != khh; incr(kk)) {
245 #ifdef GPU_DEBUG
246  if (kk != kl && kk != kh)
247  nmin += phiBinner.size(kk + hoff);
248 #endif
249  auto const* __restrict__ p = phiBinner.begin(kk + hoff);
250  auto const* __restrict__ e = phiBinner.end(kk + hoff);
251  p += first;
252  for (; p < e; p += stride) {
253  auto oi = __ldg(p);
254  assert(oi >= offsets[outer]);
255  assert(oi < offsets[outer + 1]);
256  auto mo = hh[oi].detectorIndex();
257 
259  continue; // invalid
260 
261  if (doZ0Cut && z0cutoff(oi))
262  continue;
263  auto mop = hh[oi].iphi();
264  uint16_t idphi = std::min(std::abs(int16_t(mop - mep)), std::abs(int16_t(mep - mop)));
265  if (idphi > iphicut)
266  continue;
267 
268  if (doClusterCut && cuts.zSizeCut(hh, i, oi))
269  continue;
270  if (doPtCut && ptcut(oi, idphi))
271  continue;
272 
273  auto ind = atomicAdd(nCells, 1);
274  if (ind >= maxNumOfDoublets) {
275  atomicSub(nCells, 1);
276  break;
277  } // move to SimpleVector??
278  // int layerPairId, int doubletId, int innerHitId, int outerHitId)
279  cells[ind].init(*cellNeighbors, *cellTracks, hh, pairLayerId, i, oi);
280  isOuterHitOfCell[oi].push_back(ind);
281 #ifdef GPU_DEBUG
282  if (isOuterHitOfCell[oi].full())
283  ++tooMany;
284  ++tot;
285 #endif
286  }
287  }
288 // #endif
289 #ifdef GPU_DEBUG
290  if (tooMany > 0)
291  printf("OuterHitOfCell full for %d in layer %d/%d, %d,%d %d, %d %.3f %.3f\n",
292  i,
293  inner,
294  outer,
295  nmin,
296  tot,
297  tooMany,
298  iphicut,
301 #endif
302  } // loop in block...
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const isOuterHitOfCell
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int maxNumOfDoublets
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const * cellNeighbors
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const * cellTracks
T1 atomicSub(T1 *a, T2 b)
Definition: cudaCompat.h:73
uint32_t const *__restrict__ offsets
__shared__ uint32_t innerLayerCumulativeSize[TrackerTraits::nPairs]
constexpr uint16_t last_barrel_layer
auto const &__restrict__ phiBinner
Definition: GenABIO.cc:168
Abs< T >::type abs(const T &t)
Definition: Abs.h:22
constexpr uint16_t maxNumModules
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ nCells
T __ldg(T const *x)
Definition: cudaCompat.h:137
constexpr float short2phi(short x)
Definition: approx_atan2.h:285
constexpr float minz[nPairs]
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > hh
constexpr float maxz[nPairs]
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61
constexpr float maxr[nPairs]

◆ for() [2/2]

gpuPixelDoublets::for ( int  idy = firstY)

Definition at line 53 of file gpuFishbone.h.

References cms::cudacompat::blockDim, fftjetpileupestimator_calo_uncalib_cfi::c0, cc, cells, checkTrack, d, f, firstX, hh, idy, isOuterHitOfCell, l, n, x, y, and z.

53  {
54  auto const& vc = isOuterHitOfCell[idy];
55  auto size = vc.size();
56  if (size < 2)
57  continue;
58  // if alligned kill one of the two.
59  // in principle one could try to relax the cut (only in r-z?) for jumping-doublets
60  auto const& c0 = cells[vc[0]];
61  auto xo = c0.outer_x(hh);
62  auto yo = c0.outer_y(hh);
63  auto zo = c0.outer_z(hh);
64  auto sg = 0;
65  for (int32_t ic = 0; ic < size; ++ic) {
66  auto& ci = cells[vc[ic]];
67  if (ci.unused())
68  continue; // for triplets equivalent to next
69  if (checkTrack && ci.tracks().empty())
70  continue;
71  cc[sg] = vc[ic];
72  l[sg] = ci.layerPairId();
73  d[sg] = ci.inner_detIndex(hh);
74  x[sg] = ci.inner_x(hh) - xo;
75  y[sg] = ci.inner_y(hh) - yo;
76  z[sg] = ci.inner_z(hh) - zo;
77  n[sg] = x[sg] * x[sg] + y[sg] * y[sg] + z[sg] * z[sg];
78  ++sg;
79  }
80  if (sg < 2)
81  continue;
82  // here we parallelize
83  for (int32_t ic = firstX; ic < sg - 1; ic += blockDim.x) {
84  auto& ci = cells[cc[ic]];
85  for (auto jc = ic + 1; jc < sg; ++jc) {
86  auto& cj = cells[cc[jc]];
87  // must be different detectors
88  // if (d[ic]==d[jc]) continue;
89  auto cos12 = x[ic] * x[jc] + y[ic] * y[jc] + z[ic] * z[jc];
90  if (d[ic] != d[jc] && cos12 * cos12 >= 0.99999f * (n[ic] * n[jc])) {
91  // alligned: kill farthest (prefer consecutive layers)
92  // if same layer prefer farthest (longer level arm) and make space for intermediate hit
93  bool sameLayer = l[ic] == l[jc];
94  if (n[ic] > n[jc]) {
95  if (sameLayer) {
96  cj.kill(); // closest
97  ci.setFishbone(cj.inner_hit_id(), cj.inner_z(hh), hh);
98  } else {
99  ci.kill(); // farthest
100  // break; // removed to improve reproducibility. keep it for reference and tests
101  }
102  } else {
103  if (!sameLayer) {
104  cj.kill(); // farthest
105  } else {
106  ci.kill(); // closest
107  cj.setFishbone(ci.inner_hit_id(), ci.inner_z(hh), hh);
108  // break; // removed to improve reproducibility. keep it for reference and tests
109  }
110  }
111  }
112  } // cj
113  } // ci
114  } // hits
size
Write out results.
float z[maxCellsPerHit]
Definition: gpuFishbone.h:48
float n[maxCellsPerHit]
Definition: gpuFishbone.h:48
uint16_t d[maxCellsPerHit]
Definition: gpuFishbone.h:50
float x[maxCellsPerHit]
Definition: gpuFishbone.h:48
uint32_t cc[maxCellsPerHit]
Definition: gpuFishbone.h:49
const dim3 blockDim
Definition: cudaCompat.h:30
float y[maxCellsPerHit]
Definition: gpuFishbone.h:48
GPUCACellT< TrackerTraits > * cells
Definition: gpuFishbone.h:34
uint8_t l[maxCellsPerHit]
Definition: gpuFishbone.h:51
auto const isOuterHitOfCell
Definition: gpuFishbone.h:41
double f[11][100]
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > hh
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const int32_t bool checkTrack
Definition: gpuFishbone.h:38

◆ if() [1/2]

gpuPixelDoublets::if ( = first)

Definition at line 22 of file gpuPixelDoublets.h.

References assert(), cellNeighbors, cellNeighborsContainer, cellTracks, cellTracksContainer, and mps_fire::i.

22  {
23  cellNeighbors->construct(TrackerTraits::maxNumOfActiveDoublets, cellNeighborsContainer);
24  cellTracks->construct(TrackerTraits::maxNumOfActiveDoublets, cellTracksContainer);
25  auto i = cellNeighbors->extend();
26  assert(0 == i);
27  (*cellNeighbors)[0].reset();
28  i = cellTracks->extend();
29  assert(0 == i);
30  (*cellTracks)[0].reset();
31  }
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > * cellTracks
assert(be >=bs)
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > * cellNeighborsContainer
int CellNeighborsVector< TrackerTraits > * cellNeighbors
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > CellTracks< TrackerTraits > * cellTracksContainer

◆ if() [2/2]

gpuPixelDoublets::if ( threadIdx.  y = = 0 && threadIdx.x == 0)

Variable Documentation

◆ cc

uint32_t gpuPixelDoublets::cc[maxCellsPerHit]

Definition at line 49 of file gpuFishbone.h.

Referenced by BPHWriteSpecificDecay::addTrackModes(), HcalLutAnalyzer::analyze(), PPSPixelDigiAnalyzer::analyze(), AntiElectronIDMVA6< TauType, ElectronType >::AntiElectronIDMVA6(), PhotonEnergyCorrector::applyCrackCorrection(), EgammaSCEnergyCorrectionAlgo::applyCrackCorrection(), BeamHaloPropagatorESProducer::BeamHaloPropagatorESProducer(), l1t::stage2::BMTFTokens::BMTFTokens(), FFTJetPileupEstimator::calibrateFromDB(), CaloGeometryBuilder::CaloGeometryBuilder(), CaloGeometryDBEP< T, U >::CaloGeometryDBEP(), CaloGeometryEP< T, D >::CaloGeometryEP(), l1t::stage2::CaloLayer1Tokens::CaloLayer1Tokens(), CaloSimParametersESModule::CaloSimParametersESModule(), CaloSimulationConstantsESModule::CaloSimulationConstantsESModule(), l1t::stage2::CaloTokens::CaloTokens(), l1t::stage1::CaloTokens::CaloTokens(), CaloTowerConstituentsMapBuilder::CaloTowerConstituentsMapBuilder(), CaloTowerHardcodeGeometryEP::CaloTowerHardcodeGeometryEP(), CaloTPGTranscoderULUTs::CaloTPGTranscoderULUTs(), CaloTrkProcessingBuilder::CaloTrkProcessingBuilder(), CastorDbProducer::CastorDbProducer(), cert_plot(), FWGeometryTableManager::checkRegionOfInterest(), egamma::classBasedElectronEnergy(), ClusterShapeHitFilterESProducer::ClusterShapeHitFilterESProducer(), EcalUncalibRecHitTimingCCAlgo::computeCC(), TrajSeedMatcher::Configuration::Configuration(), hitfit::Fourvec_Constrainer::constrain(), ConvertingESProducerT< Record, Target, Source >::ConvertingESProducerT(), ConvertingESProducerWithDependenciesT< CombinedRecord< DepsRecords... >, Target, Dependencies... >::ConvertingESProducerWithDependenciesT(), p2eg::Cluster::crystaliEtaFromCardRegionInfo(), p2eg::Cluster::crystaliPhiFromCardRegionInfo(), p2eg::SimpleCaloHit::crystalLocaliEta(), p2eg::SimpleCaloHit::crystalLocaliPhi(), CSCConditions::CSCConditions(), CSCGeometryESModule::CSCGeometryESModule(), CTPPSBeamParametersFromLHCInfoESSource::CTPPSBeamParametersFromLHCInfoESSource(), CTPPSInterpolatedOpticalFunctionsESSource::CTPPSInterpolatedOpticalFunctionsESSource(), magneticfield::DD4hep_VolumeBasedMagneticFieldESProducerFromDB::DD4hep_VolumeBasedMagneticFieldESProducerFromDB(), GflashHadronShowerProfile::doCholeskyReduction(), DreamSensitiveDetectorBuilder::DreamSensitiveDetectorBuilder(), DTCombinatorialExtendedPatternReco::DTCombinatorialExtendedPatternReco(), DTCombinatorialPatternReco4D::DTCombinatorialPatternReco4D(), DTConfigDBProducer::DTConfigDBProducer(), DTFakeT0ESProducer::DTFakeT0ESProducer(), DTGeometryESModule::DTGeometryESModule(), DTGeometryESProducer::DTGeometryESProducer(), DTLinearDriftFromDBAlgo::DTLinearDriftFromDBAlgo(), DTMeantimerPatternReco::DTMeantimerPatternReco(), DTMeantimerPatternReco4D::DTMeantimerPatternReco4D(), DTRecHitBaseAlgo::DTRecHitBaseAlgo(), DTRefitAndCombineReco4D::DTRefitAndCombineReco4D(), dtCalibration::DTTTrigConstantShift::DTTTrigConstantShift(), dtCalibration::DTTTrigFillWithAverage::DTTTrigFillWithAverage(), dtCalibration::DTTTrigMatchRPhi::DTTTrigMatchRPhi(), dtCalibration::DTTTrigResidualCorrection::DTTTrigResidualCorrection(), DTTTrigSyncFromDB::DTTTrigSyncFromDB(), dtCalibration::DTTTrigT0SegCorrection::DTTTrigT0SegCorrection(), dtCalibration::DTVDriftSegment::DTVDriftSegment(), VVIObjDetails::dzero(), VVIObjFDetails::dzero(), sistripvvi::VVIObjDetails::dzero(), EcalBasicClusterLocalContCorrection::EcalBasicClusterLocalContCorrection(), EcalLaserCorrectionService::EcalLaserCorrectionService(), EcalLaserCorrectionServiceMC::EcalLaserCorrectionServiceMC(), EcalRegionCablingESProducer::EcalRegionCablingESProducer(), EcalSensitiveDetectorBuilder::EcalSensitiveDetectorBuilder(), EcalSignalGenerator< EEDigitizerTraits >::EcalSignalGenerator(), EcalSimParametersESModule::EcalSimParametersESModule(), EcalTBGeometryBuilder::EcalTBGeometryBuilder(), EcalTPGDBCopy::EcalTPGDBCopy(), EG8XObjectUpdateModifier::EG8XObjectUpdateModifier(), EgammaRegressionContainer::EgammaRegressionContainer(), EGEtScaleSysModifier::EGEtScaleSysModifier(), EGExtraInfoModifierFromValueMaps< MapType, OutputType >::EGExtraInfoModifierFromValueMaps(), EGFull5x5ShowerShapeModifierFromValueMaps::EGFull5x5ShowerShapeModifierFromValueMaps(), EGRegressionModifierCondTokens::EGRegressionModifierCondTokens(), EGRegressionModifierDRN::EGRegressionModifierDRN(), EGRegressionModifierV1::EGRegressionModifierV1(), EGRegressionModifierV2::EGRegressionModifierV2(), EGRegressionModifierV3::EGRegressionModifierV3(), ElectronHcalHelper::ElectronHcalHelper(), ElectronSeedGenerator::ElectronSeedGenerator(), l1t::stage2::EMTFTokens::EMTFTokens(), LumiCalculator::endRun(), ES_TTStubAlgorithm_cbc3< T >::ES_TTStubAlgorithm_cbc3(), ES_TTStubAlgorithm_official< T >::ES_TTStubAlgorithm_official(), helper::BFieldIsolationAlgorithmSetup< Alg >::esConsumes(), RegressionHelper::ESGetTokens::ESGetTokens(), EcalClusterLazyToolsBase::ESGetTokens::ESGetTokens(), ESRecHitWorker::ESRecHitWorker(), FFTJetCorrectionESProducer< CT >::FFTJetCorrectionESProducer(), FFTJetLookupTableESProducer< CT >::FFTJetLookupTableESProducer(), FiberSensitiveDetectorBuilder::FiberSensitiveDetectorBuilder(), BPHWriteSpecificDecay::fill(), TimingSD::fillHits(), TkAccumulatingSensitiveDetector::fillHits(), MuonSensitiveDetector::fillHits(), TotemSD::fillHits(), FP420SD::fillHits(), CaloSD::fillHits(), CaloSteppingAction::fillHits(), NtupleManager::FillMeasurements(), CaloSteppingAction::fillPassiveHits(), mkfit::MkBuilder::find_tracks_in_layers(), for(), FWRecoGeometryESProducer::FWRecoGeometryESProducer(), FWTGeoRecoGeometryESProducer::FWTGeoRecoGeometryESProducer(), SiStripGainESProducer::GainGetterT< Record >::GainGetterT(), GeneralTracksImporter::GeneralTracksImporter(), GenericSimClusterMapper::GenericSimClusterMapper(), p2eg::getCard_iEtaMax(), p2eg::getCard_iEtaMin(), p2eg::getCard_iPhiMax(), p2eg::getCard_iPhiMin(), p2eg::getCard_refCrystal_iEta(), p2eg::getCard_refCrystal_iPhi(), CaloGeometryHelper::getClosestCell(), GlobalDetLayerGeometryESProducer::GlobalDetLayerGeometryESProducer(), GlobalTrackingGeometryESProducer::GlobalTrackingGeometryESProducer(), l1t::stage2::GMTTokens::GMTTokens(), GsfTrajectoryFitterESProducer::GsfTrajectoryFitterESProducer(), GsfTrajectorySmootherESProducer::GsfTrajectorySmootherESProducer(), l1t::stage2::GTTokens::GTTokens(), HcalAlignmentEP::HcalAlignmentEP(), HcalDbProducer::HcalDbProducer(), HcalDDDGeometryEP::HcalDDDGeometryEP(), HcalDDDRecConstantsESModule::HcalDDDRecConstantsESModule(), HcalDDDSimConstantsESModule::HcalDDDSimConstantsESModule(), HcalHardcodeGeometryEP::HcalHardcodeGeometryEP(), HcalParametersESModule::HcalParametersESModule(), HcalSensitiveDetectorBuilder::HcalSensitiveDetectorBuilder(), HcalSimParametersESModule::HcalSimParametersESModule(), HcalSimulationConstantsESModule::HcalSimulationConstantsESModule(), HcalTB02ParametersESModule::HcalTB02ParametersESModule(), HcalTB02SensitiveDetectorBuilder::HcalTB02SensitiveDetectorBuilder(), HcalTB06BeamDetectorBuilder::HcalTB06BeamDetectorBuilder(), HcalTB06ParametersESModule::HcalTB06ParametersESModule(), HcalTPGCoderULUT::HcalTPGCoderULUT(), HFNoseSensitiveDetectorBuilder::HFNoseSensitiveDetectorBuilder(), HGCalGeometryESProducer::HGCalGeometryESProducer(), HGCalNumberingInitialization::HGCalNumberingInitialization(), HGCalParametersESModule::HGCalParametersESModule(), HGCalSensitiveDetectorBuilder::HGCalSensitiveDetectorBuilder(), HGCalTBGeometryESProducer::HGCalTBGeometryESProducer(), HGCalTBNumberingInitialization::HGCalTBNumberingInitialization(), HGCalTBParametersESModule::HGCalTBParametersESModule(), HGCalTriggerGeometryESProducer::HGCalTriggerGeometryESProducer(), HGCRecHitNavigator< D1, hgcee, D2, hgchef, D3, hgcheb >::HGCRecHitNavigator(), HGCScintillatorSensitiveDetectorBuilder::HGCScintillatorSensitiveDetectorBuilder(), HGCSensitiveDetectorBuilder::HGCSensitiveDetectorBuilder(), HLTDQMTagAndProbeEff< TagType, TagCollType, ProbeType, ProbeCollType >::HLTDQMTagAndProbeEff(), HLTExclDiJetFilter< T >::hltFilter(), HLTTauDQML1Plotter::HLTTauDQML1Plotter(), p2eg::SimpleCaloHit::isInCard(), L1EGCrystalClusterEmulatorProducer::SimpleCaloHit::isInCard(), L1EGCrystalClusterEmulatorProducer::SimpleCaloHit::isInCardAndRegion(), p2eg::isValidCard(), ConvertingESProducerWithDependenciesT< CombinedRecord< DepsRecords... >, Target, Dependencies... >::WalkConsumes< N >::iterate(), L1CondDBPayloadWriterExt::L1CondDBPayloadWriterExt(), L1ConfigOnlineProdBase< L1HtMissScaleRcd, L1CaloEtScale >::L1ConfigOnlineProdBase(), L1ObjectKeysOnlineProdBase::L1ObjectKeysOnlineProdBase(), l1t::L1TDigiToRaw::L1TDigiToRaw(), L1TriggerKeyOnlineProd::L1TriggerKeyOnlineProd(), L1TriggerKeyOnlineProdExt::L1TriggerKeyOnlineProdExt(), MagneticFieldMapESProducer::MagneticFieldMapESProducer(), SiStripGainESProducer::make_GainGetter(), calogeometryDBEPimpl::GeometryTraits< HGCalGeometry, true >::makeToken(), calogeometryDBEPimpl::GeometryTraits< T, U::writeFlag()>::makeToken(), calogeometryDBEPimpl::GeometryTraits< T, false >::makeToken(), calogeometryDBEPimpl::AdditionalTokens< CaloTowerGeometry >::makeTokens(), calogeometryDBEPimpl::AdditionalTokens< HcalGeometry >::makeTokens(), calogeometryDBEPimpl::AdditionalTokens< HGCalGeometry >::makeTokens(), ME0GeometryESModule::ME0GeometryESModule(), MisalignedTrackerESProducer::MisalignedTrackerESProducer(), MkFitGeometryESProducer::MkFitGeometryESProducer(), MTDDetLayerGeometryESProducer::MTDDetLayerGeometryESProducer(), MTDDigiGeometryESModule::MTDDigiGeometryESModule(), MTDGeometricTimingDetESModule::MTDGeometricTimingDetESModule(), MTDParametersESModule::MTDParametersESModule(), MTDTimeCalibESProducer::MTDTimeCalibESProducer(), MultipleScatteringParametrisationMakerESProducer::MultipleScatteringParametrisationMakerESProducer(), MultiRecHitCollectorESProducer::MultiRecHitCollectorESProducer(), MuonDetLayerGeometryESProducer::MuonDetLayerGeometryESProducer(), MuonGeometryConstantsESModule::MuonGeometryConstantsESModule(), MuonOffsetESProducer::MuonOffsetESProducer(), MuonSegmentCompatibilityCut::MuonSegmentCompatibilityCut(), MuonTrackCut::MuonTrackCut(), MVAVariableHelper::MVAVariableHelper(), pat::helper::NamedUserDataLoader< pat::helper::AddUserCand >::NamedUserDataLoader(), NavigationSchoolESProducer::NavigationSchoolESProducer(), pat::ObjectModifier< T >::ObjectModifier(), OfflineToTransientBeamSpotESProducer::OfflineToTransientBeamSpotESProducer(), TwoTrackMinimumDistanceHelixLine::oneIteration(), pat::PATMETSlimmer::OneMETShift::OneMETShift(), OnlineBeamSpotESProducer::OnlineBeamSpotESProducer(), EGRegressionModifierDRN::partVars< T >::partVars(), PFClusterBuilderBase::PFClusterBuilderBase(), PFClusterEMEnergyCorrector::PFClusterEMEnergyCorrector(), PFClusterFromHGCalTrackster::PFClusterFromHGCalTrackster(), PFClusterProducer::PFClusterProducer(), PFEcalBarrelRecHitCreator::PFEcalBarrelRecHitCreator(), PFEcalEndcapRecHitCreator::PFEcalEndcapRecHitCreator(), PFMultiDepthClusterizer::PFMultiDepthClusterizer(), PFMultiDepthClusterProducer::PFMultiDepthClusterProducer(), PFRecHitCreatorBase::PFRecHitCreatorBase(), PFRecHitDualNavigator< D1, barrel, D2, endcap >::PFRecHitDualNavigator(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitECALParamsESProducer::PFRecHitECALParamsESProducer(), PFRecHitProducer::PFRecHitProducer(), PFRecHitQTestHCALThresholdVsDepth::PFRecHitQTestHCALThresholdVsDepth(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitTopologyESProducer< CAL >::PFRecHitTopologyESProducer(), PixelCPEFastESProducerT< TrackerTraits >::PixelCPEFastESProducerT(), PlotDebugFPIX_XYMap(), PlotPixelMultVtxPos(), FastTimerService::postESModule(), FastTimerService::preESModule(), DDEcalEndcapTrapX::print(), DDEcalEndcapTrap::print(), p2eg::Cluster::printClusterInfo(), printPlot(), PixelTracksProducer::produce(), CastorClusterProducer::produce(), Phase2L1CaloEGammaEmulator::produce(), JetCoreMCtruthSeedGenerator::produce(), DeepCoreSeedGenerator::produce(), L1EGCrystalClusterEmulatorProducer::produce(), trklet::ProducerChannelAssignment::ProducerChannelAssignment(), trackerTFP::ProducerDemonstrator::ProducerDemonstrator(), trackerTFP::ProducerES::ProducerES(), trackerTFP::ProducerFormatsKF::ProducerFormatsKF(), hph::ProducerHPH::ProducerHPH(), trackerTFP::ProducerLayerEncoding::ProducerLayerEncoding(), trackerDTC::ProducerLayerEncoding::ProducerLayerEncoding(), tt::ProducerSetup::ProducerSetup(), p2eg::Cluster::realEta(), p2eg::Cluster::realPhi(), MultiHitGeneratorFromChi2::refit2Hits(), HITrackingRegionProducer::regions(), HITrackingRegionForPrimaryVtxProducer::regions(), l1t::stage2::CaloLayer1Setup::registerConsumes(), l1t::stage2::GMTSetup::registerConsumes(), l1t::stage2::EMTFSetup::registerConsumes(), l1t::stage2::CaloSetup::registerConsumes(), l1t::stage2::BMTFSetup::registerConsumes(), l1t::stage1::CaloSetup::registerConsumes(), MonopoleSteppingAction::registerConsumes(), l1t::stage2::GTSetup::registerConsumes(), MaterialBudgetHcal::registerConsumes(), PrintGeomSummary::registerConsumes(), SimG4HGCalValidation::registerConsumes(), HcalTestAnalysis::registerConsumes(), PrintGeomInfoAction::registerConsumes(), MaterialBudgetHcalProducer::registerConsumes(), SimG4HcalValidation::registerConsumes(), CaloSteppingAction::registerConsumes(), SCEnergyCorrectorSemiParm::RegParam::RegParam(), RegressionHelper::RegressionHelper(), lumi::CMSRunSummary2DB::retrieveData(), SiPixelDQMRocLevelAnalyzer::RocSumOneModule(), DDEcalEndcapTrap::rotate(), DDEcalEndcapTrapX::rotate(), RPCConeBuilder::RPCConeBuilder(), RPCInverseCPPFLinkMapESProducer::RPCInverseCPPFLinkMapESProducer(), RPCInverseLBLinkMapESProducer::RPCInverseLBLinkMapESProducer(), RPCInverseOMTFLinkMapESProducer::RPCInverseOMTFLinkMapESProducer(), RPCInverseTwinMuxLinkMapESProducer::RPCInverseTwinMuxLinkMapESProducer(), DTTracoChip::run(), SCEnergyCorrectorDRN::SCEnergyCorrectorDRN(), SCEnergyCorrectorSemiParm::SCEnergyCorrectorSemiParm(), sim::sensitiveDetectorMakers(), mkfit::TrackCand::setCombCandidate(), PhoMVACut::setConsumes(), PhoFull5x5SigmaIEtaIEtaValueMapCut::setConsumes(), PhoAnyPFIsoWithEAAndQuadScalingCut::setConsumes(), PhoAnyPFIsoWithEAAndExpoScalingCut::setConsumes(), PhoAnyPFIsoWithEACut::setConsumes(), GsfEleRelPFIsoScaledCut::setConsumes(), PhoAnyPFIsoWithEAAndExpoScalingEBCut::setConsumes(), GsfEleEffAreaPFIsoCut::setConsumes(), GsfEleMVACut::setConsumes(), GsfEleEmHadD1IsoRhoCut::setConsumes(), GsfEleValueMapIsoRhoCut::setConsumes(), GsfEleDzCut::setConsumes(), GsfEleDxyCut::setConsumes(), GsfEleCalPFClusterIsoCut::setConsumes(), MuonPOGStandardCut::setConsumes(), PhoGenericRhoPtScaledCut::setConsumes(), PhoGenericQuadraticRhoPtScaledCut::setConsumes(), GsfEleTrkPtIsoRhoCut::setConsumes(), GsfEleConversionVetoCut::setConsumes(), MuonDxyCut::setConsumes(), MuonDzCut::setConsumes(), GsfEleHadronicOverEMEnergyScaledCut::setConsumes(), EcalTPGDBCopy::setConsumes(), HcalDbProducer::ServiceTokenImpl< ProductType, RecordType, LABEL, EffectiveType >::setConsumes(), HcalDbProducer::TokensForServiceHolder< RecordType, TokenHolders >::setConsumes(), HcalDbProducer::TokenAndTopologyHolder< ProductT, RecordT >::setConsumes(), VersionedSelector< T >::setConsumes(), PFBlockAlgo::setImporters(), l1t::WriterProxyT< Record, Type >::setToken(), EgammaHLTPhase2ExtraProducer::Tokens::setToken(), EgammaHLTExtraProducer::Tokens::setToken(), SCEnergyCorrectorSemiParm::setTokens(), SCEnergyCorrectorDRN::setTokens(), SCEnergyCorrectorSemiParm::RegParam::setTokens(), PFECALSuperClusterAlgo::setTokens(), SiPhase2BadStripConfigurableFakeESSource::SiPhase2BadStripConfigurableFakeESSource(), SiPhase2OuterTrackerFakeLorentzAngleESSource::SiPhase2OuterTrackerFakeLorentzAngleESSource(), SiPhase2RecHitMatcherESProducer::SiPhase2RecHitMatcherESProducer(), SiPixel2DTemplateDBObjectESProducer::SiPixel2DTemplateDBObjectESProducer(), SiPixelGainCalibrationForHLTGPUESProducer::SiPixelGainCalibrationForHLTGPUESProducer(), SiPixelGenErrorDBObjectESProducer::SiPixelGenErrorDBObjectESProducer(), SiPixelROCsStatusAndMappingWrapperESProducer::SiPixelROCsStatusAndMappingWrapperESProducer(), SiPixelTemplateDBObjectESProducer::SiPixelTemplateDBObjectESProducer(), SiStripBackPlaneCorrectionFakeESSource::SiStripBackPlaneCorrectionFakeESSource(), SiStripClusterizerConditionsESProducer::SiStripClusterizerConditionsESProducer(), SiStripClusterizerConditionsGPUESProducer::SiStripClusterizerConditionsGPUESProducer(), SiStripDelayESProducer::SiStripDelayESProducer(), SiStripLorentzAngleFakeESSource::SiStripLorentzAngleFakeESSource(), SiStripRegionConnectivity::SiStripRegionConnectivity(), SiTrackerMultiRecHitUpdatorESProducer::SiTrackerMultiRecHitUpdatorESProducer(), SkippingLayerCosmicNavigationSchoolESProducer::SkippingLayerCosmicNavigationSchoolESProducer(), SmartPropagatorESProducer::SmartPropagatorESProducer(), p2eg::stitchClusterOverRegionBoundary(), LRHelpFunctions::storeControlPlots(), storeTracks(), StripCompletePlot(), StripCPEESProducer::StripCPEESProducer(), heppy::mt2w_bisect::mt2w::teco(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerA::TestAlpakaESProducerA(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerB::TestAlpakaESProducerB(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerC::TestAlpakaESProducerC(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerD::TestAlpakaESProducerD(), edmtest::ThinningThingSelector::ThinningThingSelector(), TkDetMapESProducer::TkDetMapESProducer(), TkMSParameterizationBuilder::TkMSParameterizationBuilder(), SiStripGainSimESProducer::TokenLabel::TokenLabel(), CandidateBoostedDoubleSecondaryVertexComputer::Tokens::Tokens(), CombinedMVAV2JetTagComputer::Tokens::Tokens(), CandidateChargeBTagComputer::Tokens::Tokens(), MuonTagger::Tokens::Tokens(), ElectronTagger::Tokens::Tokens(), CharmTagger::Tokens::Tokens(), HeavyIonCSVTagger::Tokens::Tokens(), EgammaHLTPhase2ExtraProducer::Tokens::Tokens(), EgammaHLTExtraProducer::Tokens::Tokens(), SiStripDelayESProducer::TokenSign::TokenSign(), TotemGeometryESModule::TotemGeometryESModule(), TPNCor::TPNCor(), TrackerDigiGeometryESModule::TrackerDigiGeometryESModule(), TrackerGeometricDetESModule::TrackerGeometricDetESModule(), TrackerInteractionGeometryESProducer::TrackerInteractionGeometryESProducer(), TrackerParametersESModule::TrackerParametersESModule(), pflow::importers::TrackFromParentImporter< Collection, Adaptor >::TrackFromParentImporter(), TransientTrackBuilderESProducer::TransientTrackBuilderESProducer(), TSCBLBuilderWithPropagatorESProducer::TSCBLBuilderWithPropagatorESProducer(), TwoTrackMinimumDistanceHelixLine::updateCoeffs(), magneticfield::VolumeBasedMagneticFieldESProducerFromDB::VolumeBasedMagneticFieldESProducerFromDB(), BPHWriteSpecificDecay::write(), WriteESAlignments::WriteESAlignments(), and XMLIdealMagneticFieldGeometryESProducer::XMLIdealMagneticFieldGeometryESProducer().

◆ cellNeighbors

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > * gpuPixelDoublets::cellNeighbors

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by for(), and if().

◆ cellNeighborsContainer

int CellNeighborsVector<TrackerTraits> CellNeighbors<TrackerTraits>* gpuPixelDoublets::cellNeighborsContainer

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by if().

◆ cells

__device__ uint32_t GPUCACellT< TrackerTraits > * gpuPixelDoublets::cells

Definition at line 34 of file gpuFishbone.h.

Referenced by for().

◆ cellTracks

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > * gpuPixelDoublets::cellTracks

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by for(), and if().

◆ cellTracksContainer

int CellNeighborsVector<TrackerTraits> CellNeighbors<TrackerTraits> CellTracksVector<TrackerTraits> CellTracks<TrackerTraits>* gpuPixelDoublets::cellTracksContainer
Initial value:
{
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const isOuterHitOfCell
assert(be >=bs)

Definition at line 16 of file gpuPixelDoublets.h.

Referenced by if().

◆ checkTrack

GPUCACellT<TrackerTraits> uint32_t const* __restrict__ OuterHitOfCell<TrackerTraits> const int32_t bool gpuPixelDoublets::checkTrack
Initial value:
{
constexpr auto maxCellsPerHit = GPUCACellT<TrackerTraits>::maxCellsPerHit

Definition at line 38 of file gpuFishbone.h.

Referenced by for().

◆ cuts

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > CellCutsT< TrackerTraits > const & gpuPixelDoublets::cuts
Initial value:
{
doubletsFromHisto<TrackerTraits>(
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const isOuterHitOfCell
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int maxNumOfDoublets
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int CellCutsT< TrackerTraits > *const cuts
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > * cellTracks
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int nActualPairs
int CellNeighborsVector< TrackerTraits > * cellNeighbors
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > hh
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ nCells

Definition at line 50 of file gpuPixelDoublets.h.

Referenced by for().

◆ d

uint16_t gpuPixelDoublets::d[maxCellsPerHit]

Definition at line 50 of file gpuFishbone.h.

Referenced by for().

◆ doPtCut

const bool gpuPixelDoublets::doPtCut = cuts.doPtCut_

Definition at line 136 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ doZ0Cut

const bool gpuPixelDoublets::doZ0Cut = cuts.doZ0Cut_

Definition at line 135 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ first

auto gpuPixelDoublets::first = blockIdx.x * blockDim.x + threadIdx.x

Definition at line 18 of file gpuPixelDoublets.h.

Referenced by for().

◆ firstX

auto gpuPixelDoublets::firstX = threadIdx.x

◆ firstY

auto gpuPixelDoublets::firstY = threadIdx.y + blockIdx.y * blockDim.y

Definition at line 45 of file gpuFishbone.h.

Referenced by NoisyChannel::getAverage2D(), and PixelTemplateSmearerBase::smearHit().

◆ getDoubletsFromHistoMaxBlockSize

constexpr auto gpuPixelDoublets::getDoubletsFromHistoMaxBlockSize = 64

Definition at line 34 of file gpuPixelDoublets.h.

◆ getDoubletsFromHistoMinBlocksPerMP

constexpr auto gpuPixelDoublets::getDoubletsFromHistoMinBlocksPerMP = 16

Definition at line 35 of file gpuPixelDoublets.h.

◆ hardPtCut

const float gpuPixelDoublets::hardPtCut = cuts.ptCut_

◆ hh

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > gpuPixelDoublets::hh

◆ idy

auto gpuPixelDoublets::idy = blockIdx.y * blockDim.y + threadIdx.y

Definition at line 168 of file gpuPixelDoubletsAlgos.h.

Referenced by for(), and TTUTrackingAlg::runSeedBuster().

◆ innerLayerCumulativeSize

__shared__ uint32_t gpuPixelDoublets::innerLayerCumulativeSize[TrackerTraits::nPairs]

Definition at line 156 of file gpuPixelDoubletsAlgos.h.

Referenced by for(), and if().

◆ isOuterHitOfCell

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > gpuPixelDoublets::isOuterHitOfCell = isOuterHitOfCellWrap.container

Definition at line 41 of file gpuFishbone.h.

Referenced by for().

◆ isOuterHitOfCellWrap

GPUCACellT<TrackerTraits> uint32_t const* __restrict__ OuterHitOfCell<TrackerTraits> const gpuPixelDoublets::isOuterHitOfCellWrap

Definition at line 34 of file gpuFishbone.h.

◆ l

uint8_t gpuPixelDoublets::l[maxCellsPerHit]

Definition at line 51 of file gpuFishbone.h.

Referenced by for().

◆ layerSize

auto gpuPixelDoublets::layerSize = [=](uint8_t li) { return offsets[li + 1] - offsets[li]; }

Definition at line 150 of file gpuPixelDoubletsAlgos.h.

Referenced by if().

◆ maxNumOfDoublets

__device__ uint32_t gpuPixelDoublets::maxNumOfDoublets

Definition at line 43 of file gpuPixelDoublets.h.

Referenced by for().

◆ minRadius

const float gpuPixelDoublets::minRadius = hardPtCut * 87.78f

Definition at line 141 of file gpuPixelDoubletsAlgos.h.

◆ minRadius2T4

const float gpuPixelDoublets::minRadius2T4 = 4.f * minRadius * minRadius

Definition at line 142 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ n

float gpuPixelDoublets::n[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

◆ nActualPairs

uint32_t CellNeighborsVector<TrackerTraits> CellTracksVector<TrackerTraits> HitsConstView<TrackerTraits> OuterHitOfCell<TrackerTraits> int gpuPixelDoublets::nActualPairs

◆ nCells

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t * gpuPixelDoublets::nCells

Definition at line 34 of file gpuFishbone.h.

Referenced by for().

◆ nHits

int gpuPixelDoublets::nHits

Definition at line 34 of file gpuFishbone.h.

◆ ntot

__shared__ uint32_t gpuPixelDoublets::ntot

◆ offset

int32_t gpuPixelDoublets::offset = isOuterHitOfCellWrap.offset

Definition at line 42 of file gpuFishbone.h.

◆ offsets

uint32_t const* __restrict__ gpuPixelDoublets::offsets = hh.hitsLayerStart().data()

Definition at line 147 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ pairLayerId

uint32_t gpuPixelDoublets::pairLayerId = 0

Definition at line 172 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ phiBinner

auto const& __restrict__ gpuPixelDoublets::phiBinner = hh.phiBinner()

Definition at line 146 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ stride

auto gpuPixelDoublets::stride = blockDim.x

Definition at line 170 of file gpuPixelDoubletsAlgos.h.

Referenced by caHitNtupletGeneratorKernels::__attribute__(), HcalDigisProducerGPU::acquire(), cms::rocmtest::add_vectors_d(), cms::rocmtest::add_vectors_f(), PhysicsPerformanceDBWriterFromFile_WPandPayload::beginJob(), PhysicsPerformanceDBWriterFromFile_WPandPayload_IOV::beginJob(), cms::soa::SoAParametersImpl< SoAColumnType::eigen, T >::checkAlignment(), calo::multifit::compute_decomposition_unrolled(), npstat::ArrayND< Numeric >::contract(), npstat::ArrayND< Numeric >::contractLoop(), npstat::ArrayND< Numeric >::convertToLastDimCdfLoop(), for(), npstat::ArrayND< Numeric >::functorFillLoop(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_exotic(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_multiSeedIterative(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_multiSeedParallel(), cms::alpakatools::blocks_with_stride< TAcc, typename >::iterator::iterator(), cms::alpakatools::independent_groups< TAcc, typename >::iterator::iterator(), npstat::ArrayND< Numeric >::jointSliceLoop(), npstat::ArrayND< Numeric >::linearFillLoop(), npstat::ArrayND< Numeric >::makeUnit(), npstat::ArrayND< Numeric >::processSubrangeLoop(), npstat::ArrayND< Numeric >::projectInnerLoop(), npstat::ArrayND< Numeric >::projectInnerLoop2(), npstat::ArrayND< Numeric >::projectLoop(), npstat::ArrayND< Numeric >::projectLoop2(), npstat::ArrayND< Numeric >::scaleBySliceInnerLoop(), npstat::ArrayND< Numeric >::scaleBySliceLoop(), l1tVertexFinder::VertexFinder::strided_iota(), npstat::ArrayND< Numeric >::sumBelowLoop(), and ALPAKA_ACCELERATOR_NAMESPACE::testESAlgoAsync().

◆ x

float gpuPixelDoublets::x[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

◆ y

float gpuPixelDoublets::y[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

◆ z

float gpuPixelDoublets::z[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

◆ z0cut

const float gpuPixelDoublets::z0cut = cuts.z0Cut_

Definition at line 138 of file gpuPixelDoubletsAlgos.h.

Referenced by for(), and trklet::TrackletLUT::getVMRLookup().