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 66 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Cell

using caHitNtupletGeneratorKernels::Cell = typedef GPUCACellT<TrackerTraits>

Definition at line 420 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ CellNeighborsVector

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

Definition at line 43 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ CellTracksVector

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

Definition at line 46 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Counters

Definition at line 68 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitContainer

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

Definition at line 57 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitsConstView

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

Definition at line 60 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitToTuple

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

Definition at line 37 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ OuterHitOfCell

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

Definition at line 49 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Quality

Definition at line 51 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ QualityCuts

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

Definition at line 63 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ TkSoAView

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

Definition at line 54 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ TupleMultiplicity

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

Definition at line 40 of file CAHitNtupletGeneratorKernelsImpl.h.

Function Documentation

◆ __attribute__()

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

TODO : why there was quality here?

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

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

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

138  {
139  auto const &thisCell = cells[idx];
140  if (thisCell.hasFishbone() && !thisCell.isKilled())
141  atomicAdd(&c.nFishCells, 1);
142  if (thisCell.outerNeighbors().full()) //++tooManyNeighbors[thisCell.theLayerPairId];
143  printf("OuterNeighbors overflow %d in %d\n", idx, thisCell.layerPairId());
144  if (thisCell.tracks().full()) //++tooManyTracks[thisCell.theLayerPairId];
145  printf("Tracks overflow %d in %d\n", idx, thisCell.layerPairId());
146  if (thisCell.isKilled())
147  atomicAdd(&c.nKilledCells, 1);
148  if (!thisCell.unused())
149  atomicAdd(&c.nEmptyCells, 1);
150  if ((0 == hitToTuple->size(thisCell.inner_hit_id())) && (0 == hitToTuple->size(thisCell.outer_hit_id())))
151  atomicAdd(&c.nZeroTrackCells, 1);
152  }
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 87 of file CAHitNtupletGeneratorKernelsImpl.h.

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

87  {
88  atomicAdd(&c.nEvents, 1);
89  atomicAdd(&c.nHits, nHits);
90  atomicAdd(&c.nCells, *nCells);
91  atomicAdd(&c.nTuples, apc->get().m);
92  atomicAdd(&c.nFitTracks, tupleMultiplicity->size());
93  }
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 82 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 166 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 763 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 72 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 860 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(), HGCalTBPassive::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 72 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 591 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 163 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