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 CellCutsT< TrackerTraits > 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
 
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]; }
 
const uint32_t maxNumOfDoublets = cuts.maxNumberOfDoublets_
 
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]
 

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 115 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 145 of file gpuPixelDoubletsAlgos.h.

References cms::cudacompat::__ldg(), funct::abs(), assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::atomicSub(), newFWLiteAna::bin, cellNeighbors, cells, cellTracks, cuts, HLT_2022v15_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, phase1PixelTopology::minz, LaserClient_cfi::nbins, nCells, caHitNtupletGeneratorKernels::nmin, phase1PixelTopology::nPairs, offsets, SurfaceOrientation::outer, AlCaHLTBitMon_ParallelJobs::p, pairLayerId, phiBinner, phase1PixelTopology::phicuts, submitPVValidationJobs::ptcut, short2phi(), stride, compareTotals::tot, and HLTMuonOfflineAnalyzer_cff::z0Cut.

145  {
147  ;
148  --pairLayerId; // move to lower_bound ??
149 
153 
155  uint8_t outer = TrackerTraits::layerPairs[2 * pairLayerId + 1];
156  assert(outer > inner);
157 
158  auto hoff = PhiBinner::histOff(outer);
159  auto i = (0 == pairLayerId) ? j : j - innerLayerCumulativeSize[pairLayerId - 1];
160  i += offsets[inner];
161 
162  assert(i >= offsets[inner]);
163  assert(i < offsets[inner + 1]);
164 
165  // found hit corresponding to our cuda thread, now do the job
166 
167  if (hh[i].detectorIndex() > gpuClustering::maxNumModules)
168  continue; // invalid
169 
170  /* maybe clever, not effective when zoCut is on
171  auto bpos = (mi%8)/4; // if barrel is 1 for z>0
172  auto fpos = (outer>3) & (outer<7);
173  if ( ((inner<3) & (outer>3)) && bpos!=fpos) continue;
174  */
175 
176  auto mez = hh[i].zGlobal();
177 
179  continue;
180 
181  if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts.clusterCut(hh, i))
182  continue;
183 
184  auto mep = hh[i].iphi();
185  auto mer = hh[i].rGlobal();
186 
187  // all cuts: true if fails
188  constexpr float z0cut = TrackerTraits::z0Cut; // cm
189  constexpr float hardPtCut = TrackerTraits::doubletHardPt; // GeV
190  // cm (1 GeV track has 1 GeV/c / (e * 3.8T) ~ 87 cm radius in a 3.8T field)
191  constexpr float minRadius = hardPtCut * 87.78f;
192  constexpr float minRadius2T4 = 4.f * minRadius * minRadius;
193  auto ptcut = [&](int j, int16_t idphi) {
194  auto r2t4 = minRadius2T4;
195  auto ri = mer;
196  auto ro = hh[j].rGlobal();
197  auto dphi = short2phi(idphi);
198  return dphi * dphi * (r2t4 - ri * ro) > (ro - ri) * (ro - ri);
199  };
200  auto z0cutoff = [&](int j) {
201  auto zo = hh[j].zGlobal();
202  auto ro = hh[j].rGlobal();
203  auto dr = ro - mer;
204  return dr > TrackerTraits::maxr[pairLayerId] || dr < 0 || std::abs((mez * ro - mer * zo)) > z0cut * dr;
205  };
206 
207  auto iphicut = TrackerTraits::phicuts[pairLayerId];
208 
209  auto kl = PhiBinner::bin(int16_t(mep - iphicut));
210  auto kh = PhiBinner::bin(int16_t(mep + iphicut));
211  auto incr = [](auto& k) { return k = (k + 1) % PhiBinner::nbins(); };
212 
213 #ifdef GPU_DEBUG
214  int tot = 0;
215  int nmin = 0;
216  int tooMany = 0;
217 #endif
218 
219  auto khh = kh;
220  incr(khh);
221  for (auto kk = kl; kk != khh; incr(kk)) {
222 #ifdef GPU_DEBUG
223  if (kk != kl && kk != kh)
224  nmin += phiBinner.size(kk + hoff);
225 #endif
226  auto const* __restrict__ p = phiBinner.begin(kk + hoff);
227  auto const* __restrict__ e = phiBinner.end(kk + hoff);
228  p += first;
229  for (; p < e; p += stride) {
230  auto oi = __ldg(p);
231  assert(oi >= offsets[outer]);
232  assert(oi < offsets[outer + 1]);
233  auto mo = hh[oi].detectorIndex();
234 
236  continue; // invalid
237 
238  if (doZ0Cut && z0cutoff(oi))
239  continue;
240  auto mop = hh[oi].iphi();
241  uint16_t idphi = std::min(std::abs(int16_t(mop - mep)), std::abs(int16_t(mep - mop)));
242  if (idphi > iphicut)
243  continue;
244 
245  if (doClusterCut && cuts.zSizeCut(hh, i, oi))
246  continue;
247  if (doPtCut && ptcut(oi, idphi))
248  continue;
249 
250  auto ind = atomicAdd(nCells, 1);
251  if (ind >= maxNumOfDoublets) {
252  atomicSub(nCells, 1);
253  break;
254  } // move to SimpleVector??
255  // int layerPairId, int doubletId, int innerHitId, int outerHitId)
256  cells[ind].init(*cellNeighbors, *cellTracks, hh, pairLayerId, i, oi);
257  isOuterHitOfCell[oi].push_back(ind);
258 #ifdef GPU_DEBUG
259  if (isOuterHitOfCell[oi].full())
260  ++tooMany;
261  ++tot;
262 #endif
263  }
264  }
265 // #endif
266 #ifdef GPU_DEBUG
267  if (tooMany > 0)
268  printf("OuterHitOfCell full for %d in layer %d/%d, %d,%d %d, %d %.3f %.3f\n",
269  i,
270  inner,
271  outer,
272  nmin,
273  tot,
274  tooMany,
275  iphicut,
278 #endif
279  } // 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
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
constexpr int16_t phicuts[nPairs]
uint32_t const *__restrict__ offsets
const uint32_t maxNumOfDoublets
__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, findQualityFiles::size, 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
uint8_t l[maxCellsPerHit]
Definition: gpuFishbone.h:51
double f[11][100]
GPUCACellT< TrackerTraits > * cells
Definition: gpuFishbone.h:34
auto const isOuterHitOfCell
Definition: gpuFishbone.h:41
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
assert(be >=bs)
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > * cellNeighborsContainer
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > * cellTracks
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(), CSCConditions::CSCConditions(), CSCGeometryESModule::CSCGeometryESModule(), 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(), 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(), 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(), PFRecHitProducer::PFRecHitProducer(), PixelCPEFastESProducerT< TrackerTraits >::PixelCPEFastESProducerT(), PlotDebugFPIX_XYMap(), PlotPixelMultVtxPos(), FastTimerService::postESModule(), FastTimerService::preESModule(), DDEcalEndcapTrapX::print(), DDEcalEndcapTrap::print(), printPlot(), PixelTracksProducer::produce(), CastorClusterProducer::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(), MultiHitGeneratorFromChi2::refit2Hits(), HITrackingRegionProducer::regions(), HITrackingRegionForPrimaryVtxProducer::regions(), l1t::stage2::CaloLayer1Setup::registerConsumes(), l1t::stage2::EMTFSetup::registerConsumes(), l1t::stage2::CaloSetup::registerConsumes(), l1t::stage2::GMTSetup::registerConsumes(), l1t::stage2::GTSetup::registerConsumes(), l1t::stage2::BMTFSetup::registerConsumes(), l1t::stage1::CaloSetup::registerConsumes(), MonopoleSteppingAction::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(), DDEcalEndcapTrapX::rotate(), DDEcalEndcapTrap::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(), PhoAnyPFIsoWithEACut::setConsumes(), PhoAnyPFIsoWithEAAndQuadScalingCut::setConsumes(), GsfEleRelPFIsoScaledCut::setConsumes(), GsfEleEffAreaPFIsoCut::setConsumes(), GsfEleMVACut::setConsumes(), PhoAnyPFIsoWithEAAndExpoScalingCut::setConsumes(), PhoAnyPFIsoWithEAAndExpoScalingEBCut::setConsumes(), GsfEleEmHadD1IsoRhoCut::setConsumes(), GsfEleValueMapIsoRhoCut::setConsumes(), GsfEleDxyCut::setConsumes(), GsfEleDzCut::setConsumes(), GsfEleCalPFClusterIsoCut::setConsumes(), GsfEleTrkPtIsoRhoCut::setConsumes(), PhoGenericRhoPtScaledCut::setConsumes(), GsfEleConversionVetoCut::setConsumes(), MuonPOGStandardCut::setConsumes(), PhoGenericQuadraticRhoPtScaledCut::setConsumes(), MuonDzCut::setConsumes(), MuonDxyCut::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(), SiStripDelayESProducer::SiStripDelayESProducer(), SiStripLorentzAngleFakeESSource::SiStripLorentzAngleFakeESSource(), SiStripRegionConnectivity::SiStripRegionConnectivity(), SiTrackerMultiRecHitUpdatorESProducer::SiTrackerMultiRecHitUpdatorESProducer(), SkippingLayerCosmicNavigationSchoolESProducer::SkippingLayerCosmicNavigationSchoolESProducer(), SmartPropagatorESProducer::SmartPropagatorESProducer(), 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(), MuonTagger::Tokens::Tokens(), CandidateChargeBTagComputer::Tokens::Tokens(), ElectronTagger::Tokens::Tokens(), CharmTagger::Tokens::Tokens(), HeavyIonCSVTagger::Tokens::Tokens(), SiPixelQualityESProducer::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__ 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__ GPUCACellT< TrackerTraits > * gpuPixelDoublets::cells

Definition at line 34 of file gpuFishbone.h.

Referenced by for().

◆ cellTracks

__device__ 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__ 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
int CellNeighborsVector< TrackerTraits > * cellNeighbors
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > hh
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > * cellTracks
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int nActualPairs
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ nCells
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int CellCutsT< TrackerTraits > cuts

Definition at line 49 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 112 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ doZ0Cut

const bool gpuPixelDoublets::doZ0Cut = cuts.doZ0Cut_

Definition at line 111 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.

◆ hh

__device__ 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 139 of file gpuPixelDoubletsAlgos.h.

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

◆ innerLayerCumulativeSize

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

Definition at line 127 of file gpuPixelDoubletsAlgos.h.

Referenced by for(), and if().

◆ isOuterHitOfCell

__device__ 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 121 of file gpuPixelDoubletsAlgos.h.

Referenced by if().

◆ maxNumOfDoublets

const uint32_t gpuPixelDoublets::maxNumOfDoublets = cuts.maxNumberOfDoublets_

Definition at line 113 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__ 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 118 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ pairLayerId

uint32_t gpuPixelDoublets::pairLayerId = 0

Definition at line 143 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ phiBinner

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

Definition at line 117 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

◆ stride

auto gpuPixelDoublets::stride = blockDim.x

◆ 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().