CMS 3D CMS Logo

Typedefs | Functions | Variables
caHitNtupletGeneratorKernels Namespace Reference

Typedefs

template<typename TrackerTraits >
using CAParams = caHitNtupletGenerator::CAParamsT< TrackerTraits >
 
using Cell = GPUCACellT< TrackerTraits >
 
template<typename TrackerTraits >
using CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
 
template<typename TrackerTraits >
using CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
 
using Counters = caHitNtupletGenerator::Counters
 
template<typename TrackerTraits >
using HitContainer = typename TrackSoA< TrackerTraits >::HitContainer
 
template<typename TrackerTraits >
using HitsConstView = typename GPUCACellT< TrackerTraits >::HitsConstView
 
template<typename TrackerTraits >
using HitToTuple = caStructures::HitToTupleT< TrackerTraits >
 
template<typename TrackerTraits >
using OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
 
using Quality = pixelTrack::Quality
 
template<typename TrackerTraits >
using QualityCuts = pixelTrack::QualityCutsT< TrackerTraits >
 
template<typename TrackerTraits >
using TkSoAView = TrackSoAView< TrackerTraits >
 
template<typename TrackerTraits >
using TupleMultiplicity = caStructures::TupleMultiplicityT< TrackerTraits >
 

Functions

template<typename TrackerTraits >
 __attribute__ ((always_inline)) void kernel_classifyTracks(TkSoAView< TrackerTraits > tracks_view
 TODO : why there was quality here? More...
 
 assert (nCells)
 cannot be loose More...
 
 for (int idx=first, nt=(*nCells);idx< nt;idx+=gridDim.x *blockDim.x)
 
 if (0==first) = ntracks
 

Variables

TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounterapc
 
auto & c = *counters
 
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__ cells
 
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const * cellTracks
 
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 int32_t uint32_t Counterscounters
 
QualityCuts< TrackerTraits > cuts
 
uint32_t const *__restrict__ TkSoAView< TrackerTraits > bool dupPassThrough
 
auto first = threadIdx.x + blockIdx.x * blockDim.x
 
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t firstPrint
 
auto const & foundNtuplets = *ptuples
 
auto const good = pixelTrack::Quality::strict
 min quality of good More...
 
HitsConstView< TrackerTraits > hh
 
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const * hitToTuple
 
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t int32_t int iev
 
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 l1end = hh.hitsLayerStart()[1]
 
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t int32_t lastPrint
 
auto const longTqual = pixelTrack::Quality::highPurity
 
constexpr auto loose = pixelTrack::Quality::loose
 
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 int32_t uint32_t maxNumberOfDoublets
 
constexpr float maxScore = std::numeric_limits<float>::max()
 
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ nCells
 
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 int32_t nHits
 
auto nhits = hh.nHits()
 
uint16_t nmin
 
constexpr float nSigma2 = 25.f
 
auto ntracks = std::min<int>(apc->get().m, tracks_view.metadata().size() - 1)
 
TkSoAView< TrackerTraits > GPUCACellT< TrackerTraits > *__restrict__ uint32_t const CellTracksVector< TrackerTraits > cms::cuda::AtomicPairCounter CAParams< TrackerTraits > params
 
HitContainer< TrackerTraits > const *__restrict__ Quality const *__restrict__ HitToTuple< TrackerTraits > const *__restrict__ phitToTuple
 
HitContainer< TrackerTraits > const *__restrict__ ptuples
 
HitContainer< TrackerTraits > const *__restrict__ Quality const *__restrict__ quality
 
auto const reject = dupPassThrough ? loose : dup
 
constexpr uint32_t tkNotFound = std::numeric_limits<uint16_t>::max()
 
uint32_t const *__restrict__ TkSoAView< TrackerTraits > tracks_view
 
TupleMultiplicity< TrackerTraits > const * tupleMultiplicity
 
HitContainer< TrackerTraits > const *__restrict__ tuples
 

Typedef Documentation

◆ CAParams

template<typename TrackerTraits >
using caHitNtupletGeneratorKernels::CAParams = typedef caHitNtupletGenerator::CAParamsT<TrackerTraits>

Definition at line 64 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Cell

using caHitNtupletGeneratorKernels::Cell = typedef GPUCACellT<TrackerTraits>

Definition at line 418 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ CellNeighborsVector

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

Definition at line 41 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ CellTracksVector

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

Definition at line 44 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Counters

Definition at line 66 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitContainer

template<typename TrackerTraits >
using caHitNtupletGeneratorKernels::HitContainer = typedef typename TrackSoA<TrackerTraits>::HitContainer

Definition at line 55 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitsConstView

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

Definition at line 58 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitToTuple

template<typename TrackerTraits >
using caHitNtupletGeneratorKernels::HitToTuple = typedef caStructures::HitToTupleT<TrackerTraits>

Definition at line 35 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ OuterHitOfCell

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

Definition at line 47 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Quality

Definition at line 49 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ QualityCuts

template<typename TrackerTraits >
using caHitNtupletGeneratorKernels::QualityCuts = typedef pixelTrack::QualityCutsT<TrackerTraits>

Definition at line 61 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ TkSoAView

template<typename TrackerTraits >
using caHitNtupletGeneratorKernels::TkSoAView = typedef TrackSoAView<TrackerTraits>

Definition at line 52 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ TupleMultiplicity

template<typename TrackerTraits >
using caHitNtupletGeneratorKernels::TupleMultiplicity = typedef caStructures::TupleMultiplicityT<TrackerTraits>

Definition at line 38 of file CAHitNtupletGeneratorKernelsImpl.h.

Function Documentation

◆ __attribute__()

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

TODO : why there was quality here?

Definition at line 299 of file CAHitNtupletGeneratorKernelsImpl.h.

References cms::cudacompat::__ldg(), cms::cudacompat::blockDim, cms::cudacompat::blockIdx, cellNeighbors, cells, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), first, cms::cudacompat::gridDim, hh, heavyIonCSV_trainingSettings::idx, PixelPluginsPhase0_cfi::isBarrel, isOuterHitOfCell, dqmiolumiharvest::j, GPUCACellT< TrackerTraits >::kUsed, nCells, nt, params, riemannFit::stride, and cms::cudacompat::threadIdx.

306  {
308 
309  auto firstCellIndex = threadIdx.y + blockIdx.y * blockDim.y;
310  auto first = threadIdx.x;
311  auto stride = blockDim.x;
312 
313  if (0 == (firstCellIndex + first)) {
314  (*apc1) = 0;
315  (*apc2) = 0;
316  } // ready for next kernel
317 
318  constexpr uint32_t last_bpix1_detIndex = TrackerTraits::last_bpix1_detIndex;
319  constexpr uint32_t last_barrel_detIndex = TrackerTraits::last_barrel_detIndex;
320  for (int idx = firstCellIndex, nt = (*nCells); idx < nt; idx += gridDim.y * blockDim.y) {
321  auto cellIndex = idx;
322  auto &thisCell = cells[idx];
323  auto innerHitId = thisCell.inner_hit_id();
324  if (int(innerHitId) < isOuterHitOfCell.offset)
325  continue;
326  int numberOfPossibleNeighbors = isOuterHitOfCell[innerHitId].size();
327  auto vi = isOuterHitOfCell[innerHitId].data();
328 
329  auto ri = thisCell.inner_r(hh);
330  auto zi = thisCell.inner_z(hh);
331 
332  auto ro = thisCell.outer_r(hh);
333  auto zo = thisCell.outer_z(hh);
334  auto isBarrel = thisCell.inner_detIndex(hh) < last_barrel_detIndex;
335 
336  for (int j = first; j < numberOfPossibleNeighbors; j += stride) {
337  auto otherCell = __ldg(vi + j);
338  auto &oc = cells[otherCell];
339  auto r1 = oc.inner_r(hh);
340  auto z1 = oc.inner_z(hh);
341  bool aligned = Cell::areAlignedRZ(
342  r1,
343  z1,
344  ri,
345  zi,
346  ro,
347  zo,
348  params.ptmin_,
349  isBarrel ? params.CAThetaCutBarrel_ : params.CAThetaCutForward_); // 2.f*thetaCut); // FIXME tune cuts
350  if (aligned && thisCell.dcaCut(hh,
351  oc,
352  oc.inner_detIndex(hh) < last_bpix1_detIndex ? params.dcaCutInnerTriplet_
353  : params.dcaCutOuterTriplet_,
354  params.hardCurvCut_)) { // FIXME tune cuts
355  oc.addOuterNeighbor(cellIndex, *cellNeighbors);
356  thisCell.setStatusBits(Cell::StatusBit::kUsed);
357  oc.setStatusBits(Cell::StatusBit::kUsed);
358  }
359  } // loop on inner cells
360  } // loop on outer cells
361  }
const dim3 threadIdx
Definition: cudaCompat.h:29
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
const dim3 gridDim
Definition: cudaCompat.h:33
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const * cellNeighbors
const dim3 blockDim
Definition: cudaCompat.h:30
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ cells
TkSoAView< TrackerTraits > GPUCACellT< TrackerTraits > *__restrict__ uint32_t const CellTracksVector< TrackerTraits > cms::cuda::AtomicPairCounter CAParams< TrackerTraits > params
constexpr uint32_t stride
Definition: HelixFit.h:22
const dim3 blockIdx
Definition: cudaCompat.h:32
int nt
Definition: AMPTWrapper.h:42
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

◆ assert()

caHitNtupletGeneratorKernels::assert ( nCells  )

cannot be loose

◆ for()

caHitNtupletGeneratorKernels::for ( int  idx = first)

Definition at line 136 of file CAHitNtupletGeneratorKernelsImpl.h.

References cms::cudacompat::atomicAdd(), c, cells, hitToTuple, and heavyIonCSV_trainingSettings::idx.

136  {
137  auto const &thisCell = cells[idx];
138  if (thisCell.hasFishbone() && !thisCell.isKilled())
139  atomicAdd(&c.nFishCells, 1);
140  if (thisCell.outerNeighbors().full()) //++tooManyNeighbors[thisCell.theLayerPairId];
141  printf("OuterNeighbors overflow %d in %d\n", idx, thisCell.layerPairId());
142  if (thisCell.tracks().full()) //++tooManyTracks[thisCell.theLayerPairId];
143  printf("Tracks overflow %d in %d\n", idx, thisCell.layerPairId());
144  if (thisCell.isKilled())
145  atomicAdd(&c.nKilledCells, 1);
146  if (!thisCell.unused())
147  atomicAdd(&c.nEmptyCells, 1);
148  if ((0 == hitToTuple->size(thisCell.inner_hit_id())) && (0 == hitToTuple->size(thisCell.outer_hit_id())))
149  atomicAdd(&c.nZeroTrackCells, 1);
150  }
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ cells
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const * hitToTuple
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61

◆ if()

caHitNtupletGeneratorKernels::if ( = =first) = ntracks

Definition at line 85 of file CAHitNtupletGeneratorKernelsImpl.h.

References apc, cms::cudacompat::atomicAdd(), c, cms::cuda::AtomicPairCounter::get(), nCells, nHits, and tupleMultiplicity.

85  {
86  atomicAdd(&c.nEvents, 1);
87  atomicAdd(&c.nHits, nHits);
88  atomicAdd(&c.nCells, *nCells);
89  atomicAdd(&c.nTuples, apc->get().m);
90  atomicAdd(&c.nFitTracks, tupleMultiplicity->size());
91  }
TupleMultiplicity< TrackerTraits > const * tupleMultiplicity
__device__ __host__ Counters get() const
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ nCells
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter * apc
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t nHits
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61

Variable Documentation

◆ apc

cms::cuda::AtomicPairCounter * caHitNtupletGeneratorKernels::apc

◆ c

auto& caHitNtupletGeneratorKernels::c = *counters

◆ cellNeighbors

TupleMultiplicity<TrackerTraits> const HitToTuple<TrackerTraits> const cms::cuda::AtomicPairCounter GPUCACellT<TrackerTraits> const* __restrict__ uint32_t const* __restrict__ CellNeighborsVector<TrackerTraits> const* caHitNtupletGeneratorKernels::cellNeighbors

◆ cells

TkSoAView< TrackerTraits > GPUCACellT< TrackerTraits > *__restrict__ caHitNtupletGeneratorKernels::cells

◆ cellTracks

TkSoAView< TrackerTraits > GPUCACellT< TrackerTraits > *__restrict__ uint32_t const CellTracksVector< TrackerTraits > * caHitNtupletGeneratorKernels::cellTracks

◆ counters

Counters * caHitNtupletGeneratorKernels::counters
Initial value:

Definition at line 80 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by L1TStage2EMTF::analyze(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_printCounters::operator()(), CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::printCounters(), RPCTwinMuxRawToDigi::processBlock(), RPCAMCRawToDigi::processBlocks(), RPCAMCRawToDigi::processCDFHeaders(), RPCTwinMuxRawToDigi::processCDFHeaders(), RPCAMCRawToDigi::processCDFTrailers(), RPCTwinMuxRawToDigi::processCDFTrailers(), RPCCPPFUnpacker::processCPPF(), RPCTwinMuxRawToDigi::processRPCRecord(), RPCCPPFUnpacker::processRXRecord(), RPCTwinMuxRawToDigi::processTwinMux(), RPCCPPFUnpacker::produce(), RPCAMCRawToDigi::produce(), CSCTFPacker::produce(), RPCTwinMuxRawToDigi::produce(), RPCTwinMuxRawToDigi::putCounters(), and CAHitNtupletGeneratorKernels< cms::cudacompat::CPUTraits, TrackerTraits >::setCounters().

◆ cuts

QualityCuts<TrackerTraits> caHitNtupletGeneratorKernels::cuts

◆ dupPassThrough

uint16_t bool caHitNtupletGeneratorKernels::dupPassThrough

◆ first

int caHitNtupletGeneratorKernels::first = threadIdx.x + blockIdx.x * blockDim.x

Definition at line 164 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by __attribute__().

◆ firstPrint

TkSoAView<TrackerTraits> HitToTuple<TrackerTraits> const* __restrict__ int32_t caHitNtupletGeneratorKernels::firstPrint

◆ foundNtuplets

auto const& caHitNtupletGeneratorKernels::foundNtuplets = *ptuples

◆ good

auto const caHitNtupletGeneratorKernels::good = pixelTrack::Quality::strict

min quality of good

Definition at line 761 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by HGcalHitIdCheck::analyze(), HGCalTestScintHits::analyze(), HGCalTestPartialWaferRecHits::analyze(), HGCalTestPartialWaferHits::analyze(), HcalTestSimHitID::analyze(), HGCalTestGuardRing::analyze(), TkConvValidator::analyze(), PhotonValidator::analyze(), heppy::CMGMuonCleanerBySegmentsAlgo::clean(), HcalDbASCIIIO::createObject< HcalFrontEndMap >(), HcalDbASCIIIO::createObject< HcalSiPMCharacteristics >(), dumpLutDiff(), gpuClustering::for(), GctDigiToPsbText::GctDigiToPsbText(), HGCalWaferMask::goodCell(), GtPsbTextToDigi::GtPsbTextToDigi(), PFBadHcalPseudoClusterProducer::init(), EcalRecHitSimpleAlgo::makeRecHit(), MatcherByPullsAlgorithm::match(), MatcherByPullsAlgorithm::matchMany(), magneticfield::interpolation::binary_ifstream::operator bool(), pixelClustering::ClusterChargeCut< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_tripletCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_simpleTripletCleaner< TrackerTraits >::operator()(), FastFedCablingHistograms::printSummary(), CommissioningHistograms::printSummary(), GtPsbTextToDigi::produce(), modules::MuonCleanerBySegmentsT< T >::produce(), l1t::L1ComparatorRun2::produce(), RctTextToRctDigi::produce(), RctDigiToRctText::RctDigiToRctText(), RctTextToRctDigi::RctTextToRctDigi(), readCMSSWcoeff(), readCMSSWcoeffForComparison(), edm::ProductSelectorRules::Rule::Rule(), L1MuDTEtaProcessor::runEtaTrackFinder(), L1MuBMEtaProcessor::runEtaTrackFinder(), MuonSeedCleaner::SeedCandidates(), MuonOverlapSeedFromRecHits::seeds(), CosmicMuonSeedGenerator::selectSegments(), StMeasurementConditionSet::set128StripStatus(), and TkStripMeasurementDet::set128StripStatus().

◆ hh

HitsConstView<TrackerTraits> caHitNtupletGeneratorKernels::hh

◆ hitToTuple

auto & caHitNtupletGeneratorKernels::hitToTuple
Initial value:

Definition at line 70 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by for(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countSharedHit< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_rejectDuplicate< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_sharedHitCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_tripletCleaner< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_simpleTripletCleaner< TrackerTraits >::operator()().

◆ iev

TkSoAView<TrackerTraits> HitToTuple<TrackerTraits> const* __restrict__ int32_t int32_t int caHitNtupletGeneratorKernels::iev
Initial value:

Definition at line 858 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by TestCUDAProducerGPUEWTask::acquire(), ProtonTransport::addPartToHepMC(), Primary4DVertexValidation::recoPrimaryVertex::addTrack(), HcalForwardLibWriter::analyze(), L1ExtraTestAnalyzer::analyze(), Primary4DVertexValidation::analyze(), ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::classifyTuples(), CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::classifyTuples(), fwlite::Scanner< Collection >::count(), fwlite::Scanner< Collection >::countEvents(), fwlite::Scanner< Collection >::draw(), fwlite::Scanner< Collection >::draw2D(), fwlite::Scanner< Collection >::drawGraph(), fwlite::Scanner< Collection >::drawProf(), fwlite::Scanner< Collection >::fillDataSet(), Primary4DVertexValidation::matchReco2Sim(), lheh5::Events::mkEventHeader(), lheh5::Events2::mkEventHeader(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_print_found_ntuplets< TrackerTraits >::operator()(), fwlite::Scanner< Collection >::scan(), CheckSecondary::update(), SimG4FluxProducer::update(), HGCPassive::update(), TotemTestGem::update(), MaterialBudgetVolume::update(), SimG4HGCalValidation::update(), SimG4HcalValidation::update(), and RootTreeHandler::writeTree().

◆ isOuterHitOfCell

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 caHitNtupletGeneratorKernels::isOuterHitOfCell

◆ l1end

int caHitNtupletGeneratorKernels::l1end = hh.hitsLayerStart()[1]

◆ lastPrint

TkSoAView<TrackerTraits> HitToTuple<TrackerTraits> const* __restrict__ int32_t int32_t caHitNtupletGeneratorKernels::lastPrint

◆ longTqual

auto const caHitNtupletGeneratorKernels::longTqual = pixelTrack::Quality::highPurity

◆ loose

constexpr auto caHitNtupletGeneratorKernels::loose = pixelTrack::Quality::loose

◆ maxNumberOfDoublets

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 int32_t uint32_t caHitNtupletGeneratorKernels::maxNumberOfDoublets

◆ maxScore

constexpr float caHitNtupletGeneratorKernels::maxScore = std::numeric_limits<float>::max()

◆ nCells

uint32_t const * caHitNtupletGeneratorKernels::nCells

◆ nHits

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 int32_t caHitNtupletGeneratorKernels::nHits

Definition at line 70 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by if().

◆ nhits

auto caHitNtupletGeneratorKernels::nhits = hh.nHits()

◆ nmin

uint16_t caHitNtupletGeneratorKernels::nmin

◆ nSigma2

constexpr float caHitNtupletGeneratorKernels::nSigma2 = 25.f

◆ ntracks

auto caHitNtupletGeneratorKernels::ntracks = std::min<int>(apc->get().m, tracks_view.metadata().size() - 1)

◆ params

TkSoAView<TrackerTraits> GPUCACellT<TrackerTraits>* __restrict__ uint32_t const CellTracksVector<TrackerTraits> cms::cuda::AtomicPairCounter CAParams<TrackerTraits> caHitNtupletGeneratorKernels::params

◆ phitToTuple

TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ caHitNtupletGeneratorKernels::phitToTuple

◆ ptuples

HitContainer<TrackerTraits> const* __restrict__ caHitNtupletGeneratorKernels::ptuples

◆ quality

HitContainer< TrackerTraits > const *__restrict__ Quality *__restrict__ caHitNtupletGeneratorKernels::quality

Definition at line 589 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ reject

auto const caHitNtupletGeneratorKernels::reject = dupPassThrough ? loose : dup

◆ tkNotFound

constexpr uint32_t caHitNtupletGeneratorKernels::tkNotFound = std::numeric_limits<uint16_t>::max()

◆ tracks_view

TkSoAView< TrackerTraits > caHitNtupletGeneratorKernels::tracks_view
Initial value:

Definition at line 161 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::classifyTuples(), CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::classifyTuples(), gpuVertexFinder::for(), ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::launchKernels(), CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::launchKernels(), gpuVertexFinder::Producer< TrackerTraits >::make(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::LoadTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fishboneCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_earlyDuplicateRemover< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fastDuplicateRemover< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_find_ntuplets< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countMultiplicity< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillMultiplicity< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_classifyTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillHitDetIndices< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillNLayers< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_rejectDuplicate< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_sharedHitCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_tripletCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_simpleTripletCleaner< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_print_found_ntuplets< TrackerTraits >::operator()().

◆ tupleMultiplicity

TupleMultiplicity< TrackerTraits > * caHitNtupletGeneratorKernels::tupleMultiplicity

◆ tuples

HitContainer<TrackerTraits> const* __restrict__ caHitNtupletGeneratorKernels::tuples