CMS 3D CMS Logo

Classes | Typedefs | Functions | Variables
gpuVertexFinder Namespace Reference

Classes

class  Producer
 
struct  WorkSpace
 

Typedefs

using Hist = cms::cuda::HistoContainer< uint8_t, 256, 16000, 8, uint16_t >
 
using TkSoA = pixelTrack::TrackSoA
 
using ZVertices = ZVertexSoA
 

Functions

__device__ __attribute__ ((always_inline)) void clusterTracksByDensity(gpuVertexFinder
 
 __syncthreads ()
 
 assert (soa)
 
 assert (nvFinal<=nvIntermediate)
 
 assert (pdata)
 
 assert (zt)
 
 assert ((int) nt<=hist.capacity())
 
 assert (hist.size()==nt)
 
 assert (foundClusters< ZVertices::MAXVTX)
 
size d for d tracks hist hist capacity ()
 
hist finalize (hws)
 
 for (int idx=first, nt=tracks.nTracks();idx< nt;idx+=gridDim.x *blockDim.x)
 
 for (auto j=threadIdx.x;j< Hist::totbins();j+=blockDim.x) = float(nn[i]) / chi2[i]
 
 if (verbose &&0==threadIdx.x) printf("params %d %f %f %f\
 
 if (1==nvFinal)
 
 if (threadIdx.x< 32) hws[threadIdx.x]=0
 
pws init ()
 
size d for d tracks hist nbins ()
 
 while (__syncthreads_or(more))
 

Variables

float *__restrict__ chi2 = data.chi2
 
__device__ WorkSpace float chi2Max
 
WorkSpace int float float float chi2max
 
auto &__restrict__ data = *pdata
 
WorkSpace int float eps
 
auto er2mx = errmax * errmax
 
WorkSpace int float float errmax
 
float const *__restrict__ ezt2 = ws.ezt2
 
auto first = blockIdx.x * blockDim.x + threadIdx.x
 
auto const & fit = tracks.stateAtBS
 
__shared__ unsigned int foundClusters = 0
 
__shared__ Hist hist
 
__shared__ Hist::Counter hws [32]
 
int32_t *__restrict__ iv = ws.iv
 
uint8_t *__restrict__ izt = ws.izt
 
__device__ WorkSpace float maxChi2
 
constexpr float maxChi2ForFinalFit = 5000.f
 
constexpr float maxChi2ForFirstFit = 50.f
 
constexpr float maxChi2ForSplit = 9.f
 
WorkSpace int minT
 
bool more = true
 
size d for d tracks n
 
__shared__ int nloops = 0
 
int32_t *__restrict__ nn = data.ndof
 
__shared__ int noise
 
auto nt = ws.ntrks
 
uint32_t & nvFinal = data.nvFinal
 
uint32_t & nvIntermediate = ws.nvIntermediate
 
ZVertexSoA WorkSpace float float ptMax
 
ZVertexSoA WorkSpace float ptMin
 
float const *__restrict__ ptt2 = ws.ptt2
 
float *__restrict__ ptv2 = data.ptv2
 
WorkSpacepws
 
auto const * quality = tracks.qualityData()
 
ZVertexSoAsoa
 
uint16_t *__restrict__ sortInd = data.sortInd
 
auto const & tracks = *ptracks
 
auto &__restrict__ ws = *pws
 
float *__restrict__ wv = data.wv
 
float const *__restrict__ zt = ws.zt
 
float *__restrict__ zv = data.zv
 

Typedef Documentation

◆ Hist

typedef cms::cuda::HistoContainer< uint8_t, 256, 16000, 8, uint16_t > gpuVertexFinder::Hist

Definition at line 47 of file gpuClusterTracksDBSCAN.h.

◆ TkSoA

Definition at line 13 of file gpuVertexFinder.h.

◆ ZVertices

Definition at line 12 of file gpuVertexFinder.h.

Function Documentation

◆ __attribute__()

gpuVertexFinder::__attribute__ ( (always_inline)  )
inline

Definition at line 20 of file gpuClusterTracksByDensity.h.

References __syncthreads(), funct::abs(), assert(), cms::cudacompat::atomicInc(), cms::cudacompat::blockDim, cms::cuda::OneToManyAssoc::capacity(), chi2max, data, eps, er2mx, errmax, ezt2, foundClusters, hist, hws, mps_fire::i, createfilelist::int, iv, izt, dqmiolumiharvest::j, heppy_loop::loop, visualization-live-secondInstance_cfg::m, SiStripPI::max, ZVertexSoA::MAXTRACKS, ZVertexSoA::MAXVTX, SiStripPI::min, minT, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins(), nn, nt, nvFinal, nvIntermediate, cms::cuda::OneToManyAssoc::off, pws, cms::cudacompat::threadIdx, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::totbins(), verbose, ws, and zt.

26  {
27  using namespace gpuVertexFinder;
28  constexpr bool verbose = false; // in principle the compiler should optmize out if false
29 
30  if (verbose && 0 == threadIdx.x)
31  printf("params %d %f %f %f\n", minT, eps, errmax, chi2max);
32 
33  auto er2mx = errmax * errmax;
34 
35  auto& __restrict__ data = *pdata;
36  auto& __restrict__ ws = *pws;
37  auto nt = ws.ntrks;
38  float const* __restrict__ zt = ws.zt;
39  float const* __restrict__ ezt2 = ws.ezt2;
40 
41  uint32_t& nvFinal = data.nvFinal;
42  uint32_t& nvIntermediate = ws.nvIntermediate;
43 
44  uint8_t* __restrict__ izt = ws.izt;
45  int32_t* __restrict__ nn = data.ndof;
46  int32_t* __restrict__ iv = ws.iv;
47 
48  assert(pdata);
49  assert(zt);
50 
52  __shared__ Hist hist;
53  __shared__ typename Hist::Counter hws[32];
54  for (auto j = threadIdx.x; j < Hist::totbins(); j += blockDim.x) {
55  hist.off[j] = 0;
56  }
57  __syncthreads();
58 
59  if (verbose && 0 == threadIdx.x)
60  printf("booked hist with %d bins, size %d for %d tracks\n", hist.nbins(), hist.capacity(), nt);
61 
62  assert((int)nt <= hist.capacity());
63 
64  // fill hist (bin shall be wider than "eps")
65  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
67  int iz = int(zt[i] * 10.); // valid if eps<=0.1
68  // iz = std::clamp(iz, INT8_MIN, INT8_MAX); // sorry c++17 only
69  iz = std::min(std::max(iz, INT8_MIN), INT8_MAX);
70  izt[i] = iz - INT8_MIN;
71  assert(iz - INT8_MIN >= 0);
72  assert(iz - INT8_MIN < 256);
73  hist.count(izt[i]);
74  iv[i] = i;
75  nn[i] = 0;
76  }
77  __syncthreads();
78  if (threadIdx.x < 32)
79  hws[threadIdx.x] = 0; // used by prefix scan...
80  __syncthreads();
81  hist.finalize(hws);
82  __syncthreads();
83  assert(hist.size() == nt);
84  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
85  hist.fill(izt[i], uint16_t(i));
86  }
87  __syncthreads();
88 
89  // count neighbours
90  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
91  if (ezt2[i] > er2mx)
92  continue;
93  auto loop = [&](uint32_t j) {
94  if (i == j)
95  return;
96  auto dist = std::abs(zt[i] - zt[j]);
97  if (dist > eps)
98  return;
99  if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
100  return;
101  nn[i]++;
102  };
103 
104  cms::cuda::forEachInBins(hist, izt[i], 1, loop);
105  }
106 
107  __syncthreads();
108 
109  // find closest above me .... (we ignore the possibility of two j at same distance from i)
110  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
111  float mdist = eps;
112  auto loop = [&](uint32_t j) {
113  if (nn[j] < nn[i])
114  return;
115  if (nn[j] == nn[i] && zt[j] >= zt[i])
116  return; // if equal use natural order...
117  auto dist = std::abs(zt[i] - zt[j]);
118  if (dist > mdist)
119  return;
120  if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
121  return; // (break natural order???)
122  mdist = dist;
123  iv[i] = j; // assign to cluster (better be unique??)
124  };
125  cms::cuda::forEachInBins(hist, izt[i], 1, loop);
126  }
127 
128  __syncthreads();
129 
130 #ifdef GPU_DEBUG
131  // mini verification
132  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
133  if (iv[i] != int(i))
134  assert(iv[iv[i]] != int(i));
135  }
136  __syncthreads();
137 #endif
138 
139  // consolidate graph (percolate index of seed)
140  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
141  auto m = iv[i];
142  while (m != iv[m])
143  m = iv[m];
144  iv[i] = m;
145  }
146 
147 #ifdef GPU_DEBUG
148  __syncthreads();
149  // mini verification
150  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
151  if (iv[i] != int(i))
152  assert(iv[iv[i]] != int(i));
153  }
154 #endif
155 
156 #ifdef GPU_DEBUG
157  // and verify that we did not spit any cluster...
158  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
159  auto minJ = i;
160  auto mdist = eps;
161  auto loop = [&](uint32_t j) {
162  if (nn[j] < nn[i])
163  return;
164  if (nn[j] == nn[i] && zt[j] >= zt[i])
165  return; // if equal use natural order...
166  auto dist = std::abs(zt[i] - zt[j]);
167  if (dist > mdist)
168  return;
169  if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
170  return;
171  mdist = dist;
172  minJ = j;
173  };
174  cms::cuda::forEachInBins(hist, izt[i], 1, loop);
175  // should belong to the same cluster...
176  assert(iv[i] == iv[minJ]);
177  assert(nn[i] <= nn[iv[i]]);
178  }
179  __syncthreads();
180 #endif
181 
182  __shared__ unsigned int foundClusters;
183  foundClusters = 0;
184  __syncthreads();
185 
186  // find the number of different clusters, identified by a tracks with clus[i] == i and density larger than threshold;
187  // mark these tracks with a negative id.
188  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
189  if (iv[i] == int(i)) {
190  if (nn[i] >= minT) {
191  auto old = atomicInc(&foundClusters, 0xffffffff);
192  iv[i] = -(old + 1);
193  } else { // noise
194  iv[i] = -9998;
195  }
196  }
197  }
198  __syncthreads();
199 
201 
202  // propagate the negative id to all the tracks in the cluster.
203  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
204  if (iv[i] >= 0) {
205  // mark each track in a cluster with the same id as the first one
206  iv[i] = iv[iv[i]];
207  }
208  }
209  __syncthreads();
210 
211  // adjust the cluster id to be a positive value starting from 0
212  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
213  iv[i] = -iv[i] - 1;
214  }
215 
217 
218  if (verbose && 0 == threadIdx.x)
219  printf("found %d proto vertices\n", foundClusters);
220  }
const dim3 threadIdx
Definition: cudaCompat.h:29
int32_t *__restrict__ iv
WorkSpace int float eps
bool verbose
auto &__restrict__ data
const dim3 blockDim
Definition: cudaCompat.h:30
static constexpr uint32_t MAXVTX
Definition: ZVertexSoA.h:12
float const *__restrict__ zt
auto &__restrict__ ws
WorkSpace int float float float chi2max
T1 atomicInc(T1 *a, T2 b)
Definition: cudaCompat.h:48
Abs< T >::type abs(const T &t)
Definition: Abs.h:22
typename Base::Counter Counter
float const *__restrict__ ezt2
__shared__ Hist::Counter hws[32]
uint8_t *__restrict__ izt
WorkSpace int float float errmax
__shared__ Hist hist
static constexpr uint32_t nbins()
static constexpr uint32_t totbins()
int32_t *__restrict__ nn
static constexpr uint32_t MAXTRACKS
Definition: ZVertexSoA.h:11
constexpr auto capacity() const
__shared__ unsigned int foundClusters

◆ __syncthreads()

gpuVertexFinder::__syncthreads ( )
inline

◆ assert() [1/7]

gpuVertexFinder::assert ( soa  )

◆ assert() [2/7]

gpuVertexFinder::assert ( nvFinal<=  nvIntermediate)

◆ assert() [3/7]

gpuVertexFinder::assert ( pdata  )

◆ assert() [4/7]

gpuVertexFinder::assert ( zt  )

◆ assert() [5/7]

gpuVertexFinder::assert ( (int) nt<=hist.capacity()  )

◆ assert() [6/7]

gpuVertexFinder::assert ( hist.  size() = =nt)

◆ assert() [7/7]

gpuVertexFinder::assert ( )

◆ capacity()

size d for d tracks hist hist gpuVertexFinder::capacity ( )

◆ finalize()

hist gpuVertexFinder::finalize ( hws  )

◆ for() [1/2]

gpuVertexFinder::for ( int  idx = first)

Definition at line 29 of file gpuVertexFinder.cc.

References assert(), cms::cudacompat::atomicAdd(), data, pixelTrack::highPurity, ZVertexSoA::idv, heavyIonCSV_trainingSettings::idx, SiStripPI::min, nHits, DiDispStaMuonMonitor_cfi::pt, ptMax, ptMin, pws, quality, soa, and tracks.

29  {
30  auto nHits = tracks.nHits(idx);
31  assert(nHits >= 3);
32 
33  // initialize soa...
34  soa->idv[idx] = -1;
35 
36  if (tracks.isTriplet(idx))
37  continue; // no triplets
39  continue;
40 
41  auto pt = tracks.pt(idx);
42 
43  if (pt < ptMin)
44  continue;
45 
46  // clamp pt
47  pt = std::min(pt, ptMax);
48 
49  auto& data = *pws;
50  auto it = atomicAdd(&data.ntrks, 1);
51  data.itrk[it] = idx;
52  data.zt[it] = tracks.zip(idx);
53  data.ezt2[it] = fit.covariance(idx)(14);
54  data.ptt2[it] = pt * pt;
55  }
constexpr float ptMin
auto const * quality
ZVertexSoA * soa
auto const & tracks
cannot be loose
caConstants::TupleMultiplicity const CAHitNtupletGeneratorKernelsGPU::HitToTuple const cms::cuda::AtomicPairCounter GPUCACell const *__restrict__ uint32_t const *__restrict__ gpuPixelDoublets::CellNeighborsVector const gpuPixelDoublets::CellTracksVector const GPUCACell::OuterHitOfCell const int32_t nHits
char data[epos_bytes_allocation]
Definition: EPOS_Wrapper.h:79
int16_t idv[MAXTRACKS]
Definition: ZVertexSoA.h:14
ZVertexSoA WorkSpace float float ptMax
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61

◆ for() [2/2]

gpuVertexFinder::for ( ) = float(nn[i]) / chi2[i]

◆ if() [1/3]

gpuVertexFinder::if ( verbose &&  0 = = threadIdx.x)

◆ if() [2/3]

gpuVertexFinder::if ( = nvFinal)

Definition at line 53 of file gpuSortByPt2.h.

References sortInd, and cms::cudacompat::threadIdx.

53  {
54  if (threadIdx.x == 0)
55  sortInd[0] = 0;
56  return;
57  }
const dim3 threadIdx
Definition: cudaCompat.h:29
uint16_t *__restrict__ sortInd
Definition: gpuSortByPt2.h:27

◆ if() [3/3]

gpuVertexFinder::if ( )
pure virtual

◆ init()

pws gpuVertexFinder::init ( )

◆ nbins()

size d for d tracks hist gpuVertexFinder::nbins ( )

◆ while()

gpuVertexFinder::while ( __syncthreads_or(more )

Definition at line 109 of file gpuClusterTracksIterative.h.

References funct::abs(), assert(), cms::cudacompat::atomicMin(), cms::cuda::be, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::bin(), cms::cudacompat::blockDim, chi2max, eps, ezt2, hist, mps_fire::i, iv, izt, dqmiolumiharvest::j, dqmdumpme::k, heppy_loop::loop, visualization-live-secondInstance_cfg::m, SiStripPI::min, minT, more, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins(), nloops, nn, nt, AlCaHLTBitMon_ParallelJobs::p, cms::cudacompat::threadIdx, and zt.

Referenced by ThreeThresholdAlgorithm::addToCandidate(), EcalSimParametersFromDD::build(), cppFunctionSkipper::delLines(), cppFunctionSkipper::filterFile(), commentSkipper.buildFileCommentSkipper::filterMultilineComment(), mkfit::MkBuilder::find_tracks_in_layers(), mkfit::MkBuilder::findTracksStandard(), dqm_interfaces.DQMcommunicator::get_runs_list(), and histoStyle::savePlots().

109  {
110  if (1 == nloops % 2) {
111  for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
112  auto m = iv[i];
113  while (m != iv[m])
114  m = iv[m];
115  iv[i] = m;
116  }
117  } else {
118  more = false;
119  for (auto k = threadIdx.x; k < hist.size(); k += blockDim.x) {
120  auto p = hist.begin() + k;
121  auto i = (*p);
122  auto be = std::min(Hist::bin(izt[i]) + 1, int(hist.nbins() - 1));
123  if (nn[i] < minT)
124  continue; // DBSCAN core rule
125  auto loop = [&](uint32_t j) {
126  assert(i != j);
127  if (nn[j] < minT)
128  return; // DBSCAN core rule
129  auto dist = std::abs(zt[i] - zt[j]);
130  if (dist > eps)
131  return;
132  if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
133  return;
134  auto old = atomicMin(&iv[j], iv[i]);
135  if (old != iv[i]) {
136  // end the loop only if no changes were applied
137  more = true;
138  }
139  atomicMin(&iv[i], old);
140  };
141  ++p;
142  for (; p < hist.end(be); ++p)
143  loop(*p);
144  } // for i
145  }
146  if (threadIdx.x == 0)
147  ++nloops;
148  } // while
const dim3 threadIdx
Definition: cudaCompat.h:29
int32_t *__restrict__ iv
const dim3 blockDim
Definition: cudaCompat.h:30
assert(be >=bs)
float const *__restrict__ zt
Abs< T >::type abs(const T &t)
Definition: Abs.h:22
float const *__restrict__ ezt2
int nt
Definition: AMPTWrapper.h:42
__shared__ Hist hist
static constexpr uint32_t nbins()
T1 atomicMin(T1 *a, T2 b)
Definition: cudaCompat.h:85

Variable Documentation

◆ chi2

float const *__restrict__ gpuVertexFinder::chi2 = data.chi2

Definition at line 28 of file gpuFitVertices.h.

◆ chi2Max

WorkSpace float gpuVertexFinder::chi2Max
Initial value:
{
constexpr bool verbose = false
bool verbose

Definition at line 18 of file gpuFitVertices.h.

◆ chi2max

WorkSpace int float float float gpuVertexFinder::chi2max
Initial value:
{
constexpr bool verbose = false
bool verbose

Definition at line 23 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ data

auto &__restrict__ gpuVertexFinder::data = *pdata

Definition at line 31 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and for().

◆ eps

WorkSpace int float gpuVertexFinder::eps

Definition at line 18 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ er2mx

auto gpuVertexFinder::er2mx = errmax * errmax

Definition at line 29 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__().

◆ errmax

WorkSpace int float float gpuVertexFinder::errmax

Definition at line 18 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__().

◆ ezt2

float const *__restrict__ gpuVertexFinder::ezt2 = ws.ezt2

Definition at line 35 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ first

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

Definition at line 28 of file gpuVertexFinder.cc.

◆ fit

auto const& gpuVertexFinder::fit = tracks.stateAtBS

Definition at line 25 of file gpuVertexFinder.cc.

◆ foundClusters

auto gpuVertexFinder::foundClusters = 0

Definition at line 199 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__().

◆ hist

__shared__ Hist gpuVertexFinder::hist

Definition at line 48 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), DDSpecificsHasNamedValueFilter::accept(), DDSpecificsMatchesValueFilter::accept(), DDSpecificsFilter::accept_impl(), PVValHelper::add(), DTT0Calibration::analyze(), HGCalHitValidation::analyzeHGCalSimHit(), APVCyclePhaseMonitor::beginRun(), Book::book(), HLTGenValHistCollFilter::book1D(), HLTDQMFilterEffHists< ProbeType >::book1D(), HLTGenValHistCollFilter::book2D(), HLTDQMFilterEffHists< ProbeType >::book2D(), GEMBaseValidation::bookDetectorOccupancy(), L1TStage2EMTF::bookHistograms(), GEMDQMEfficiencySourceBase::bookNumerator1D(), GEMDQMEfficiencySourceBase::bookNumerator2D(), GEMBaseValidation::bookPIDHist(), GEMBaseValidation::bookZROccupancy(), BTagEntry::BTagEntry(), TrackerG4SimHitNumberingScheme::buildAll(), ClusMultInvestPlots(), Combined2DHisto(), CombinedHisto(), HGCalSimHitValidation::createHisto(), DeadTimeAPVCycle(), fwlite::Scanner< Collection >::draw(), fwlite::Scanner< Collection >::draw2D(), PreparePVTrends::drawConstantWithErr(), fwlite::Scanner< Collection >::drawProf(), DTTMax::DTTMax(), L1TEMTFEventInfoClient::dumpContentMonitorElements(), L1TEventInfoClient::dumpContentMonitorElements(), PrintG4Solids::dumpSummary(), EcalPedHists::endJob(), EcalURecHitHists::endJob(), EcalCosmicsHists::endJob(), DQMHistNormalizer::endRun(), l1tVertexFinder::VertexFinder::fastHisto(), l1tVertexFinder::VertexFinder::fastHistoEmulation(), TotemT2Segmentation::fill(), ExpressionHisto< T >::fill(), DigiCollectionProfiler< T >::fill(), TrigObjTnPHistColl::HistColl::fill(), GEMDQMBase::MEMapInfT< M, K >::Fill(), helper::ScannerBase::fill1D(), helper::ScannerBase::fill2D(), GEMDQMBase::MEMapInfT< M, K >::FillBits(), HLTDQMFilterEffHists< ProbeType >::fillHists(), HLTGenValHistCollFilter::fillHists(), GEMDQMEfficiencySourceBase::fillMEWithinLimits(), edm::Principal::fillPrincipal(), helper::ScannerBase::fillProf(), Rivet::HiggsTemplateCrossSections::finalize(), LA_Filler_Fitter::find_rebin(), PlotAlignmentValidation::fitGauss(), DiMuonMassBiasClient::fitLineShape(), SplitVertexResolution::fitResiduals(), TrackerOfflineValidationSummary::fitResiduals(), PVValHelper::fitResiduals(), TrackerOfflineValidation::fitResiduals(), SplitVertexResolution::fitResiduals_v0(), for(), TrackerOfflineValidation::Fwhm(), PreparePVTrends::getBiases(), EgHLTOfflineSummaryClient::getEgHLTSumHist_(), MuonGEMBaseHarvestor::getElement(), HcalObjRepresent::HcalDataContainer< Items, Item >::GetProjection(), L1TPhase2CorrelatorOffline::getQuantile(), HcalObjRepresent::HcalDataContainer< Items, Item >::GetRange(), fwlite::Scanner< Collection >::getSameH1(), fwlite::Scanner< Collection >::getSameH2(), fwlite::Scanner< Collection >::getSameProf(), SiStripSummaryCreator::getSummaryME(), hTMaxCell::GetT0Factor(), PPSAlignmentHarvester::getTH1DFromTGraphErrors(), PreparePVTrends::getUnrolledHisto(), HGCalTBMB::HGCalTBMB(), HGCalTBMBAnalyzer::HGCalTBMBAnalyzer(), FWEveViewManager::highlightAdded(), HistoAnalyzer< C >::HistoAnalyzer(), MuonResidualsFitter::histogramChi2GaussianFit(), CTPPSCommonDQMSource::GlobalPlots::Init(), EcalDisplaysByEvent::init2DEcalHist(), EcalDisplaysByEvent::init3DEcalHist(), EcalURecHitHists::initHists(), EcalCosmicsHists::initHists(), ExpressionHisto< T >::initialize(), TopDiLeptonOffline::MonitorEnsemble::loggerBinLabels(), main(), LA_Filler_Fitter::make_and_fit_symmchi2(), SiStripSpyDisplayModule::MakeDigiHist_(), DTGeometryValidate::makeHistogram(), GEMGeometryValidate::makeHistogram(), RPCGeometryValidate::makeHistogram(), ME0GeometryValidate::makeHistogram(), CSCGeometryValidate::makeHistogram(), ValidateGeometry::makeHistogram(), Phase1PixelMaps::makeNicePlotStyle(), SiStripPI::makeNicePlotStyle(), SiPixelPI::makeNicePlotStyle(), AlignmentPI::makeNicePlotStyle(), AlignmentPI::makeNiceStats(), SiStripPI::makeNiceStyle(), SiStripSpyDisplayModule::MakeProcessedRawDigiHist_(), SiStripSpyDisplayModule::MakeRawDigiHist_(), DDExpandedView::mergedSpecificsV(), DQMGenericClient::normalizeToEntries(), DQMRivetClient::normalizeToIntegral(), DQMRivetClient::normalizeToLumi(), TrigObjTnPHistColl::HistFiller::operator()(), CSCAlignmentCorrections::plot(), MuonResidualsFitter::plotsimple(), MuonResidualsFitter::plotweighted(), SymmetryFit::pol2_from_pol2(), SymmetryFit::pol2_from_pol3(), printPlot(), EcalPedHists::readEBdigis(), EcalPedHists::readEEdigis(), edm::OneLumiPoolSource::readLuminosityBlockAuxiliary_(), PlotAlignmentValidation::residual_by_moduleID(), PatternOptimizerBase::savePatternsInRoot(), DQMRivetClient::scaleByFactor(), PlotAlignmentValidation::scaleXaxis(), SiPixelUtility::setDrawingOption(), SiStripHistoPlotter::setDrawingOption(), PlotAlignmentValidation::setHistStyle(), DDFilteredView::setScope(), PlotAlignmentValidation::setTitleStyle(), DDExpandedView::specificsV(), PlotAlignmentValidation::storeHistogramInRootfile(), tfwliteselectortest::ThingsTSelector::terminate(), tfwliteselectortest::ThingsTSelector2::terminate(), __class__< T >::terminate(), th1ToFormulaBinTree(), th1ToFormulaLin(), while(), sistrip::EnsembleCalibrationLA::write_ensembles_plots(), sistrip::MeasureLA::write_report_plots(), and sistrip::EnsembleCalibrationLA::write_samples_plots().

◆ hws

__shared__ Hist::Counter gpuVertexFinder::hws

Definition at line 49 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__().

◆ iv

int32_t *__restrict__ gpuVertexFinder::iv = ws.iv

Definition at line 42 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), RPCLinkSynchroStat::add(), TAPD::addEntry(), FBaseSimEvent::addSimTrack(), reco::IsoDeposit::algoWithin(), SiPixelPhase1MonitorVertexSoA::analyze(), HcalRecHitsAnalyzer::analyze(), IsoTrig::analyze(), PrimaryVertexAnalyzer4PUSlimmed::analyze(), IsolatedTracksNxN::analyze(), Primary4DVertexValidation::analyze(), XtalDedxAnalysis::analyzeHits(), HGCalTimingAnalyzer::analyzeSimTracks(), HGCalTBAnalyzer::analyzeSimTracks(), HLTMuonDimuonL3Filter::applyDiMuonSelection(), Basic2DVector< float >::Basic2DVector(), Basic3DVector< align::Scalar >::Basic3DVector(), magneticfield::MagGeoBuilder::build(), MagGeoBuilderFromDDD::build(), HGCalCellUV::cellUVFromXY1(), HGCalCellUV::cellUVFromXY3(), PFPileUpAlgo::chargedHadronVertex(), PrimaryVertexAssignment::chargedHadronVertex(), PFIsolationEstimator::chargedHadronVertex(), IsoTrig::chgIsolation(), PileupJetIdAlgo::computeIdVariables(), MVAJetPuId::computeIdVariables(), SeedMvaEstimator::computeMva(), gen::TauolappInterface::decay(), reco::IsoDeposit::depositAndCountWithin(), TMultiDimFet::EvalControl(), SiPixelActionExecutor::fillFEDErrorSummary(), SiPixelActionExecutor::fillGrandBarrelSummaryHistos(), SiPixelActionExecutor::fillGrandEndcapSummaryHistos(), GsfTrackProducerBase::fillMode(), DeepBoostedJetTagInfoProducer::fillParticleFeatures(), SiPixelActionExecutor::fillSummary(), TrackingFailureFilter::filter(), PythiaDauVFilter::filter(), PythiaAllDauVFilter::filter(), PythiaDauVFilterMatchID::filter(), FFTJetPFPileupCleaner::findSomeVertexWFakes(), DivisiveVertexFinder::findVertexesAlt(), CaloTowerGeometry::getSummary(), CaloSubdetectorGeometry::getSummary(), HcalGeometry::getSummary(), HGCalTypes::getUnpackedV(), reco::Photon::hcalTowerSumEt(), reco::GsfElectron::hcalTowerSumEt(), reco::Photon::hcalTowerSumEtBc(), reco::GsfElectron::hcalTowerSumEtBc(), HGCalCellUV::HGCalCellUV(), edm::OneToMany< std::vector< Trajectory >, std::vector< TrajectorySeed >, unsigned int >::insert(), edm::OneToOneGeneric< std::vector< TrackCandidate >, std::vector< Trajectory >, unsigned int >::insert(), edm::OneToManyWithQualityGeneric< TrackingVertexCollection, edm::View< reco::Vertex >, double, unsigned int >::insert(), FFTJetPFPileupCleaner::isAcceptableVtx(), edm::isTransientEqual(), G4SimEvent::load(), TMultiDimFet::MakeCandidates(), Primary4DVertexValidation::matchReco2Sim(), JetPlusTrackCorrector::matchTracks(), multiTrajectoryStateMode::momentumFromModeLocal(), reco::IsoDeposit::nearestDR(), HGCalTypes::packTypeUV(), HitParentTest::parentSimTrack(), spr::parentSimTrack(), multiTrajectoryStateMode::positionFromModeLocal(), DDHGCalMixLayer::positionMix(), DDHGCalMixRotatedLayer::positionMix(), HGCalMixLayer::positionMix(), HGCalMixRotatedLayer::positionMix(), DDHGCalEEAlgo::positionSensitive(), DDHGCalEEFileAlgo::positionSensitive(), DDHGCalSiliconModule::positionSensitive(), DDHGCalSiliconRotatedModule::positionSensitive(), DDHGCalHEAlgo::positionSensitive(), DDHGCalHEFileAlgo::positionSensitive(), HGCalEEFileAlgo::positionSensitive(), HGCalSiliconModule::positionSensitive(), HGCalSiliconRotatedModule::positionSensitive(), HGCalEEAlgo::PositionSensitive(), HGCalHEFileAlgo::positionSensitive(), HGCalHEAlgo::positionSensitive(), SeedGeneratorFromProtoTracksEDProducer::produce(), JetTracksAssociationDRVertexAssigned::produce(), PFV0Producer::produce(), FFTJetVertexAdder::produce(), PrimaryVertexProducer::produce(), ConvBremSeedProducer::produce(), LmfSource::readEventWithinFile(), ZdcSimpleRecAlgoImpl::reco2(), rpcrawtodigi::RecordCD::RecordCD(), PointSeededTrackingRegionsProducer::regions(), PATPrimaryVertexSelector::select(), MultiTrackSelector::select(), HIMultiTrackSelector::select(), EMShower::setIntervals(), CaloNumberingScheme::setVerbosity(), ZdcNumberingScheme::setVerbosity(), spr::simTrackAtOrigin(), edm::OneToMany< std::vector< Trajectory >, std::vector< TrajectorySeed >, unsigned int >::val(), edm::OneToOneGeneric< std::vector< TrackCandidate >, std::vector< Trajectory >, unsigned int >::val(), edm::OneToManyWithQualityGeneric< TrackingVertexCollection, edm::View< reco::Vertex >, double, unsigned int >::val(), spr::validSimTrack(), HitParentTest::validSimTrack(), ConfigurableTrimmedVertexFinder::vertexCandidates(), HGCalWaferIndex::waferV(), while(), and ZdcNumberingScheme::ZdcNumberingScheme().

◆ izt

uint8_t *__restrict__ gpuVertexFinder::izt = ws.izt

Definition at line 40 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ maxChi2

WorkSpace float gpuVertexFinder::maxChi2
Initial value:
{
constexpr bool verbose = false
bool verbose

Definition at line 15 of file gpuSplitVertices.h.

◆ maxChi2ForFinalFit

constexpr float gpuVertexFinder::maxChi2ForFinalFit = 5000.f

Definition at line 16 of file gpuVertexFinder.cc.

Referenced by gpuVertexFinder::Producer::make().

◆ maxChi2ForFirstFit

constexpr float gpuVertexFinder::maxChi2ForFirstFit = 50.f

Definition at line 15 of file gpuVertexFinder.cc.

Referenced by gpuVertexFinder::Producer::make().

◆ maxChi2ForSplit

constexpr float gpuVertexFinder::maxChi2ForSplit = 9.f

Definition at line 19 of file gpuVertexFinder.cc.

Referenced by gpuVertexFinder::Producer::make().

◆ minT

WorkSpace int gpuVertexFinder::minT

Definition at line 18 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ more

bool gpuVertexFinder::more = true

◆ n

size d for d tracks gpuVertexFinder::n

Definition at line 56 of file gpuClusterTracksDBSCAN.h.

◆ nloops

gpuVertexFinder::nloops = 0

Definition at line 102 of file gpuClusterTracksIterative.h.

Referenced by while().

◆ nn

int32_t const *__restrict__ gpuVertexFinder::nn = data.ndof

Definition at line 41 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ noise

gpuVertexFinder::noise

Definition at line 50 of file gpuFitVertices.h.

◆ nt

auto gpuVertexFinder::nt = ws.ntrks

Definition at line 33 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and while().

◆ nvFinal

uint32_t & gpuVertexFinder::nvFinal = data.nvFinal

Definition at line 37 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__().

◆ nvIntermediate

uint32_t & gpuVertexFinder::nvIntermediate = ws.nvIntermediate

Definition at line 38 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__().

◆ ptMax

ZVertexSoA WorkSpace float float gpuVertexFinder::ptMax
Initial value:
{
uint32_t const *__restrict__ TkSoA const *__restrict__ ptracks

Definition at line 21 of file gpuVertexFinder.cc.

Referenced by for(), and gpuVertexFinder::Producer::make().

◆ ptMin

ZVertexSoA WorkSpace float gpuVertexFinder::ptMin

Definition at line 21 of file gpuVertexFinder.cc.

Referenced by for(), and gpuVertexFinder::Producer::make().

◆ ptt2

float const* __restrict__ gpuVertexFinder::ptt2 = ws.ptt2

Definition at line 22 of file gpuSortByPt2.h.

◆ ptv2

float* __restrict__ gpuVertexFinder::ptv2 = data.ptv2

Definition at line 26 of file gpuSortByPt2.h.

◆ pws

WorkSpace * gpuVertexFinder::pws
Initial value:
{
auto& __restrict__ data = *pdata
char data[epos_bytes_allocation]
Definition: EPOS_Wrapper.h:79

Definition at line 18 of file gpuClusterTracksDBSCAN.h.

Referenced by __attribute__(), and for().

◆ quality

auto const* gpuVertexFinder::quality = tracks.qualityData()

Definition at line 26 of file gpuVertexFinder.cc.

Referenced by for().

◆ soa

ZVertexSoA* gpuVertexFinder::soa

◆ sortInd

uint16_t* __restrict__ gpuVertexFinder::sortInd = data.sortInd

Definition at line 27 of file gpuSortByPt2.h.

Referenced by if().

◆ tracks

auto const& gpuVertexFinder::tracks = *ptracks

Definition at line 24 of file gpuVertexFinder.cc.

Referenced by for().

◆ ws

auto &__restrict__ gpuVertexFinder::ws = *pws

◆ wv

float *__restrict__ gpuVertexFinder::wv = data.wv

◆ zt

float const *__restrict__ gpuVertexFinder::zt = ws.zt

◆ zv

float *__restrict__ gpuVertexFinder::zv = data.zv