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

◆ Cell

using caHitNtupletGeneratorKernels::Cell = typedef GPUCACellT<TrackerTraits>

Definition at line 429 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ CellNeighborsVector

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

Definition at line 42 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ CellTracksVector

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

Definition at line 45 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Counters

Definition at line 67 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitContainer

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

Definition at line 56 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitsConstView

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

Definition at line 59 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ HitToTuple

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

Definition at line 36 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ OuterHitOfCell

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

Definition at line 48 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ Quality

Definition at line 50 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ QualityCuts

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

Definition at line 62 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ TkSoAView

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

Definition at line 53 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ TupleMultiplicity

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

Definition at line 39 of file CAHitNtupletGeneratorKernelsImpl.h.

Function Documentation

◆ __attribute__()

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

TODO : why there was quality here?

Definition at line 310 of file CAHitNtupletGeneratorKernelsImpl.h.

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

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

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

147  {
148  auto const &thisCell = cells[idx];
149  if (thisCell.hasFishbone() && !thisCell.isKilled())
150  atomicAdd(&c.nFishCells, 1);
151  if (thisCell.outerNeighbors().full()) //++tooManyNeighbors[thisCell.theLayerPairId];
152  printf("OuterNeighbors overflow %d in %d\n", idx, thisCell.layerPairId());
153  if (thisCell.tracks().full()) //++tooManyTracks[thisCell.theLayerPairId];
154  printf("Tracks overflow %d in %d\n", idx, thisCell.layerPairId());
155  if (thisCell.isKilled())
156  atomicAdd(&c.nKilledCells, 1);
157  if (!thisCell.unused())
158  atomicAdd(&c.nEmptyCells, 1);
159  if ((0 == hitToTuple->size(thisCell.inner_hit_id())) && (0 == hitToTuple->size(thisCell.outer_hit_id())))
160  atomicAdd(&c.nZeroTrackCells, 1);
161  }
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 86 of file CAHitNtupletGeneratorKernelsImpl.h.

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

Referenced by VIDSelectorBase.VIDSelectorBase::__call__(), Mixins._TypedParameterizable::__findDefaultsFor(), RPCLinkSynchroStat::add(), vid_id_tools::addVIDSelectionToPATProducer(), EcalMatacqAnalyzer::analyze(), SiPhase2BadStripChannelReader::analyze(), CSCAFEBConnectAnalysis::analyze(), BTagPerformanceAnalyzerMC::analyze(), HGCGeometryValidation::analyze(), CTPPSProtonReconstructionEfficiencyEstimatorMC::analyze(), HLTrigReport::analyze(), reco::PFDisplacedVertexCandidate::associatedElements(), reco::PFBlock::associatedElements(), OMTFProcessor< GoldenPatternType >::averagePatterns(), HistogramManager::book(), HLTObjectMonitorProtonLead::bookHistograms(), HLTObjectMonitor::bookHistograms(), CosmicMuonTrajectoryBuilder::build(), magneticfield::MagGeoBuilder::buildInterpolator(), ticl::PatternRecognitionbyCLUE3D< TILES >::calculateLocalDensity(), batchmanager.BatchManager::CheckBatchScript(), CondDBESSource::CondDBESSource(), CaloTowersCreationAlgo::convert(), popcon::EcalChannelStatusHandler::cosmicsAnalysis(), DeadROCCounter_Phase1::countBadROCBarrel(), DeadROCCounter_Phase1::countBadROCForward(), python.cmstools::createBranchBuffer(), FWEveDigitSetScalableMarkerGL::DirectDraw(), getPayloadData::discover_plugins(), evf::FastMonitoringService::doSnapshot(), SiPixelLorentzAnglePCLHarvester::dqmEndJob(), SiPixelStatusHarvester::dqmEndRun(), getInfo::dumpIDs(), getInfo::dumpSNs(), EMap::EMap(), hcaldqm::DigiRunSummary::endJob(), ElectronCalibration::endJob(), OccupancyPlots::endRun(), l1t::CorrCondition::evaluateCondition(), l1t::CorrWithOverlapRemovalCondition::evaluateCondition(), trklet::TripletEngine::execute(), runMETCorrectionsAndUncertainties.RunMETCorrectionsAndUncertainties::extractMET(), ecaldqm::fetchAndFill(), dqm::impl::MonitorElement::Fill(), OMTFProcessor< GoldenPatternType >::fillCounts(), FWECALCaloDataDetailViewBuilder::fillData(), NtupleManager::FillFitParameters(), TrackerValidationVariables::fillHitQuantities(), reco::HcalNoiseInfoProducer::fillOtherSummaryVariables(), DTTimingExtractor::fillTiming(), TrackingQualityChecker::fillTrackingStatus(), LogErrorEventFilter::filter(), batchHippy.MyBatchManager::finalize(), ElectronMCTruthFinder::find(), PhotonMCTruthFinder::find(), sistrip::SpyUtilities::findAPVErrorBits(), ContainmentCorrectionAnalyzer::findMcTruth(), plotBaryCentre_VS_BeamSpot::findRunIndex(), GEMVFATStatus::GEMVFATStatus(), genlkupmap(), gen::HydjetHadronizer::get_particles(), python.Lumis::getByLabel(), python.Runs::getByLabel(), python.Events::getByLabel(), compareTotals::getCanvasMainPad(), HcalLutManager::getCompressionLutXmlFromAsciiMaster(), HcalLutManager::getCompressionLutXmlFromCoder(), EcalElectronicsMapper::getDCCId(), HistoManager::getDetIdsForType(), G2GainsValidator::getExpressGT(), HcalLutManager::getHEFineGrainLUTs(), HcalObjRepresent::HcalDataContainer< Items, Item >::getItemFromValCont(), HcalLutManager::getLinearizationLutXmlFromAsciiMasterEmap(), HcalLutManager::getLinearizationLutXmlFromCoder(), HcalLutManager::getLinearizationLutXmlFromCoderEmap(), HcalLutManager::getLutXmlFromAsciiMaster(), MT2Analyzer.MT2Analyzer::getMT2Hemi(), popcon::EcalPulseShapesHandler::getNewObjects(), DTKeyedConfigHandler::getNewObjects(), FastLineRecognition::getOneLine(), OMTFinput::getRefHits(), CSCOfflineMonitor::getSignal(), HcalLutManager::getZdcLutXml(), FWGUIEventFilter::HandleKey(), VIDSelectorBase.VIDSelectorBase::initialize(), SiPixelFedCablingMap::initializeRocs(), SiPixelTemplate::interpolate(), helpers::jetCollectionString(), getInfo::listRuns(), LoadEPDB::LoadEPDB(), HGCalGeomParameters::loadGeometryHexagon8(), HGCalGeomParameters::loadGeometryHexagonModule(), LzmaDec_DecodeToDic(), run_AlCaRecoTriggerBitsUpdateWorkflow::main(), checkRuns::main(), checkPayloads::main(), MultipleCompare::main(), BeamSpotWorkflow::main(), objects.IsoTrackAnalyzer.IsoTrackAnalyzer::makeIsoTrack(), FastTrackerRecHitMaskProducer_cfi::maskProducerFromClusterRemover(), batchHippy.MyBatchManager::mkdir(), batchmanager.BatchManager::mkdir(), GenSpecificAlgo::mkSpecificGenMETData(), MSLayersKeeperX0DetLayer::MSLayersKeeperX0DetLayer(), HGCalDDDConstants::numberCellsHexagon(), DQMNet::onMessage(), CheckHitPattern::operator()(), omtf::operator<<(), operator<<(), hitfit::operator<<(), l1t::TriggerMenuParser::parseCorrelation(), l1t::TriggerMenuParser::parseCorrelationThreeBody(), l1t::TriggerMenuParser::parseCorrelationWithOverlapRemoval(), bigModule::plot(), subModule::plot(), plotBaryCentre_VS_BeamSpot::plotbarycenter(), TritonService::preBeginJob(), MatrixInjector.MatrixInjector::prepare(), cms::DDFilteredView::printFilter(), studyJets::printJetProperties(), AlCaHLTBitMon_QueryRunRegistry::printLumi(), TTUSectorORLogic::process(), PPSAlignmentWorker::SectorData::process(), DDLSpecPar::processElement(), OMTFProcessor< GoldenPatternType >::processInput(), RecHitProcessor::processLook(), PATTauHybridProducer::produce(), trklet::KFin::produce(), PuppiProducer::produce(), L1Comparator::produce(), ecaldqm::PedestalClient::producePlots(), spr::propagateHCAL(), SiPixelTemplate::qbin(), LMap::impl::read(), plotBaryCentre_VS_BeamSpot::readBaryCentreAnalyzerTree(), FedRawDataInputSource::readWorker(), cms::DDNamespace::rotation(), production_tasks.CheckDatasetExists::run(), plotBaryCentre_VS_BeamSpot::Run(), ecaldqm::LedTask::runOnUncalibRecHits(), ecaldqm::LaserTask::runOnUncalibRecHits(), HGCDigitizerBase::runSimple(), CondDBESSource::setIntervalFor(), cms::DDNamespace::solid(), batchmanager.BatchManager::SubmitJobs(), vid_id_tools::switchOnVIDElectronIdProducer(), vid_id_tools::switchOnVIDMuonIdProducer(), vid_id_tools::switchOnVIDPhotonIdProducer(), pfTools::switchToPFMET(), dqm::impl::MonitorElement::syncCoreObject(), SiPixelTemplate::temperrors(), data_formats::to_datatables(), coreTools.RemoveMCMatching::toolCode(), trackTools.MakePATTrackCandidates::toolCode(), CaloTowerConstituentsMap::towerOf(), HGVHistoProducerAlgo::tracksters_to_SimTracksters(), CosmicMuonTrajectoryBuilder::trajectories(), omtf::translatePact2Omtf(), CSCDDUEventData::unpack_data(), reco::KalmanGhostTrackUpdater::update(), CountProcessesAction::update(), run_AlCaRecoTriggerBitsUpdateWorkflow::updateBits(), and OMTFPatternMaker::writeMergedGPs().

86  {
87  atomicAdd(&c.nEvents, 1);
88  atomicAdd(&c.nHits, nHits);
89  atomicAdd(&c.nCells, *nCells);
90  atomicAdd(&c.nTuples, apc->get().m);
91  atomicAdd(&c.nFitTracks, tupleMultiplicity->size());
92  }
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

Definition at line 84 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by for(), and if().

◆ 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

Definition at line 71 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by __attribute__(), and for().

◆ cellTracks

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

Definition at line 71 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ counters

Counters * caHitNtupletGeneratorKernels::counters

◆ cuts

QualityCuts<TrackerTraits> caHitNtupletGeneratorKernels::cuts
Initial value:

Definition at line 475 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ dupPassThrough

uint16_t bool caHitNtupletGeneratorKernels::dupPassThrough

◆ first

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

Definition at line 175 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by __attribute__().

◆ firstPrint

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

Definition at line 865 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ foundNtuplets

auto const& caHitNtupletGeneratorKernels::foundNtuplets = *ptuples

◆ good

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

min quality of good

Definition at line 772 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by HGcalHitIdCheck::analyze(), HGCalTestScintHits::analyze(), HGCalTestPartialWaferRecHits::analyze(), HGCalTestPartialWaferHits::analyze(), HcalTestSimHitID::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(), 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 71 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by for().

◆ iev

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

◆ 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

Definition at line 71 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by __attribute__().

◆ l1end

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

Definition at line 726 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ lastPrint

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

Definition at line 865 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ longTqual

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

Definition at line 722 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ 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

Definition at line 71 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ maxScore

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

Definition at line 30 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by FastTrajectoryCleaner::clean().

◆ 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 71 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

Definition at line 31 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ ntracks

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

Definition at line 574 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ params

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

Definition at line 381 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by __attribute__().

◆ phitToTuple

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

◆ ptuples

HitContainer<TrackerTraits> const* __restrict__ caHitNtupletGeneratorKernels::ptuples

Definition at line 600 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ quality

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

Definition at line 600 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ reject

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

◆ tkNotFound

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

Definition at line 29 of file CAHitNtupletGeneratorKernelsImpl.h.

◆ tracks_view

TkSoAView< TrackerTraits > caHitNtupletGeneratorKernels::tracks_view

◆ tupleMultiplicity

TupleMultiplicity< TrackerTraits > * caHitNtupletGeneratorKernels::tupleMultiplicity
Initial value:

Definition at line 71 of file CAHitNtupletGeneratorKernelsImpl.h.

Referenced by if().

◆ tuples

HitContainer<TrackerTraits> const* __restrict__ caHitNtupletGeneratorKernels::tuples

Definition at line 637 of file CAHitNtupletGeneratorKernelsImpl.h.