gpuPixelDoublets Namespace Reference

Classes
struct	CellCutsT

Typedefs
template<typename TrackerTraits >
using	CellNeighbors = caStructures::CellNeighborsT< TrackerTraits >
template<typename TrackerTraits >
using	CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
template<typename TrackerTraits >
using	CellTracks = caStructures::CellTracksT< TrackerTraits >
template<typename TrackerTraits >
using	CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
template<typename TrackerTraits >
using	HitsConstView = typename GPUCACellT< TrackerTraits >::HitsConstView
template<typename TrackerTraits >
using	OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
using	PhiBinner = typename TrackingRecHitSoA< TrackerTraits >::PhiBinner

Functions
template<typename TrackerTraits >
	__attribute__ ((always_inline)) void fishbone(HitsConstView< TrackerTraits > hh
template<typename TrackerTraits >
__device__	__attribute__ ((always_inline)) void doubletsFromHisto(uint32_t nPairs
	__syncthreads ()
	assert (offsets)
	for (int i=first;i< nHits - isOuterHitOfCell.offset;i+=gridDim.x *blockDim.x) isOuterHitOfCell.container[i].reset()
	for (int idy=firstY, nt=nHits - offset;idy< nt;idy+=gridDim.y *blockDim.y)
	if (0==first)
	if (threadIdx.y==0 &&threadIdx.x==0)

Variables
uint32_t	cc [maxCellsPerHit]
int CellNeighborsVector< TrackerTraits > *	cellNeighbors
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > *	cellNeighborsContainer
GPUCACellT< TrackerTraits > *	cells
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > *	cellTracks
int CellNeighborsVector< TrackerTraits > CellNeighbors< TrackerTraits > CellTracksVector< TrackerTraits > CellTracks< TrackerTraits > *	cellTracksContainer
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const int32_t bool	checkTrack
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int CellCutsT< TrackerTraits > *const	cuts
uint16_t	d [maxCellsPerHit]
const bool	doPtCut = cuts.doPtCut_
const bool	doZ0Cut = cuts.doZ0Cut_
int	first = blockIdx.x * blockDim.x + threadIdx.x
auto	firstX = threadIdx.x
auto	firstY = threadIdx.y + blockIdx.y * blockDim.y
constexpr auto	getDoubletsFromHistoMaxBlockSize = 64
constexpr auto	getDoubletsFromHistoMinBlocksPerMP = 16
const float	hardPtCut = cuts.ptCut_
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits >	hh
auto	idy = blockIdx.y * blockDim.y + threadIdx.y
__shared__ uint32_t	innerLayerCumulativeSize [TrackerTraits::nPairs]
auto const	isOuterHitOfCell = isOuterHitOfCellWrap.container
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const	isOuterHitOfCellWrap
uint8_t	l [maxCellsPerHit]
auto	layerSize = [=](uint8_t li) { return offsets[li + 1] - offsets[li]; }
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int const int	maxNumOfDoublets
const float	minRadius = hardPtCut * 87.78f
const float	minRadius2T4 = 4.f * minRadius * minRadius
float	n [maxCellsPerHit]
uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > int	nActualPairs
GPUCACellT< TrackerTraits > uint32_t const *__restrict__	nCells
GPUCACellT< TrackerTraits > uint32_t const *__restrict__ OuterHitOfCell< TrackerTraits > const int32_t	nHits
__shared__ uint32_t	ntot
int32_t	offset = isOuterHitOfCellWrap.offset
uint32_t const *__restrict__	offsets = hh.hitsLayerStart().data()
uint32_t	pairLayerId = 0
auto const &__restrict__	phiBinner = hh.phiBinner()
auto	stride = blockDim.x
float	x [maxCellsPerHit]
float	y [maxCellsPerHit]
float	z [maxCellsPerHit]
const float	z0cut = cuts.z0Cut_

Typedef Documentation

CellNeighbors

template<typename TrackerTraits >

using gpuPixelDoublets::CellNeighbors = typedef caStructures::CellNeighborsT<TrackerTraits>

Definition at line 20 of file gpuFishbone.h.

CellNeighborsVector

template<typename TrackerTraits >

using gpuPixelDoublets::CellNeighborsVector = typedef caStructures::CellNeighborsVectorT<TrackerTraits>

Definition at line 24 of file gpuFishbone.h.

CellTracks

template<typename TrackerTraits >

using gpuPixelDoublets::CellTracks = typedef caStructures::CellTracksT<TrackerTraits>

Definition at line 22 of file gpuFishbone.h.

CellTracksVector

template<typename TrackerTraits >

using gpuPixelDoublets::CellTracksVector = typedef caStructures::CellTracksVectorT<TrackerTraits>

Definition at line 26 of file gpuFishbone.h.

HitsConstView

template<typename TrackerTraits >

using gpuPixelDoublets::HitsConstView = typedef typename GPUCACellT<TrackerTraits>::HitsConstView

Definition at line 30 of file gpuFishbone.h.

OuterHitOfCell

template<typename TrackerTraits >

using gpuPixelDoublets::OuterHitOfCell = typedef caStructures::OuterHitOfCellT<TrackerTraits>

Definition at line 28 of file gpuFishbone.h.

PhiBinner

using gpuPixelDoublets::PhiBinner = typedef typename TrackingRecHitSoA<TrackerTraits>::PhiBinner

Definition at line 144 of file gpuPixelDoubletsAlgos.h.

Function Documentation

__attribute__() [1/2]

template<typename TrackerTraits >

gpuPixelDoublets::__attribute__ ( (always_inline)  )

inline

__attribute__() [2/2]

template<typename TrackerTraits >

__device__ gpuPixelDoublets::__attribute__ ( (always_inline)  )

inline

__syncthreads()

gpuPixelDoublets::__syncthreads ( )

inline

Definition at line 132 of file cudaCompat.h.

{}

assert()

gpuPixelDoublets::assert

(

offsets

)

Referenced by CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::buildDoublets(), gpuPixelDoublets::CellCutsT< TrackerTraits >::CellCutsT(), for(), and if().

for() [1/2]

gpuPixelDoublets::for

(

)

Definition at line 174 of file gpuPixelDoubletsAlgos.h.

References cms::cudacompat::__ldg(), funct::abs(), assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::atomicSub(), newFWLiteAna::bin, cellNeighbors, cells, cellTracks, cuts, HLT_2024v11_cff::doClusterCut, doPtCut, doZ0Cut, MillePedeFileConverter_cfg::e, first, full, hh, mps_fire::i, SurfaceOrientation::inner, innerLayerCumulativeSize, isOuterHitOfCell, dqmiolumiharvest::j, dqmdumpme::k, GetRecoTauVFromDQM_MC_cff::kk, pixelTopology::last_barrel_layer, TrajectorySeedProducer_cfi::layerPairs, gpuClustering::maxNumModules, maxNumOfDoublets, phase1PixelTopology::maxr, phase1PixelTopology::maxz, SiStripPI::min, minRadius2T4, phase1PixelTopology::minz, LaserClient_cfi::nbins, nCells, TrackingDataMCValidation_Standalone_cff::nmin, phase1PixelTopology::nPairs, offsets, SurfaceOrientation::outer, AlCaHLTBitMon_ParallelJobs::p, pairLayerId, phiBinner, submitPVValidationJobs::ptcut, short2phi(), stride, compareTotals::tot, and z0cut.

                                                              {
      while (j >= innerLayerCumulativeSize[pairLayerId++])
        ;
      --pairLayerId;  // move to lower_bound ??

      assert(pairLayerId < nPairs);
      assert(j < innerLayerCumulativeSize[pairLayerId]);
      assert(0 == pairLayerId || j >= innerLayerCumulativeSize[pairLayerId - 1]);

      uint8_t inner = TrackerTraits::layerPairs[2 * pairLayerId];
      uint8_t outer = TrackerTraits::layerPairs[2 * pairLayerId + 1];
      assert(outer > inner);

      auto hoff = PhiBinner::histOff(outer);
      auto i = (0 == pairLayerId) ? j : j - innerLayerCumulativeSize[pairLayerId - 1];
      i += offsets[inner];

      assert(i >= offsets[inner]);
      assert(i < offsets[inner + 1]);

      // found hit corresponding to our cuda thread, now do the job

      if (hh[i].detectorIndex() > gpuClustering::maxNumModules)
        continue;  // invalid

      /* maybe clever, not effective when zoCut is on
      auto bpos = (mi%8)/4;  // if barrel is 1 for z>0
      auto fpos = (outer>3) & (outer<7);
      if ( ((inner<3) & (outer>3)) && bpos!=fpos) continue;
      */

      auto mez = hh[i].zGlobal();

      if (mez < TrackerTraits::minz[pairLayerId] || mez > TrackerTraits::maxz[pairLayerId])
        continue;

      if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts.clusterCut(hh, i))
        continue;

      auto mep = hh[i].iphi();
      auto mer = hh[i].rGlobal();

      auto ptcut = [&](int j, int16_t idphi) {
        auto r2t4 = minRadius2T4;
        auto ri = mer;
        auto ro = hh[j].rGlobal();
        auto dphi = short2phi(idphi);
        return dphi * dphi * (r2t4 - ri * ro) > (ro - ri) * (ro - ri);
      };
      auto z0cutoff = [&](int j) {
        auto zo = hh[j].zGlobal();
        auto ro = hh[j].rGlobal();
        auto dr = ro - mer;
        return dr > TrackerTraits::maxr[pairLayerId] || dr < 0 || std::abs((mez * ro - mer * zo)) > z0cut * dr;
      };

      auto iphicut = cuts.phiCuts[pairLayerId];

      auto kl = PhiBinner::bin(int16_t(mep - iphicut));
      auto kh = PhiBinner::bin(int16_t(mep + iphicut));
      auto incr = [](auto& k) { return k = (k + 1) % PhiBinner::nbins(); };

#ifdef GPU_DEBUG
      int tot = 0;
      int nmin = 0;
      int tooMany = 0;
#endif

      auto khh = kh;
      incr(khh);
      for (auto kk = kl; kk != khh; incr(kk)) {
#ifdef GPU_DEBUG
        if (kk != kl && kk != kh)
          nmin += phiBinner.size(kk + hoff);
#endif
        auto const* __restrict__ p = phiBinner.begin(kk + hoff);
        auto const* __restrict__ e = phiBinner.end(kk + hoff);
        p += first;
        for (; p < e; p += stride) {
          auto oi = __ldg(p);
          assert(oi >= offsets[outer]);
          assert(oi < offsets[outer + 1]);
          auto mo = hh[oi].detectorIndex();

          if (mo > gpuClustering::maxNumModules)
            continue;  //    invalid

          if (doZ0Cut && z0cutoff(oi))
            continue;
          auto mop = hh[oi].iphi();
          uint16_t idphi = std::min(std::abs(int16_t(mop - mep)), std::abs(int16_t(mep - mop)));
          if (idphi > iphicut)
            continue;

          if (doClusterCut && cuts.zSizeCut(hh, i, oi))
            continue;
          if (doPtCut && ptcut(oi, idphi))
            continue;

          auto ind = atomicAdd(nCells, 1);
          if (ind >= maxNumOfDoublets) {
            atomicSub(nCells, 1);
            break;
          }  // move to SimpleVector??
          // int layerPairId, int doubletId, int innerHitId, int outerHitId)
          cells[ind].init(*cellNeighbors, *cellTracks, hh, pairLayerId, i, oi);
          isOuterHitOfCell[oi].push_back(ind);
#ifdef GPU_DEBUG
          if (isOuterHitOfCell[oi].full())
            ++tooMany;
          ++tot;
#endif
        }
      }
//      #endif
#ifdef GPU_DEBUG
      if (tooMany > 0 || tot > 0)
        printf("OuterHitOfCell for %d in layer %d/%d, %d,%d %d, %d %.3f %.3f %s\n",
               i,
               inner,
               outer,
               nmin,
               tot,
               tooMany,
               iphicut,
               TrackerTraits::minz[pairLayerId],
               TrackerTraits::maxz[pairLayerId],
               tooMany > 0 ? "FULL!!" : "not full.");
#endif
    }  // loop in block...

for() [2/2]

gpuPixelDoublets::for

(

int

idy = firstY

)

Definition at line 53 of file gpuFishbone.h.

References cms::cudacompat::blockDim, fftjetpileupestimator_calo_uncalib_cfi::c0, cc, cells, checkTrack, d, f, firstX, hh, idy, isOuterHitOfCell, l, n, x, y, and z.

                                                                                         {
      auto const& vc = isOuterHitOfCell[idy];
      auto size = vc.size();
      if (size < 2)
        continue;
      // if alligned kill one of the two.
      // in principle one could try to relax the cut (only in r-z?) for jumping-doublets
      auto const& c0 = cells[vc[0]];
      auto xo = c0.outer_x(hh);
      auto yo = c0.outer_y(hh);
      auto zo = c0.outer_z(hh);
      auto sg = 0;
      for (int32_t ic = 0; ic < size; ++ic) {
        auto& ci = cells[vc[ic]];
        if (ci.unused())
          continue;  // for triplets equivalent to next
        if (checkTrack && ci.tracks().empty())
          continue;
        cc[sg] = vc[ic];
        l[sg] = ci.layerPairId();
        d[sg] = ci.inner_detIndex(hh);
        x[sg] = ci.inner_x(hh) - xo;
        y[sg] = ci.inner_y(hh) - yo;
        z[sg] = ci.inner_z(hh) - zo;
        n[sg] = x[sg] * x[sg] + y[sg] * y[sg] + z[sg] * z[sg];
        ++sg;
      }
      if (sg < 2)
        continue;
      // here we parallelize
      for (int32_t ic = firstX; ic < sg - 1; ic += blockDim.x) {
        auto& ci = cells[cc[ic]];
        for (auto jc = ic + 1; jc < sg; ++jc) {
          auto& cj = cells[cc[jc]];
          // must be different detectors
          //        if (d[ic]==d[jc]) continue;
          auto cos12 = x[ic] * x[jc] + y[ic] * y[jc] + z[ic] * z[jc];
          if (d[ic] != d[jc] && cos12 * cos12 >= 0.99999f * (n[ic] * n[jc])) {
            // alligned:  kill farthest (prefer consecutive layers)
            // if same layer prefer farthest (longer level arm) and make space for intermediate hit
            bool sameLayer = l[ic] == l[jc];
            if (n[ic] > n[jc]) {
              if (sameLayer) {
                cj.kill();  // closest
                ci.setFishbone(cj.inner_hit_id(), cj.inner_z(hh), hh);
              } else {
                ci.kill();  // farthest
                // break;  // removed to improve reproducibility. keep it for reference and tests
              }
            } else {
              if (!sameLayer) {
                cj.kill();  // farthest
              } else {
                ci.kill();  // closest
                cj.setFishbone(ci.inner_hit_id(), ci.inner_z(hh), hh);
                // break;  // removed to improve reproducibility. keep it for reference    and tests
              }
            }
          }
        }  // cj
      }    // ci
    }      // hits

if() [1/2]

gpuPixelDoublets::if

(

0

= = first

)

Definition at line 22 of file gpuPixelDoublets.h.

References assert(), cellNeighbors, cellNeighborsContainer, cellTracks, cellTracksContainer, and mps_fire::i.

                    {
      cellNeighbors->construct(TrackerTraits::maxNumOfActiveDoublets, cellNeighborsContainer);
      cellTracks->construct(TrackerTraits::maxNumOfActiveDoublets, cellTracksContainer);
      auto i = cellNeighbors->extend();
      assert(0 == i);
      (*cellNeighbors)[0].reset();
      i = cellTracks->extend();
      assert(0 == i);
      (*cellTracks)[0].reset();
    }

if() [2/2]

gpuPixelDoublets::if

(

threadIdx.

y = = 0 && threadIdx.x == 0

)

Definition at line 158 of file gpuPixelDoubletsAlgos.h.

References mps_fire::i, innerLayerCumulativeSize, TrajectorySeedProducer_cfi::layerPairs, layerSize, phase1PixelTopology::nPairs, and ntot.

                                              {
      innerLayerCumulativeSize[0] = layerSize(TrackerTraits::layerPairs[0]);
      for (uint32_t i = 1; i < nPairs; ++i) {
        innerLayerCumulativeSize[i] = innerLayerCumulativeSize[i - 1] + layerSize(TrackerTraits::layerPairs[2 * i]);
      }
      ntot = innerLayerCumulativeSize[nPairs - 1];
    }

Variable Documentation

cc

uint32_t gpuPixelDoublets::cc[maxCellsPerHit]

Definition at line 49 of file gpuFishbone.h.

Referenced by BPHWriteSpecificDecay::addTrackModes(), HcalLutAnalyzer::analyze(), PPSPixelDigiAnalyzer::analyze(), AntiElectronIDMVA6< TauType, ElectronType >::AntiElectronIDMVA6(), PhotonEnergyCorrector::applyCrackCorrection(), EgammaSCEnergyCorrectionAlgo::applyCrackCorrection(), BeamHaloPropagatorESProducer::BeamHaloPropagatorESProducer(), l1t::stage2::BMTFTokens::BMTFTokens(), FFTJetPileupEstimator::calibrateFromDB(), CaloGeometryBuilder::CaloGeometryBuilder(), CaloGeometryDBEP< T, U >::CaloGeometryDBEP(), CaloGeometryEP< T, D >::CaloGeometryEP(), l1t::stage2::CaloLayer1Tokens::CaloLayer1Tokens(), CaloSimParametersESModule::CaloSimParametersESModule(), CaloSimulationConstantsESModule::CaloSimulationConstantsESModule(), l1t::stage2::CaloTokens::CaloTokens(), l1t::stage1::CaloTokens::CaloTokens(), CaloTowerConstituentsMapBuilder::CaloTowerConstituentsMapBuilder(), CaloTowerHardcodeGeometryEP::CaloTowerHardcodeGeometryEP(), CaloTPGTranscoderULUTs::CaloTPGTranscoderULUTs(), CaloTrkProcessingBuilder::CaloTrkProcessingBuilder(), CastorDbProducer::CastorDbProducer(), cert_plot(), FWGeometryTableManager::checkRegionOfInterest(), egamma::classBasedElectronEnergy(), ClusterShapeHitFilterESProducer::ClusterShapeHitFilterESProducer(), EcalUncalibRecHitTimingCCAlgo::computeCC(), TrajSeedMatcher::Configuration::Configuration(), hitfit::Fourvec_Constrainer::constrain(), ConvertingESProducerT< Record, Target, Source >::ConvertingESProducerT(), ConvertingESProducerWithDependenciesT< CombinedRecord< DepsRecords... >, Target, Dependencies... >::ConvertingESProducerWithDependenciesT(), p2eg::Cluster::crystaliEtaFromCardRegionInfo(), p2eg::Cluster::crystaliPhiFromCardRegionInfo(), p2eg::SimpleCaloHit::crystalLocaliEta(), p2eg::SimpleCaloHit::crystalLocaliPhi(), CSCConditions::CSCConditions(), CSCGeometryESModule::CSCGeometryESModule(), CTPPSBeamParametersFromLHCInfoESSource::CTPPSBeamParametersFromLHCInfoESSource(), CTPPSInterpolatedOpticalFunctionsESSource::CTPPSInterpolatedOpticalFunctionsESSource(), magneticfield::DD4hep_VolumeBasedMagneticFieldESProducerFromDB::DD4hep_VolumeBasedMagneticFieldESProducerFromDB(), GflashHadronShowerProfile::doCholeskyReduction(), DreamSensitiveDetectorBuilder::DreamSensitiveDetectorBuilder(), DTCombinatorialExtendedPatternReco::DTCombinatorialExtendedPatternReco(), DTCombinatorialPatternReco4D::DTCombinatorialPatternReco4D(), DTConfigDBProducer::DTConfigDBProducer(), DTFakeT0ESProducer::DTFakeT0ESProducer(), DTGeometryESModule::DTGeometryESModule(), DTGeometryESProducer::DTGeometryESProducer(), DTLinearDriftFromDBAlgo::DTLinearDriftFromDBAlgo(), DTMeantimerPatternReco::DTMeantimerPatternReco(), DTMeantimerPatternReco4D::DTMeantimerPatternReco4D(), DTRecHitBaseAlgo::DTRecHitBaseAlgo(), DTRefitAndCombineReco4D::DTRefitAndCombineReco4D(), dtCalibration::DTTTrigConstantShift::DTTTrigConstantShift(), dtCalibration::DTTTrigFillWithAverage::DTTTrigFillWithAverage(), dtCalibration::DTTTrigMatchRPhi::DTTTrigMatchRPhi(), dtCalibration::DTTTrigResidualCorrection::DTTTrigResidualCorrection(), DTTTrigSyncFromDB::DTTTrigSyncFromDB(), dtCalibration::DTTTrigT0SegCorrection::DTTTrigT0SegCorrection(), dtCalibration::DTVDriftSegment::DTVDriftSegment(), VVIObjDetails::dzero(), VVIObjFDetails::dzero(), sistripvvi::VVIObjDetails::dzero(), EcalBasicClusterLocalContCorrection::EcalBasicClusterLocalContCorrection(), ALPAKA_ACCELERATOR_NAMESPACE::EcalElectronicsMappingHostESProducer::EcalElectronicsMappingHostESProducer(), EcalLaserCorrectionService::EcalLaserCorrectionService(), EcalLaserCorrectionServiceMC::EcalLaserCorrectionServiceMC(), ALPAKA_ACCELERATOR_NAMESPACE::EcalMultifitConditionsHostESProducer::EcalMultifitConditionsHostESProducer(), EcalRegionCablingESProducer::EcalRegionCablingESProducer(), EcalSensitiveDetectorBuilder::EcalSensitiveDetectorBuilder(), EcalSignalGenerator< EEDigitizerTraits >::EcalSignalGenerator(), EcalSimParametersESModule::EcalSimParametersESModule(), EcalTBGeometryBuilder::EcalTBGeometryBuilder(), EcalTPGDBCopy::EcalTPGDBCopy(), EG8XObjectUpdateModifier::EG8XObjectUpdateModifier(), EgammaRegressionContainer::EgammaRegressionContainer(), EGEtScaleSysModifier::EGEtScaleSysModifier(), EGExtraInfoModifierFromValueMaps< MapType, OutputType >::EGExtraInfoModifierFromValueMaps(), EGFull5x5ShowerShapeModifierFromValueMaps::EGFull5x5ShowerShapeModifierFromValueMaps(), EGRegressionModifierCondTokens::EGRegressionModifierCondTokens(), EGRegressionModifierDRN::EGRegressionModifierDRN(), EGRegressionModifierV1::EGRegressionModifierV1(), EGRegressionModifierV2::EGRegressionModifierV2(), EGRegressionModifierV3::EGRegressionModifierV3(), ElectronHcalHelper::ElectronHcalHelper(), ElectronSeedGenerator::ElectronSeedGenerator(), l1t::stage2::EMTFTokens::EMTFTokens(), LumiCalculator::endRun(), ES_TTStubAlgorithm_cbc3< T >::ES_TTStubAlgorithm_cbc3(), ES_TTStubAlgorithm_official< T >::ES_TTStubAlgorithm_official(), helper::BFieldIsolationAlgorithmSetup< Alg >::esConsumes(), RegressionHelper::ESGetTokens::ESGetTokens(), EcalClusterLazyToolsBase::ESGetTokens::ESGetTokens(), ESRecHitWorker::ESRecHitWorker(), FFTJetCorrectionESProducer< CT >::FFTJetCorrectionESProducer(), FFTJetLookupTableESProducer< CT >::FFTJetLookupTableESProducer(), FiberSensitiveDetectorBuilder::FiberSensitiveDetectorBuilder(), BPHWriteSpecificDecay::fill(), TimingSD::fillHits(), TkAccumulatingSensitiveDetector::fillHits(), MuonSensitiveDetector::fillHits(), TotemSD::fillHits(), FP420SD::fillHits(), CaloSD::fillHits(), CaloSteppingAction::fillHits(), NtupleManager::FillMeasurements(), CaloSteppingAction::fillPassiveHits(), mkfit::MkBuilder::find_tracks_in_layers(), for(), FWRecoGeometryESProducer::FWRecoGeometryESProducer(), FWTGeoRecoGeometryESProducer::FWTGeoRecoGeometryESProducer(), SiStripGainESProducer::GainGetterT< Record >::GainGetterT(), GeneralTracksImporter::GeneralTracksImporter(), GenericSimClusterMapper::GenericSimClusterMapper(), p2eg::getCard_iEtaMax(), p2eg::getCard_iEtaMin(), p2eg::getCard_iPhiMax(), p2eg::getCard_iPhiMin(), p2eg::getCard_refCrystal_iEta(), p2eg::getCard_refCrystal_iPhi(), CaloGeometryHelper::getClosestCell(), GlobalDetLayerGeometryESProducer::GlobalDetLayerGeometryESProducer(), GlobalTrackingGeometryESProducer::GlobalTrackingGeometryESProducer(), l1t::stage2::GMTTokens::GMTTokens(), GsfTrajectoryFitterESProducer::GsfTrajectoryFitterESProducer(), GsfTrajectorySmootherESProducer::GsfTrajectorySmootherESProducer(), l1t::stage2::GTTokens::GTTokens(), HcalAlignmentEP::HcalAlignmentEP(), HcalDbProducer::HcalDbProducer(), HcalDDDGeometryEP::HcalDDDGeometryEP(), HcalDDDRecConstantsESModule::HcalDDDRecConstantsESModule(), HcalDDDSimConstantsESModule::HcalDDDSimConstantsESModule(), HcalHardcodeGeometryEP::HcalHardcodeGeometryEP(), HcalParametersESModule::HcalParametersESModule(), HcalSensitiveDetectorBuilder::HcalSensitiveDetectorBuilder(), HcalSimParametersESModule::HcalSimParametersESModule(), HcalSimulationConstantsESModule::HcalSimulationConstantsESModule(), HcalTB02ParametersESModule::HcalTB02ParametersESModule(), HcalTB02SensitiveDetectorBuilder::HcalTB02SensitiveDetectorBuilder(), HcalTB06BeamDetectorBuilder::HcalTB06BeamDetectorBuilder(), HcalTB06ParametersESModule::HcalTB06ParametersESModule(), HcalTPGCoderULUT::HcalTPGCoderULUT(), HFNoseSensitiveDetectorBuilder::HFNoseSensitiveDetectorBuilder(), HGCalGeometryESProducer::HGCalGeometryESProducer(), HGCalNumberingInitialization::HGCalNumberingInitialization(), HGCalParametersESModule::HGCalParametersESModule(), HGCalSensitiveDetectorBuilder::HGCalSensitiveDetectorBuilder(), HGCalTBGeometryESProducer::HGCalTBGeometryESProducer(), HGCalTBNumberingInitialization::HGCalTBNumberingInitialization(), HGCalTBParametersESModule::HGCalTBParametersESModule(), HGCalTriggerGeometryESProducer::HGCalTriggerGeometryESProducer(), HGCRecHitNavigator< D1, hgcee, D2, hgchef, D3, hgcheb >::HGCRecHitNavigator(), HGCScintillatorSensitiveDetectorBuilder::HGCScintillatorSensitiveDetectorBuilder(), HGCSensitiveDetectorBuilder::HGCSensitiveDetectorBuilder(), HLTDQMTagAndProbeEff< TagType, TagCollType, ProbeType, ProbeCollType >::HLTDQMTagAndProbeEff(), HLTExclDiJetFilter< T >::hltFilter(), HLTTauDQML1Plotter::HLTTauDQML1Plotter(), p2eg::SimpleCaloHit::isInCard(), L1EGCrystalClusterEmulatorProducer::SimpleCaloHit::isInCard(), L1EGCrystalClusterEmulatorProducer::SimpleCaloHit::isInCardAndRegion(), p2eg::isValidCard(), ConvertingESProducerWithDependenciesT< CombinedRecord< DepsRecords... >, Target, Dependencies... >::WalkConsumes< N >::iterate(), L1CondDBPayloadWriterExt::L1CondDBPayloadWriterExt(), L1ConfigOnlineProdBase< L1HtMissScaleRcd, L1CaloEtScale >::L1ConfigOnlineProdBase(), L1ObjectKeysOnlineProdBase::L1ObjectKeysOnlineProdBase(), l1t::L1TDigiToRaw::L1TDigiToRaw(), L1TriggerKeyOnlineProd::L1TriggerKeyOnlineProd(), L1TriggerKeyOnlineProdExt::L1TriggerKeyOnlineProdExt(), MagneticFieldMapESProducer::MagneticFieldMapESProducer(), SiStripGainESProducer::make_GainGetter(), calogeometryDBEPimpl::GeometryTraits< HGCalGeometry, true >::makeToken(), calogeometryDBEPimpl::GeometryTraits< T, U::writeFlag()>::makeToken(), calogeometryDBEPimpl::GeometryTraits< T, false >::makeToken(), calogeometryDBEPimpl::AdditionalTokens< CaloTowerGeometry >::makeTokens(), calogeometryDBEPimpl::AdditionalTokens< HcalGeometry >::makeTokens(), calogeometryDBEPimpl::AdditionalTokens< HGCalGeometry >::makeTokens(), ME0GeometryESModule::ME0GeometryESModule(), MisalignedTrackerESProducer::MisalignedTrackerESProducer(), MkFitGeometryESProducer::MkFitGeometryESProducer(), MTDDetLayerGeometryESProducer::MTDDetLayerGeometryESProducer(), MTDDigiGeometryESModule::MTDDigiGeometryESModule(), MTDGeometricTimingDetESModule::MTDGeometricTimingDetESModule(), MTDParametersESModule::MTDParametersESModule(), MTDTimeCalibESProducer::MTDTimeCalibESProducer(), MultipleScatteringParametrisationMakerESProducer::MultipleScatteringParametrisationMakerESProducer(), MultiRecHitCollectorESProducer::MultiRecHitCollectorESProducer(), MuonDetLayerGeometryESProducer::MuonDetLayerGeometryESProducer(), MuonGeometryConstantsESModule::MuonGeometryConstantsESModule(), MuonOffsetESProducer::MuonOffsetESProducer(), MuonSegmentCompatibilityCut::MuonSegmentCompatibilityCut(), MuonTrackCut::MuonTrackCut(), MVAVariableHelper::MVAVariableHelper(), pat::helper::NamedUserDataLoader< pat::helper::AddUserCand >::NamedUserDataLoader(), NavigationSchoolESProducer::NavigationSchoolESProducer(), pat::ObjectModifier< T >::ObjectModifier(), OfflineToTransientBeamSpotESProducer::OfflineToTransientBeamSpotESProducer(), TwoTrackMinimumDistanceHelixLine::oneIteration(), pat::PATMETSlimmer::OneMETShift::OneMETShift(), OnlineBeamSpotESProducer::OnlineBeamSpotESProducer(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()(), EGRegressionModifierDRN::partVars< T >::partVars(), PFClusterBuilderBase::PFClusterBuilderBase(), PFClusterEMEnergyCorrector::PFClusterEMEnergyCorrector(), PFClusterFromHGCalTrackster::PFClusterFromHGCalTrackster(), PFClusterProducer::PFClusterProducer(), PFEcalBarrelRecHitCreator::PFEcalBarrelRecHitCreator(), PFEcalEndcapRecHitCreator::PFEcalEndcapRecHitCreator(), PFMultiDepthClusterizer::PFMultiDepthClusterizer(), PFMultiDepthClusterProducer::PFMultiDepthClusterProducer(), PFRecHitCreatorBase::PFRecHitCreatorBase(), PFRecHitDualNavigator< D1, barrel, D2, endcap >::PFRecHitDualNavigator(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitECALParamsESProducer::PFRecHitECALParamsESProducer(), PFRecHitProducer::PFRecHitProducer(), PFRecHitQTestHCALThresholdVsDepth::PFRecHitQTestHCALThresholdVsDepth(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitTopologyESProducer< CAL >::PFRecHitTopologyESProducer(), PixelCPEFastESProducerT< TrackerTraits >::PixelCPEFastESProducerT(), ALPAKA_ACCELERATOR_NAMESPACE::PixelCPEFastParamsESProducerAlpaka< TrackerTraits >::PixelCPEFastParamsESProducerAlpaka(), PlotDebugFPIX_XYMap(), PlotPixelMultVtxPos(), FastTimerService::postESModule(), FastTimerService::preESModule(), DDEcalEndcapTrapX::print(), DDEcalEndcapTrap::print(), p2eg::Cluster::printClusterInfo(), printPlot(), PixelTracksProducer::produce(), CastorClusterProducer::produce(), Phase2L1CaloEGammaEmulator::produce(), JetCoreMCtruthSeedGenerator::produce(), DeepCoreSeedGenerator::produce(), L1EGCrystalClusterEmulatorProducer::produce(), trklet::ProducerChannelAssignment::ProducerChannelAssignment(), trackerTFP::ProducerDemonstrator::ProducerDemonstrator(), trackerTFP::ProducerES::ProducerES(), trackerTFP::ProducerFormatsKF::ProducerFormatsKF(), hph::ProducerHPH::ProducerHPH(), trackerTFP::ProducerLayerEncoding::ProducerLayerEncoding(), trackerDTC::ProducerLayerEncoding::ProducerLayerEncoding(), tt::ProducerSetup::ProducerSetup(), p2eg::Cluster::realEta(), p2eg::Cluster::realPhi(), MultiHitGeneratorFromChi2::refit2Hits(), HITrackingRegionProducer::regions(), HITrackingRegionForPrimaryVtxProducer::regions(), l1t::stage2::CaloSetup::registerConsumes(), l1t::stage2::CaloLayer1Setup::registerConsumes(), l1t::stage2::GMTSetup::registerConsumes(), l1t::stage2::EMTFSetup::registerConsumes(), l1t::stage1::CaloSetup::registerConsumes(), l1t::stage2::BMTFSetup::registerConsumes(), l1t::stage2::GTSetup::registerConsumes(), MonopoleSteppingAction::registerConsumes(), MaterialBudgetHcal::registerConsumes(), PrintGeomSummary::registerConsumes(), SimG4HGCalValidation::registerConsumes(), HcalTestAnalysis::registerConsumes(), PrintGeomInfoAction::registerConsumes(), MaterialBudgetHcalProducer::registerConsumes(), SimG4HcalValidation::registerConsumes(), CaloSteppingAction::registerConsumes(), SCEnergyCorrectorSemiParm::RegParam::RegParam(), RegressionHelper::RegressionHelper(), lumi::CMSRunSummary2DB::retrieveData(), SiPixelDQMRocLevelAnalyzer::RocSumOneModule(), DDEcalEndcapTrap::rotate(), DDEcalEndcapTrapX::rotate(), RPCConeBuilder::RPCConeBuilder(), RPCInverseCPPFLinkMapESProducer::RPCInverseCPPFLinkMapESProducer(), RPCInverseLBLinkMapESProducer::RPCInverseLBLinkMapESProducer(), RPCInverseOMTFLinkMapESProducer::RPCInverseOMTFLinkMapESProducer(), RPCInverseTwinMuxLinkMapESProducer::RPCInverseTwinMuxLinkMapESProducer(), DTTracoChip::run(), SCEnergyCorrectorDRN::SCEnergyCorrectorDRN(), SCEnergyCorrectorSemiParm::SCEnergyCorrectorSemiParm(), sim::sensitiveDetectorMakers(), mkfit::TrackCand::setCombCandidate(), PhoMVACut::setConsumes(), PhoFull5x5SigmaIEtaIEtaValueMapCut::setConsumes(), PhoAnyPFIsoWithEAAndQuadScalingCut::setConsumes(), GsfEleRelPFIsoScaledCut::setConsumes(), PhoAnyPFIsoWithEACut::setConsumes(), GsfEleMVACut::setConsumes(), PhoAnyPFIsoWithEAAndExpoScalingEBCut::setConsumes(), GsfEleEffAreaPFIsoCut::setConsumes(), PhoAnyPFIsoWithEAAndExpoScalingCut::setConsumes(), GsfEleValueMapIsoRhoCut::setConsumes(), GsfEleEmHadD1IsoRhoCut::setConsumes(), GsfEleDxyCut::setConsumes(), GsfEleDzCut::setConsumes(), GsfEleCalPFClusterIsoCut::setConsumes(), PhoGenericQuadraticRhoPtScaledCut::setConsumes(), GsfEleTrkPtIsoRhoCut::setConsumes(), MuonPOGStandardCut::setConsumes(), PhoGenericRhoPtScaledCut::setConsumes(), GsfEleConversionVetoCut::setConsumes(), MuonDxyCut::setConsumes(), MuonDzCut::setConsumes(), GsfEleHadronicOverEMEnergyScaledCut::setConsumes(), EcalTPGDBCopy::setConsumes(), HcalDbProducer::ServiceTokenImpl< ProductType, RecordType, LABEL, EffectiveType >::setConsumes(), HcalDbProducer::TokensForServiceHolder< RecordType, TokenHolders >::setConsumes(), HcalDbProducer::TokenAndTopologyHolder< ProductT, RecordT >::setConsumes(), VersionedSelector< T >::setConsumes(), PFBlockAlgo::setImporters(), l1t::WriterProxyT< Record, Type >::setToken(), EgammaHLTPhase2ExtraProducer::Tokens::setToken(), EgammaHLTExtraProducer::Tokens::setToken(), SCEnergyCorrectorSemiParm::setTokens(), SCEnergyCorrectorDRN::setTokens(), SCEnergyCorrectorSemiParm::RegParam::setTokens(), PFECALSuperClusterAlgo::setTokens(), SiPhase2BadStripConfigurableFakeESSource::SiPhase2BadStripConfigurableFakeESSource(), SiPhase2OuterTrackerFakeLorentzAngleESSource::SiPhase2OuterTrackerFakeLorentzAngleESSource(), SiPhase2RecHitMatcherESProducer::SiPhase2RecHitMatcherESProducer(), SiPixel2DTemplateDBObjectESProducer::SiPixel2DTemplateDBObjectESProducer(), ALPAKA_ACCELERATOR_NAMESPACE::SiPixelCablingSoAESProducer::SiPixelCablingSoAESProducer(), SiPixelGainCalibrationForHLTGPUESProducer::SiPixelGainCalibrationForHLTGPUESProducer(), ALPAKA_ACCELERATOR_NAMESPACE::SiPixelGainCalibrationForHLTSoAESProducer::SiPixelGainCalibrationForHLTSoAESProducer(), SiPixelGenErrorDBObjectESProducer::SiPixelGenErrorDBObjectESProducer(), SiPixelROCsStatusAndMappingWrapperESProducer::SiPixelROCsStatusAndMappingWrapperESProducer(), SiPixelTemplateDBObjectESProducer::SiPixelTemplateDBObjectESProducer(), SiStripBackPlaneCorrectionFakeESSource::SiStripBackPlaneCorrectionFakeESSource(), SiStripClusterizerConditionsESProducer::SiStripClusterizerConditionsESProducer(), SiStripClusterizerConditionsGPUESProducer::SiStripClusterizerConditionsGPUESProducer(), SiStripDelayESProducer::SiStripDelayESProducer(), SiStripLorentzAngleFakeESSource::SiStripLorentzAngleFakeESSource(), SiStripRegionConnectivity::SiStripRegionConnectivity(), SiTrackerMultiRecHitUpdatorESProducer::SiTrackerMultiRecHitUpdatorESProducer(), SkippingLayerCosmicNavigationSchoolESProducer::SkippingLayerCosmicNavigationSchoolESProducer(), SmartPropagatorESProducer::SmartPropagatorESProducer(), p2eg::stitchClusterOverRegionBoundary(), LRHelpFunctions::storeControlPlots(), storeTracks(), StripCompletePlot(), StripCPEESProducer::StripCPEESProducer(), SuperClusterImporter::SuperClusterImporter(), heppy::mt2w_bisect::mt2w::teco(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerA::TestAlpakaESProducerA(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerB::TestAlpakaESProducerB(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerC::TestAlpakaESProducerC(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerD::TestAlpakaESProducerD(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerE::TestAlpakaESProducerE(), edmtest::ThinningThingSelector::ThinningThingSelector(), TkDetMapESProducer::TkDetMapESProducer(), TkMSParameterizationBuilder::TkMSParameterizationBuilder(), SiStripGainSimESProducer::TokenLabel::TokenLabel(), CandidateBoostedDoubleSecondaryVertexComputer::Tokens::Tokens(), CombinedMVAV2JetTagComputer::Tokens::Tokens(), CandidateChargeBTagComputer::Tokens::Tokens(), MuonTagger::Tokens::Tokens(), ElectronTagger::Tokens::Tokens(), CharmTagger::Tokens::Tokens(), HeavyIonCSVTagger::Tokens::Tokens(), EgammaHLTPhase2ExtraProducer::Tokens::Tokens(), EgammaHLTExtraProducer::Tokens::Tokens(), SiStripDelayESProducer::TokenSign::TokenSign(), TotemGeometryESModule::TotemGeometryESModule(), TPNCor::TPNCor(), TrackerDigiGeometryESModule::TrackerDigiGeometryESModule(), TrackerGeometricDetESModule::TrackerGeometricDetESModule(), TrackerInteractionGeometryESProducer::TrackerInteractionGeometryESProducer(), TrackerParametersESModule::TrackerParametersESModule(), pflow::importers::TrackFromParentImporter< Collection, Adaptor >::TrackFromParentImporter(), TransientTrackBuilderESProducer::TransientTrackBuilderESProducer(), TSCBLBuilderWithPropagatorESProducer::TSCBLBuilderWithPropagatorESProducer(), TwoTrackMinimumDistanceHelixLine::updateCoeffs(), magneticfield::VolumeBasedMagneticFieldESProducerFromDB::VolumeBasedMagneticFieldESProducerFromDB(), BPHWriteSpecificDecay::write(), WriteESAlignments::WriteESAlignments(), and XMLIdealMagneticFieldGeometryESProducer::XMLIdealMagneticFieldGeometryESProducer().

cellNeighbors

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > * gpuPixelDoublets::cellNeighbors

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by for(), and if().

cellNeighborsContainer

int CellNeighborsVector<TrackerTraits> CellNeighbors<TrackerTraits>* gpuPixelDoublets::cellNeighborsContainer

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by if(), and ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::InitDoublets< TrackerTraits >::operator()().

cells

__device__ uint32_t GPUCACellT< TrackerTraits > * gpuPixelDoublets::cells

Definition at line 34 of file gpuFishbone.h.

Referenced by for().

cellTracks

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > * gpuPixelDoublets::cellTracks

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by for(), and if().

cellTracksContainer

int CellNeighborsVector<TrackerTraits> CellNeighbors<TrackerTraits> CellTracksVector<TrackerTraits> CellTracks<TrackerTraits>* gpuPixelDoublets::cellTracksContainer

Initial value:

{
    assert(isOuterHitOfCell.container)

Definition at line 16 of file gpuPixelDoublets.h.

Referenced by if(), and ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::InitDoublets< TrackerTraits >::operator()().

checkTrack

GPUCACellT<TrackerTraits> uint32_t const* __restrict__ OuterHitOfCell<TrackerTraits> const int32_t bool gpuPixelDoublets::checkTrack

Initial value:

{
    constexpr auto maxCellsPerHit = GPUCACellT<TrackerTraits>::maxCellsPerHit

Definition at line 38 of file gpuFishbone.h.

Referenced by for(), and ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()().

cuts

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > CellCutsT< TrackerTraits > const & gpuPixelDoublets::cuts

Initial value:

{

    doubletsFromHisto<TrackerTraits>(
        nActualPairs, maxNumOfDoublets, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, *cuts)

Definition at line 50 of file gpuPixelDoublets.h.

Referenced by for().

d

uint16_t gpuPixelDoublets::d[maxCellsPerHit]

Definition at line 50 of file gpuFishbone.h.

Referenced by for().

doPtCut

const bool gpuPixelDoublets::doPtCut = cuts.doPtCut_

Definition at line 136 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

doZ0Cut

const bool gpuPixelDoublets::doZ0Cut = cuts.doZ0Cut_

Definition at line 135 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

first

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

Definition at line 18 of file gpuPixelDoublets.h.

Referenced by for().

firstX

auto gpuPixelDoublets::firstX = threadIdx.x

Definition at line 46 of file gpuFishbone.h.

Referenced by for(), NoisyChannel::getAverage2D(), and PixelTemplateSmearerBase::smearHit().

firstY

auto gpuPixelDoublets::firstY = threadIdx.y + blockIdx.y * blockDim.y

Definition at line 45 of file gpuFishbone.h.

Referenced by NoisyChannel::getAverage2D(), and PixelTemplateSmearerBase::smearHit().

getDoubletsFromHistoMaxBlockSize

constexpr auto gpuPixelDoublets::getDoubletsFromHistoMaxBlockSize = 64

Definition at line 34 of file gpuPixelDoublets.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::buildDoublets().

getDoubletsFromHistoMinBlocksPerMP

constexpr auto gpuPixelDoublets::getDoubletsFromHistoMinBlocksPerMP = 16

Definition at line 35 of file gpuPixelDoublets.h.

hardPtCut

const float gpuPixelDoublets::hardPtCut = cuts.ptCut_

Definition at line 139 of file gpuPixelDoubletsAlgos.h.

hh

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > gpuPixelDoublets::hh

Definition at line 43 of file gpuPixelDoublets.h.

Referenced by CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::buildDoublets(), and for().

idy

auto gpuPixelDoublets::idy = blockIdx.y * blockDim.y + threadIdx.y

Definition at line 168 of file gpuPixelDoubletsAlgos.h.

Referenced by for(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()(), and TTUTrackingAlg::runSeedBuster().

innerLayerCumulativeSize

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

Definition at line 156 of file gpuPixelDoubletsAlgos.h.

Referenced by for(), and if().

isOuterHitOfCell

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > OuterHitOfCell< TrackerTraits > gpuPixelDoublets::isOuterHitOfCell = isOuterHitOfCellWrap.container

Definition at line 41 of file gpuFishbone.h.

Referenced by for().

isOuterHitOfCellWrap

GPUCACellT<TrackerTraits> uint32_t const* __restrict__ OuterHitOfCell<TrackerTraits> const gpuPixelDoublets::isOuterHitOfCellWrap

Definition at line 34 of file gpuFishbone.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()().

l

uint8_t gpuPixelDoublets::l[maxCellsPerHit]

Definition at line 51 of file gpuFishbone.h.

Referenced by for().

layerSize

auto gpuPixelDoublets::layerSize = [=](uint8_t li) { return offsets[li + 1] - offsets[li]; }

Definition at line 150 of file gpuPixelDoubletsAlgos.h.

Referenced by if().

maxNumOfDoublets

__device__ uint32_t gpuPixelDoublets::maxNumOfDoublets

Definition at line 43 of file gpuPixelDoublets.h.

Referenced by for().

minRadius

const float gpuPixelDoublets::minRadius = hardPtCut * 87.78f

Definition at line 141 of file gpuPixelDoubletsAlgos.h.

minRadius2T4

const float gpuPixelDoublets::minRadius2T4 = 4.f * minRadius * minRadius

Definition at line 142 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

n

float gpuPixelDoublets::n[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

nActualPairs

uint32_t CellNeighborsVector<TrackerTraits> CellTracksVector<TrackerTraits> HitsConstView<TrackerTraits> OuterHitOfCell<TrackerTraits> int gpuPixelDoublets::nActualPairs

Definition at line 43 of file gpuPixelDoublets.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::buildDoublets(), CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::buildDoublets(), caHitNtupletGenerator::ParamsT< TrackerTraits, pixelTopology::isPhase1Topology< TrackerTraits > >::nPairs(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGenerator::ParamsT< TrackerTraits, pixelTopology::isPhase1Topology< TrackerTraits > >::nPairs(), caHitNtupletGenerator::ParamsT< TrackerTraits, pixelTopology::isPhase2Topology< TrackerTraits > >::nPairs(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGenerator::ParamsT< TrackerTraits, pixelTopology::isPhase2Topology< TrackerTraits > >::nPairs(), and ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::GetDoubletsFromHisto< TrackerTraits >::operator()().

nCells

__device__ uint32_t GPUCACellT< TrackerTraits > uint32_t * gpuPixelDoublets::nCells

Definition at line 34 of file gpuFishbone.h.

Referenced by for().

nHits

int gpuPixelDoublets::nHits

Definition at line 34 of file gpuFishbone.h.

ntot

__shared__ uint32_t gpuPixelDoublets::ntot

Definition at line 157 of file gpuPixelDoubletsAlgos.h.

Referenced by if().

offset

int32_t gpuPixelDoublets::offset = isOuterHitOfCellWrap.offset

Definition at line 42 of file gpuFishbone.h.

offsets

uint32_t const* __restrict__ gpuPixelDoublets::offsets = hh.hitsLayerStart().data()

Definition at line 147 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

pairLayerId

uint32_t gpuPixelDoublets::pairLayerId = 0

Definition at line 172 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

phiBinner

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

Definition at line 146 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

stride

auto gpuPixelDoublets::stride = blockDim.x

Definition at line 170 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

x

float gpuPixelDoublets::x[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

y

float gpuPixelDoublets::y[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

z

float gpuPixelDoublets::z[maxCellsPerHit]

Definition at line 48 of file gpuFishbone.h.

Referenced by for().

z0cut

const float gpuPixelDoublets::z0cut = cuts.z0Cut_

Definition at line 138 of file gpuPixelDoubletsAlgos.h.

Referenced by for().