gpuPixelDoublets Namespace Reference

Typedefs
using	CellNeighbors = caConstants::CellNeighbors
using	CellNeighborsVector = caConstants::CellNeighborsVector
using	CellTracks = caConstants::CellTracks
using	CellTracksVector = caConstants::CellTracksVector
using	PhiBinner = TrackingRecHit2DSOAView::PhiBinner

Functions
	__attribute__ ((always_inline)) void fishbone(GPUCACell
	__syncthreads ()
	assert (offsets)
	assert (nPairs<=nPairsMax)
	doubletsFromHisto (layerPairs, nActualPairs, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, phicuts, minz, maxz, maxr, ideal_cond, doClusterCut, doZ0Cut, doPtCut, maxNumOfDoublets)
	for (auto j=idy;j< ntot;j+=blockDim.y *gridDim.y)
	if (threadIdx.y==0 &&threadIdx.x==0)

Variables
uint32_t CellNeighborsVector *	cellNeighbors
__device__ uint32_t GPUCACell *	cells
uint32_t CellNeighborsVector CellTracksVector *	cellTracks
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int bool bool	doClusterCut
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int bool bool bool bool	doPtCut
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int bool bool bool	doZ0Cut
constexpr float	dzdrFact = 8 * 0.0285 / 0.015
auto	first = threadIdx.x
constexpr auto	getDoubletsFromHistoMaxBlockSize = 64
constexpr auto	getDoubletsFromHistoMinBlocksPerMP = 16
__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__	hh
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__	hhp
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int bool	ideal_cond
auto	idy = blockIdx.y * blockDim.y + threadIdx.y
__shared__ uint32_t	innerLayerCumulativeSize [nPairsMax]
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell	isOuterHitOfCell
bool	isOuterLadder = ideal_cond
__constant__ const uint8_t	layerPairs [2 *nPairs]
auto	layerSize = [=](uint8_t li) { return offsets[li + 1] - offsets[li]; }
constexpr int	maxDYPred = 20
constexpr int	maxDYsize = 20
constexpr int	maxDYsize12 = 28
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int bool bool bool bool uint32_t	maxNumOfDoublets
__constant__ float const	maxr [nPairs]
__constant__ float const	maxz [nPairs]
constexpr int	minYsizeB2 = 28
__constant__ float const	minz [nPairs]
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int	nActualPairs
uint32_t *	nCells
constexpr int	nPairs = nPairsForTriplets + 4
constexpr int	nPairsForQuadruplets = 13
constexpr int	nPairsForTriplets = nPairsForQuadruplets + 2
const int	nPairsMax = caConstants::maxNumberOfLayerPairs
__shared__ uint32_t	ntot
uint32_t const *__restrict__	offsets = hh.hitsLayerStart()
uint32_t	pairLayerId = 0
constexpr int16_t	phi0p05 = 522
constexpr int16_t	phi0p06 = 626
constexpr int16_t	phi0p07 = 730
auto const &__restrict__	phiBinner = hh.phiBinner()
__constant__ const int16_t	phicuts [nPairs]
auto	stride = blockDim.x

Typedef Documentation

typedef caConstants::CellNeighbors gpuPixelDoublets::CellNeighbors

Definition at line 63 of file gpuPixelDoublets.h.

typedef caConstants::CellNeighborsVector gpuPixelDoublets::CellNeighborsVector

Definition at line 65 of file gpuPixelDoublets.h.

typedef caConstants::CellTracks gpuPixelDoublets::CellTracks

Definition at line 64 of file gpuPixelDoublets.h.

typedef caConstants::CellTracksVector gpuPixelDoublets::CellTracksVector

Definition at line 66 of file gpuPixelDoublets.h.

using gpuPixelDoublets::PhiBinner = typedef TrackingRecHit2DSOAView::PhiBinner

Definition at line 53 of file gpuPixelDoubletsAlgos.h.

Function Documentation

__device__ gpuPixelDoublets::__attribute__ ( (always_inline)  )

inline

Definition at line 19 of file gpuFishbone.h.

References cms::cudacompat::blockDim, cms::cudacompat::blockIdx, caConstants::OuterHitOfCell::container, ztail::d, validate-o2o-wbm::f, cms::cudacompat::gridDim, hh, hhp, idy, isOuterHitOfCell, cmsLHEtoEOSManager::l, GPUCACell::layerPairId, GPUCACell::maxCellsPerHit, caConstants::maxCellsPerHit, dqmiodumpmetadata::n, nt, hltrates_dqm_sourceclient-live_cfg::offset, caConstants::OuterHitOfCell::offset, findQualityFiles::size, and cms::cudacompat::threadIdx.

                                            {
    constexpr auto maxCellsPerHit = GPUCACell::maxCellsPerHit;

    auto const& hh = *hhp;

    auto const isOuterHitOfCell = isOuterHitOfCellWrap.container;
    int32_t offset = isOuterHitOfCellWrap.offset;

    // x runs faster...
    auto firstY = threadIdx.y + blockIdx.y * blockDim.y;
    auto firstX = threadIdx.x;

    float x[maxCellsPerHit], y[maxCellsPerHit], z[maxCellsPerHit], n[maxCellsPerHit];
    uint32_t cc[maxCellsPerHit];
    uint16_t d[maxCellsPerHit];
    uint8_t l[maxCellsPerHit];

    for (int idy = firstY, nt = nHits - offset; idy < nt; idy += gridDim.y * blockDim.y) {
      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());
              } else {
                ci.kill();  // farthest
                break;
              }
            } else {
              if (!sameLayer) {
                cj.kill();  // farthest
              } else {
                ci.kill();  // closest
                cj.setFishbone(ci.inner_hit_id());
                break;
              }
            }
          }
        }  // cj
      }    // ci
    }      // hits
  }

gpuPixelDoublets::__syncthreads ( )

inline

Definition at line 108 of file cudaCompat.h.

{}

gpuPixelDoublets::assert

(

offsets

)

Referenced by for().

gpuPixelDoublets::assert

(

nPairs<=

nPairsMax

)

gpuPixelDoublets::doubletsFromHisto	(	layerPairs	,
		nActualPairs	,
		cells	,
		nCells	,
		cellNeighbors	,
		cellTracks	,
		hh	,
		isOuterHitOfCell	,
		phicuts	,
		minz	,
		maxz	,
		maxr	,
		ideal_cond	,
		doClusterCut	,
		doZ0Cut	,
		doPtCut	,
		maxNumOfDoublets
	)

gpuPixelDoublets::for

(

)

Definition at line 83 of file gpuPixelDoubletsAlgos.h.

References cms::cudacompat::__ldg(), funct::abs(), assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::atomicSub(), cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::bin(), runTauDisplay::dr, PVValHelper::dy, alignCSCRings::e, validate-o2o-wbm::f, first, full, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::histOff(), mps_fire::i, SurfaceOrientation::inner, dqmiolumiharvest::j, isotrackApplyRegressor::k, GetRecoTauVFromDQM_MC_cff::kk, maxDYPred, maxDYsize, gpuClustering::maxNumModules, min(), cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins(), nmin, SurfaceOrientation::outer, AlCaHLTBitMon_ParallelJobs::p, pairLayerId, gpuVertexFinder::printf(), submitPVValidationJobs::ptcut, short2phi(), and stride.

                                                              {
      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 = layerPairs[2 * pairLayerId];
      uint8_t outer = layerPairs[2 * pairLayerId + 1];
      assert(outer > inner);

      auto hoff = PhiBinner::histOff(outer);

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

      // printf("Hit in Layer %d %d %d %d\n", i, inner, pairLayerId, j);

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

      // found hit corresponding to our cuda thread, now do the job
      auto mi = hh.detectorIndex(i);
      if (mi > 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.zGlobal(i);

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

      int16_t mes = -1;  // make compiler happy
      if (doClusterCut) {
        // if ideal treat inner ladder as outer
        if (inner == 0)
          assert(mi < 96);
        isOuterLadder = ideal_cond ? true : 0 == (mi / 8) % 2;  // only for B1/B2/B3 B4 is opposite, FPIX:noclue...

        // in any case we always test mes>0 ...
        mes = inner > 0 || isOuterLadder ? hh.clusterSizeY(i) : -1;

        if (inner == 0 && outer > 3)  // B1 and F1
          if (mes > 0 && mes < minYsizeB1)
            continue;                 // only long cluster  (5*8)
        if (inner == 1 && outer > 3)  // B2 and F1
          if (mes > 0 && mes < minYsizeB2)
            continue;
      }
      auto mep = hh.iphi(i);
      auto mer = hh.rGlobal(i);

      // all cuts: true if fails
      constexpr float z0cut = 12.f;      // cm
      constexpr float hardPtCut = 0.5f;  // GeV
      // cm (1 GeV track has 1 GeV/c / (e * 3.8T) ~ 87 cm radius in a 3.8T field)
      constexpr float minRadius = hardPtCut * 87.78f;
      constexpr float minRadius2T4 = 4.f * minRadius * minRadius;
      auto ptcut = [&](int j, int16_t idphi) {
        auto r2t4 = minRadius2T4;
        auto ri = mer;
        auto ro = hh.rGlobal(j);
        auto dphi = short2phi(idphi);
        return dphi * dphi * (r2t4 - ri * ro) > (ro - ri) * (ro - ri);
      };
      auto z0cutoff = [&](int j) {
        auto zo = hh.zGlobal(j);
        auto ro = hh.rGlobal(j);
        auto dr = ro - mer;
        return dr > maxr[pairLayerId] || dr < 0 || std::abs((mez * ro - mer * zo)) > z0cut * dr;
      };

      auto zsizeCut = [&](int j) {
        auto onlyBarrel = outer < 4;
        auto so = hh.clusterSizeY(j);
        auto dy = inner == 0 ? maxDYsize12 : maxDYsize;
        // in the barrel cut on difference in size
        // in the endcap on the prediction on the first layer (actually in the barrel only: happen to be safe for endcap as well)
        // FIXME move pred cut to z0cutoff to optmize loading of and computaiton ...
        auto zo = hh.zGlobal(j);
        auto ro = hh.rGlobal(j);
        return onlyBarrel ? mes > 0 && so > 0 && std::abs(so - mes) > dy
                          : (inner < 4) && mes > 0 &&
                                std::abs(mes - int(std::abs((mez - zo) / (mer - ro)) * dzdrFact + 0.5f)) > maxDYPred;
      };

      auto iphicut = 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.detectorIndex(oi);
          if (mo > gpuClustering::maxNumModules)
            continue;  //    invalid

          if (doZ0Cut && z0cutoff(oi))
            continue;

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

          if (doClusterCut && zsizeCut(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
        }
      }
#ifdef GPU_DEBUG
      if (tooMany > 0)
        printf("OuterHitOfCell full for %d in layer %d/%d, %d,%d %d\n", i, inner, outer, nmin, tot, tooMany);
#endif
    }  // loop in block...

gpuPixelDoublets::if

(

threadIdx.

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

)

Definition at line 68 of file gpuPixelDoubletsAlgos.h.

References mps_fire::i, layerSize, and nPairs.

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

Variable Documentation

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector * gpuPixelDoublets::cellNeighbors

Definition at line 99 of file gpuPixelDoublets.h.

__device__ uint32_t GPUCACell* gpuPixelDoublets::cells

Definition at line 26 of file gpuPixelDoubletsAlgos.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector * gpuPixelDoublets::cellTracks

Definition at line 99 of file gpuPixelDoublets.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ float const *__restrict__ bool bool gpuPixelDoublets::doClusterCut

Definition at line 99 of file gpuPixelDoublets.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ float const *__restrict__ bool bool bool bool gpuPixelDoublets::doPtCut

Definition at line 99 of file gpuPixelDoublets.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ float const *__restrict__ bool bool bool gpuPixelDoublets::doZ0Cut

Definition at line 99 of file gpuPixelDoublets.h.

constexpr float gpuPixelDoublets::dzdrFact = 8 * 0.0285 / 0.015

Definition at line 49 of file gpuPixelDoubletsAlgos.h.

auto gpuPixelDoublets::first = threadIdx.x

Definition at line 79 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

constexpr auto gpuPixelDoublets::getDoubletsFromHistoMaxBlockSize = 64

Definition at line 91 of file gpuPixelDoublets.h.

constexpr auto gpuPixelDoublets::getDoubletsFromHistoMinBlocksPerMP = 16

Definition at line 92 of file gpuPixelDoublets.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const& __restrict__ gpuPixelDoublets::hh

Definition at line 26 of file gpuPixelDoubletsAlgos.h.

Referenced by __attribute__().

uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const* __restrict__ gpuPixelDoublets::hhp

Definition at line 99 of file gpuPixelDoublets.h.

Referenced by __attribute__().

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ float const *__restrict__ bool gpuPixelDoublets::ideal_cond

Definition at line 99 of file gpuPixelDoublets.h.

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

Definition at line 78 of file gpuPixelDoubletsAlgos.h.

Referenced by __attribute__(), and TTUTrackingAlg::runSeedBuster().

__shared__ uint32_t gpuPixelDoublets::innerLayerCumulativeSize[nPairsMax]

Definition at line 66 of file gpuPixelDoubletsAlgos.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell gpuPixelDoublets::isOuterHitOfCell

Definition at line 99 of file gpuPixelDoublets.h.

Referenced by __attribute__().

bool gpuPixelDoublets::isOuterLadder = ideal_cond

Definition at line 51 of file gpuPixelDoubletsAlgos.h.

__constant__ const uint8_t gpuPixelDoublets::layerPairs[2 *nPairs]

Initial value:

= {
      0, 1, 0, 4, 0, 7,              
      1, 2, 1, 4, 1, 7,              
      4, 5, 7, 8,                    
      2, 3, 2, 4, 2, 7, 5, 6, 8, 9,  
      0, 2, 1, 3,                    
      0, 5, 0, 8,                    
      4, 6, 7, 9                     
  }

Definition at line 18 of file gpuPixelDoublets.h.

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

Definition at line 59 of file gpuPixelDoubletsAlgos.h.

Referenced by if().

constexpr int gpuPixelDoublets::maxDYPred = 20

Definition at line 48 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

constexpr int gpuPixelDoublets::maxDYsize = 20

Definition at line 47 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

constexpr int gpuPixelDoublets::maxDYsize12 = 28

Definition at line 46 of file gpuPixelDoubletsAlgos.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ float const *__restrict__ bool bool bool bool uint32_t gpuPixelDoublets::maxNumOfDoublets

Initial value:

{
    auto const& __restrict__ hh = *hhp

Definition at line 109 of file gpuPixelDoublets.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ float const *__restrict__ gpuPixelDoublets::maxr

Initial value:

= {
      20., 9., 9., 20., 7., 7., 5., 5., 20., 6., 6., 5., 5., 20., 20., 9., 9., 9., 9.}

Definition at line 57 of file gpuPixelDoublets.h.

Referenced by SiPixelDQMRocLevelAnalyzer::RocSumOneModule().

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ float const *__restrict__ gpuPixelDoublets::maxz

Initial value:

= {
      20., 30., 0., 22., 30., -10., 70., 70., 22., 30., -15., 70., 70., 20., 22., 30., 0., 70., 70.}

Definition at line 55 of file gpuPixelDoublets.h.

Referenced by MultiHitGeneratorFromChi2::hitSets(), reco::modules::ParameterAdapter< CalIsolationAlgo< T, C > >::make(), and PixelRegion::PixelRegion().

constexpr int gpuPixelDoublets::minYsizeB2 = 28

Definition at line 45 of file gpuPixelDoubletsAlgos.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ float const *__restrict__ gpuPixelDoublets::minz

Initial value:

= {
      -20., 0., -30., -22., 10., -30., -70., -70., -22., 15., -30, -70., -70., -20., -22., 0, -30., -70., -70.}

Definition at line 53 of file gpuPixelDoublets.h.

Referenced by MultiHitGeneratorFromChi2::hitSets(), reco::modules::ParameterAdapter< CalIsolationAlgo< T, C > >::make(), TrackExtrapolator::propagateTrackToVolume(), and CalibrationScanAlgorithm::tuneSimultaneously().

uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const* __restrict__ GPUCACell::OuterHitOfCell int gpuPixelDoublets::nActualPairs

Definition at line 99 of file gpuPixelDoublets.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

__device__ uint32_t GPUCACell uint32_t * gpuPixelDoublets::nCells

Definition at line 99 of file gpuPixelDoublets.h.

__device__ uint32_t gpuPixelDoublets::nPairs = nPairsForTriplets + 4

Definition at line 12 of file gpuPixelDoublets.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), if(), and L1MuGMTHWFileReader::readNextEvent().

constexpr int gpuPixelDoublets::nPairsForQuadruplets = 13

Definition at line 10 of file gpuPixelDoublets.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

constexpr int gpuPixelDoublets::nPairsForTriplets = nPairsForQuadruplets + 2

Definition at line 11 of file gpuPixelDoublets.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

const int gpuPixelDoublets::nPairsMax = caConstants::maxNumberOfLayerPairs

Definition at line 64 of file gpuPixelDoubletsAlgos.h.

__shared__ uint32_t gpuPixelDoublets::ntot

Definition at line 67 of file gpuPixelDoubletsAlgos.h.

Referenced by algorithm(), MaterialBudgetVolumeAnalysis::analyze(), HGCalRecHitValidation::analyze(), HGCalDigiValidation::analyze(), DDHCalBarrelAlgo::constructGeneralVolume(), SiStripBadComponentInfo::fillBadComponentMaps(), HcalBarrelAlgo::HcalBarrelAlgo(), HGCalGeomParameters::loadGeometryHexagon8(), HGCalGeomParameters::loadWaferHexagon(), DDHGCalEEAlgo::positionSensitive(), DDHGCalModule::positionSensitive(), DDHGCalEEFileAlgo::positionSensitive(), DDHGCalHEAlgo::positionSensitive(), DDHGCalHEFileAlgo::positionSensitive(), HGCalEEFileAlgo::positionSensitive(), HGCalEEAlgo::PositionSensitive(), HGCalHEFileAlgo::positionSensitive(), HGCalHEAlgo::positionSensitive(), HLTPFJetIDProducer::produce(), TrajectorySegmentBuilder::segments(), L1EmulatorErrorFlagClient::setSummary(), HDRShower::thetaFunction(), l1tpf_impl::RegionMapper::totAndMaxInput(), and l1tpf_impl::RegionMapper::totAndMaxOutput().

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

Definition at line 56 of file gpuPixelDoubletsAlgos.h.

uint32_t gpuPixelDoublets::pairLayerId = 0

Definition at line 82 of file gpuPixelDoubletsAlgos.h.

Referenced by for().

constexpr int16_t gpuPixelDoublets::phi0p05 = 522

Definition at line 28 of file gpuPixelDoublets.h.

constexpr int16_t gpuPixelDoublets::phi0p06 = 626

Definition at line 29 of file gpuPixelDoublets.h.

constexpr int16_t gpuPixelDoublets::phi0p07 = 730

Definition at line 30 of file gpuPixelDoublets.h.

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

Definition at line 55 of file gpuPixelDoubletsAlgos.h.

__device__ uint32_t GPUCACell uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const &__restrict__ GPUCACell::OuterHitOfCell int16_t const *__restrict__ gpuPixelDoublets::phicuts

Initial value:

{phi0p05,
                                             phi0p07,
                                             phi0p07,
                                             phi0p05,
                                             phi0p06,
                                             phi0p06,
                                             phi0p05,
                                             phi0p05,
                                             phi0p06,
                                             phi0p06,
                                             phi0p06,
                                             phi0p05,
                                             phi0p05,
                                             phi0p05,
                                             phi0p05,
                                             phi0p05,
                                             phi0p05,
                                             phi0p05,
                                             phi0p05}

Definition at line 32 of file gpuPixelDoublets.h.

auto gpuPixelDoublets::stride = blockDim.x

Definition at line 80 of file gpuPixelDoubletsAlgos.h.

Referenced by HcalDigisProducerGPU::acquire(), PhysicsPerformanceDBWriterFromFile_WPandPayload::beginJob(), PhysicsPerformanceDBWriterFromFile_WPandPayload_IOV::beginJob(), calo::multifit::compute_decomposition_unrolled(), npstat::ArrayND< Numeric, StackLen, StackDim >::contract(), npstat::ArrayND< Numeric, StackLen, StackDim >::contractLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::convertToLastDimCdfLoop(), for(), npstat::ArrayND< Numeric, StackLen, StackDim >::functorFillLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::jointSliceLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::linearFillLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::makeUnit(), npstat::ArrayND< Numeric, StackLen, StackDim >::processSubrangeLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::projectInnerLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::projectInnerLoop2(), npstat::ArrayND< Numeric, StackLen, StackDim >::projectLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::projectLoop2(), npstat::ArrayND< Numeric, StackLen, StackDim >::scaleBySliceInnerLoop(), npstat::ArrayND< Numeric, StackLen, StackDim >::scaleBySliceLoop(), l1tVertexFinder::VertexFinder::strided_iota(), and npstat::ArrayND< Numeric, StackLen, StackDim >::sumBelowLoop().