CMS 3D CMS Logo
List of all members | Public Types | Public Member Functions | Public Attributes | Static Public Attributes
GPUCACellT< TrackerTraits > Class Template Reference

#include <GPUCACell.h>

Public Types

using CellNeighbors = caStructures::CellNeighborsT< TrackerTraits >
 
using CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
 
using CellTracks = caStructures::CellTracksT< TrackerTraits >
 
using CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
 
using hindex_type = typename TrackerTraits::hindex_type
 
using HitContainer = typename TrackSoA< TrackerTraits >::HitContainer
 
using HitsConstView = TrackingRecHitSoAConstView< TrackerTraits >
 
using OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
 
using OuterHitOfCellContainer = caStructures::OuterHitOfCellContainerT< TrackerTraits >
 
using PtrAsInt = unsigned long long
 
using Quality = pixelTrack::Quality
 
enum  StatusBit : uint16_t { StatusBit::kUsed = 1, StatusBit::kInTrack = 2, StatusBit::kKilled = 1 << 15 }
 
using tindex_type = typename TrackerTraits::tindex_type
 
using TmpTuple = cms::cuda::VecArray< uint32_t, TrackerTraits::maxDepth >
 

Public Member Functions

__device__ __attribute__ ((always_inline)) void init(CellNeighborsVector &cellNeighbors
 
__device__ __attribute__ ((always_inline)) int addOuterNeighbor(typename TrackerTraits
 
__device__ __attribute__ ((always_inline)) int addTrack(tindex_type t
 
__device__ __attribute__ ((always_inline)) CellTracks &tracks()
 
__device__ __attribute__ ((always_inline)) CellTracks const &tracks() const
 
__device__ __attribute__ ((always_inline)) CellNeighbors &outerNeighbors()
 
__device__ __attribute__ ((always_inline)) CellNeighbors const &outerNeighbors() const
 
__device__ __attribute__ ((always_inline)) float inner_x(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_x(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_y(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_y(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_z(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_z(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_r(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_r(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) auto inner_iphi(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) auto outer_iphi(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_detIndex(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_detIndex(const HitsConstView &hh) const
 
__device__ __attribute__ ((always_inline)) static bool areAlignedRZ(float r1
 
 assert (outerNeighbors().empty())
 
 assert (tracks().empty())
 
__device__ bool check_alignment (const HitsConstView &hh, GPUCACellT const &otherCell, const float ptmin, const float hardCurvCut, const float caThetaCutBarrel, const float caThetaCutForward, const float dcaCutInnerTriplet, const float dcaCutOuterTriplet) const
 
 GPUCACellT ()=default
 
constexpr unsigned int inner_hit_id () const
 
constexpr unsigned int outer_hit_id () const
 
__device__ void print_cell () const
 
return tracks ().push_back(t)
 

Public Attributes

__device__ CellTracksVector & cellTracks
 
float distance_13_squared = radius_diff * radius_diff + (z1 - zo) * (z1 - zo)
 
__device__ CellTracksVector const HitsConstView & hh
 
__device__ CellTracksVector const HitsConstView int hindex_type innerHitId
 
__device__ CellTracksVector const HitsConstView int layerPairId
 
__device__ CellTracksVector const HitsConstView int hindex_type hindex_type outerHitId
 
float pMin = ptmin * std::sqrt(distance_13_squared)
 
__device__ float float float float float const float ptmin
 
__device__ float float ri
 
__device__ float float float float ro
 
float tan_12_13_half_mul_distance_13_squared = fabs(z1 * (ri - ro) + zi * (ro - r1) + zo * (r1 - ri))
 
 theFishboneId = invalidHitId
 
 theInnerR = hh[innerHitId].rGlobal()
 
 theInnerZ = hh[innerHitId].zGlobal()
 
 theLayerPairId_ = layerPairId
 
 theOuterHitId = outerHitId
 
 theOuterNeighbors = &cellNeighbors[0]
 
 theStatus_ = 0
 
__device__ float float float float float const float const float thetaCut
 
 theTracks = &cellTracks[0]
 
__device__ float z1
 
__device__ float float float zi
 
__device__ float float float float float zo
 

Static Public Attributes

static constexpr auto bad = pixelTrack::Quality::bad
 
static constexpr auto invalidHitId = std::numeric_limits<hindex_type>::max()
 
static constexpr auto maxCellsPerHit = TrackerTraits::maxCellsPerHit
 

Detailed Description

template<typename TrackerTraits>
class GPUCACellT< TrackerTraits >

Definition at line 22 of file GPUCACell.h.

Member Typedef Documentation

◆ CellNeighbors

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::CellNeighbors = caStructures::CellNeighborsT<TrackerTraits>

Definition at line 29 of file GPUCACell.h.

◆ CellNeighborsVector

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::CellNeighborsVector = caStructures::CellNeighborsVectorT<TrackerTraits>

Definition at line 31 of file GPUCACell.h.

◆ CellTracks

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::CellTracks = caStructures::CellTracksT<TrackerTraits>

Definition at line 30 of file GPUCACell.h.

◆ CellTracksVector

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::CellTracksVector = caStructures::CellTracksVectorT<TrackerTraits>

Definition at line 32 of file GPUCACell.h.

◆ hindex_type

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::hindex_type = typename TrackerTraits::hindex_type

Definition at line 35 of file GPUCACell.h.

◆ HitContainer

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::HitContainer = typename TrackSoA<TrackerTraits>::HitContainer

Definition at line 41 of file GPUCACell.h.

◆ HitsConstView

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::HitsConstView = TrackingRecHitSoAConstView<TrackerTraits>

Definition at line 34 of file GPUCACell.h.

◆ OuterHitOfCell

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::OuterHitOfCell = caStructures::OuterHitOfCellT<TrackerTraits>

Definition at line 28 of file GPUCACell.h.

◆ OuterHitOfCellContainer

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::OuterHitOfCellContainer = caStructures::OuterHitOfCellContainerT<TrackerTraits>

Definition at line 27 of file GPUCACell.h.

◆ PtrAsInt

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::PtrAsInt = unsigned long long

Definition at line 24 of file GPUCACell.h.

◆ Quality

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::Quality = pixelTrack::Quality

Definition at line 42 of file GPUCACell.h.

◆ tindex_type

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::tindex_type = typename TrackerTraits::tindex_type

Definition at line 36 of file GPUCACell.h.

◆ TmpTuple

template<typename TrackerTraits>
using GPUCACellT< TrackerTraits >::TmpTuple = cms::cuda::VecArray<uint32_t, TrackerTraits::maxDepth>

Definition at line 39 of file GPUCACell.h.

Member Enumeration Documentation

◆ StatusBit

template<typename TrackerTraits>
enum GPUCACellT::StatusBit : uint16_t
strong
Enumerator
kUsed 
kInTrack 
kKilled 

Definition at line 45 of file GPUCACell.h.

45 : uint16_t { kUsed = 1, kInTrack = 2, kKilled = 1 << 15 };

Constructor & Destructor Documentation

◆ GPUCACellT()

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::GPUCACellT ( )
default

Member Function Documentation

◆ __attribute__() [1/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) &
inline

◆ __attribute__() [2/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  )
inline

Definition at line 72 of file GPUCACell.h.

References cms::cudacompat::__threadfence(), cms::cudacompat::atomicCAS(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::cellNeighbors, relativeConstraints::empty, mps_fire::i, submitPVValidationJobs::t, GPUCACellT< TrackerTraits >::theOuterNeighbors, and SiPixelPI::zero.

73  {
74  // use smart cache
75  if (outerNeighbors().empty()) {
76  auto i = cellNeighbors.extend(); // maybe wasted....
77  if (i > 0) {
78  cellNeighbors[i].reset();
79  __threadfence();
80 #ifdef __CUDACC__
81  auto zero = (PtrAsInt)(&cellNeighbors[0]);
82  atomicCAS((PtrAsInt*)(&theOuterNeighbors),
83  zero,
84  (PtrAsInt)(&cellNeighbors[i])); // if fails we cannot give "i" back...
85 #else
86  theOuterNeighbors = &cellNeighbors[i];
87 #endif
88  } else
89  return -1;
90  }
91  __threadfence();
92  return outerNeighbors().push_back(t);
93  }
mps_fire.i
i
Definition: mps_fire.py:429
cms::cudacompat::atomicCAS
T1 atomicCAS(T1 *address, T1 compare, T2 val)
Definition: cudaCompat.h:36
ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::cellNeighbors
ALPAKA_FN_ACC ALPAKA_FN_INLINE void uint32_t const uint32_t CACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > * cellNeighbors
Definition: CAPixelDoubletsAlgos.h:133
GPUCACellT::PtrAsInt
unsigned long long PtrAsInt
Definition: GPUCACell.h:24
cms::cudacompat::__threadfence
void __threadfence()
Definition: cudaCompat.h:133
GPUCACellT::theOuterNeighbors
theOuterNeighbors
Definition: GPUCACell.h:66

◆ __attribute__() [3/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  )
inline

◆ __attribute__() [4/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) &
inline

Definition at line 114 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theTracks.

114 { return *theTracks; }

◆ __attribute__() [5/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 115 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theTracks.

115 { return *theTracks; }

◆ __attribute__() [6/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) &
inline

Definition at line 116 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theOuterNeighbors.

116 { return *theOuterNeighbors; }
GPUCACellT::theOuterNeighbors
theOuterNeighbors
Definition: GPUCACell.h:66

◆ __attribute__() [7/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 117 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theOuterNeighbors.

117 { return *theOuterNeighbors; }
GPUCACellT::theOuterNeighbors
theOuterNeighbors
Definition: GPUCACell.h:66

◆ __attribute__() [8/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 118 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh.

118 { return hh[theInnerHitId].xGlobal(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [9/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 119 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh, and GPUCACellT< TrackerTraits >::theOuterHitId.

119 { return hh[theOuterHitId].xGlobal(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [10/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 120 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh.

120 { return hh[theInnerHitId].yGlobal(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [11/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 121 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh, and GPUCACellT< TrackerTraits >::theOuterHitId.

121 { return hh[theOuterHitId].yGlobal(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [12/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 122 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theInnerZ.

122 { return theInnerZ; }

◆ __attribute__() [13/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 124 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh, and GPUCACellT< TrackerTraits >::theOuterHitId.

124 { return hh[theOuterHitId].zGlobal(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [14/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 125 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theInnerR.

125 { return theInnerR; }

◆ __attribute__() [15/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 127 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh, and GPUCACellT< TrackerTraits >::theOuterHitId.

127 { return hh[theOuterHitId].rGlobal(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [16/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 129 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh.

129 { return hh[theInnerHitId].iphi(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [17/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 130 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh, and GPUCACellT< TrackerTraits >::theOuterHitId.

130 { return hh[theOuterHitId].iphi(); }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [18/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 132 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh.

132  {
133  return hh[theInnerHitId].detectorIndex();
134  }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [19/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 135 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::hh, and GPUCACellT< TrackerTraits >::theOuterHitId.

135  {
136  return hh[theOuterHitId].detectorIndex();
137  }
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50

◆ __attribute__() [20/20]

template<typename TrackerTraits>
__device__ GPUCACellT< TrackerTraits >::__attribute__ ( (always_inline)  )
inline

◆ assert() [1/2]

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::assert ( outerNeighbors().empty()  )

◆ assert() [2/2]

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::assert ( tracks().empty()  )

◆ check_alignment()

template<typename TrackerTraits>
__device__ bool GPUCACellT< TrackerTraits >::check_alignment ( const HitsConstView &  hh,
GPUCACellT< TrackerTraits > const &  otherCell,
const float  ptmin,
const float  hardCurvCut,
const float  caThetaCutBarrel,
const float  caThetaCutForward,
const float  dcaCutInnerTriplet,
const float  dcaCutOuterTriplet 
) const
inline

Definition at line 149 of file GPUCACell.h.

References particleFlowDisplacedVertexCandidate_cfi::dcaCut, HLT_2023v12_cff::dcaCutInnerTriplet, HLT_2023v12_cff::dcaCutOuterTriplet, HLT_2023v12_cff::hardCurvCut, GPUCACellT< TrackerTraits >::hh, PixelPluginsPhase0_cfi::isBarrel, GPUCACellT< TrackerTraits >::ptmin, diffTwoXMLs::r1, GPUCACellT< TrackerTraits >::ri, GPUCACellT< TrackerTraits >::ro, GPUCACellT< TrackerTraits >::z1, GPUCACellT< TrackerTraits >::zi, and GPUCACellT< TrackerTraits >::zo.

156  {
157  // detIndex of the layerStart for the Phase1 Pixel Detector:
158  // [BPX1, BPX2, BPX3, BPX4, FP1, FP2, FP3, FN1, FN2, FN3, LAST_VALID]
159  // [ 0, 96, 320, 672, 1184, 1296, 1408, 1520, 1632, 1744, 1856]
160  auto ri = inner_r(hh);
161  auto zi = inner_z(hh);
162 
163  auto ro = outer_r(hh);
164  auto zo = outer_z(hh);
165 
166  auto r1 = otherCell.inner_r(hh);
167  auto z1 = otherCell.inner_z(hh);
168  auto isBarrel = otherCell.outer_detIndex(hh) < TrackerTraits::last_barrel_detIndex;
169  bool aligned = areAlignedRZ(r1,
170  z1,
171  ri,
172  zi,
173  ro,
174  zo,
175  ptmin,
176  isBarrel ? caThetaCutBarrel : caThetaCutForward); // 2.f*thetaCut); // FIXME tune cuts
177  return (aligned && dcaCut(hh,
178  otherCell,
179  otherCell.inner_detIndex(hh) < TrackerTraits::last_bpix1_detIndex ? dcaCutInnerTriplet
180  : dcaCutOuterTriplet,
181  hardCurvCut)); // FIXME tune cuts
182  }
particleFlowDisplacedVertexCandidate_cfi.dcaCut
dcaCut
Definition: particleFlowDisplacedVertexCandidate_cfi.py:17
GPUCACellT::ri
__device__ float float ri
Definition: GPUCACell.h:185
HLT_2023v12_cff.dcaCutInnerTriplet
dcaCutInnerTriplet
Definition: HLT_2023v12_cff.py:8580
GPUCACellT::zo
__device__ float float float float float zo
Definition: GPUCACell.h:185
PixelPluginsPhase0_cfi.isBarrel
isBarrel
Definition: PixelPluginsPhase0_cfi.py:17
GPUCACellT::zi
__device__ float float float zi
Definition: GPUCACell.h:185
GPUCACellT::hh
__device__ CellTracksVector const HitsConstView & hh
Definition: GPUCACell.h:50
GPUCACellT::ptmin
__device__ float float float float float const float ptmin
Definition: GPUCACell.h:185
GPUCACellT::z1
__device__ float z1
Definition: GPUCACell.h:185
HLT_2023v12_cff.dcaCutOuterTriplet
dcaCutOuterTriplet
Definition: HLT_2023v12_cff.py:8581
HLT_2023v12_cff.hardCurvCut
hardCurvCut
Definition: HLT_2023v12_cff.py:8579
GPUCACellT::ro
__device__ float float float float ro
Definition: GPUCACell.h:185

◆ inner_hit_id()

template<typename TrackerTraits>
constexpr unsigned int GPUCACellT< TrackerTraits >::inner_hit_id ( ) const
inline

Definition at line 139 of file GPUCACell.h.

139 { return theInnerHitId; }

◆ outer_hit_id()

template<typename TrackerTraits>
constexpr unsigned int GPUCACellT< TrackerTraits >::outer_hit_id ( ) const
inline

Definition at line 140 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theOuterHitId.

140 { return theOuterHitId; }

◆ print_cell()

template<typename TrackerTraits>
__device__ void GPUCACellT< TrackerTraits >::print_cell ( ) const
inline

Definition at line 142 of file GPUCACell.h.

References GPUCACellT< TrackerTraits >::theLayerPairId_, and GPUCACellT< TrackerTraits >::theOuterHitId.

142  {
143  printf("printing cell: on layerPair: %d, innerHitId: %d, outerHitId: %d \n",
144  theLayerPairId_,
145  theInnerHitId,
146  theOuterHitId);
147  }

◆ tracks()

template<typename TrackerTraits>
return GPUCACellT< TrackerTraits >::tracks ( )

Member Data Documentation

◆ bad

template<typename TrackerTraits>
constexpr auto GPUCACellT< TrackerTraits >::bad = pixelTrack::Quality::bad
static

Definition at line 43 of file GPUCACell.h.

◆ cellTracks

template<typename TrackerTraits>
__device__ CellTracksVector & GPUCACellT< TrackerTraits >::cellTracks
Initial value:
{
if (tracks().empty()) {
auto i = cellTracks.extend();
if (i > 0) {
cellTracks[i].reset();
__threadfence();
theTracks = &cellTracks[i];
} else
return -1;
}
__threadfence()

Definition at line 50 of file GPUCACell.h.

◆ distance_13_squared

template<typename TrackerTraits>
float GPUCACellT< TrackerTraits >::distance_13_squared = radius_diff * radius_diff + (z1 - zo) * (z1 - zo)

Definition at line 187 of file GPUCACell.h.

◆ hh

template<typename TrackerTraits>
__device__ CellTracksVector const HitsConstView& GPUCACellT< TrackerTraits >::hh

Definition at line 50 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__(), and GPUCACellT< TrackerTraits >::check_alignment().

◆ innerHitId

template<typename TrackerTraits>
__device__ CellTracksVector const HitsConstView int hindex_type GPUCACellT< TrackerTraits >::innerHitId

Definition at line 50 of file GPUCACell.h.

◆ invalidHitId

template<typename TrackerTraits>
constexpr auto GPUCACellT< TrackerTraits >::invalidHitId = std::numeric_limits<hindex_type>::max()
static

Definition at line 37 of file GPUCACell.h.

◆ layerPairId

template<typename TrackerTraits>
__device__ CellTracksVector const HitsConstView int GPUCACellT< TrackerTraits >::layerPairId

Definition at line 50 of file GPUCACell.h.

◆ maxCellsPerHit

template<typename TrackerTraits>
constexpr auto GPUCACellT< TrackerTraits >::maxCellsPerHit = TrackerTraits::maxCellsPerHit
static

Definition at line 26 of file GPUCACell.h.

◆ outerHitId

template<typename TrackerTraits>
__device__ CellTracksVector const HitsConstView int hindex_type hindex_type GPUCACellT< TrackerTraits >::outerHitId
Initial value:
{
theInnerHitId = innerHitId

Definition at line 54 of file GPUCACell.h.

◆ pMin

template<typename TrackerTraits>
float GPUCACellT< TrackerTraits >::pMin = ptmin * std::sqrt(distance_13_squared)

Definition at line 189 of file GPUCACell.h.

◆ ptmin

template<typename TrackerTraits>
__device__ float float float float float const float GPUCACellT< TrackerTraits >::ptmin

Definition at line 185 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::check_alignment().

◆ ri

template<typename TrackerTraits>
__device__ float float GPUCACellT< TrackerTraits >::ri

Definition at line 185 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::check_alignment().

◆ ro

template<typename TrackerTraits>
__device__ float float float float GPUCACellT< TrackerTraits >::ro

Definition at line 185 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::check_alignment().

◆ tan_12_13_half_mul_distance_13_squared

template<typename TrackerTraits>
float GPUCACellT< TrackerTraits >::tan_12_13_half_mul_distance_13_squared = fabs(z1 * (ri - ro) + zi * (ro - r1) + zo * (r1 - ri))

Definition at line 192 of file GPUCACell.h.

◆ theFishboneId

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theFishboneId = invalidHitId

Definition at line 59 of file GPUCACell.h.

◆ theInnerR

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theInnerR = hh[innerHitId].rGlobal()

Definition at line 63 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__().

◆ theInnerZ

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theInnerZ = hh[innerHitId].zGlobal()

Definition at line 62 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__().

◆ theLayerPairId_

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theLayerPairId_ = layerPairId

Definition at line 57 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::print_cell().

◆ theOuterHitId

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theOuterHitId = outerHitId

Definition at line 56 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__(), GPUCACellT< TrackerTraits >::outer_hit_id(), and GPUCACellT< TrackerTraits >::print_cell().

◆ theOuterNeighbors

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theOuterNeighbors = &cellNeighbors[0]

Definition at line 66 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__().

◆ theStatus_

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theStatus_ = 0

Definition at line 58 of file GPUCACell.h.

◆ thetaCut

template<typename TrackerTraits>
__device__ float float float float float const float const float GPUCACellT< TrackerTraits >::thetaCut
Initial value:
{
float radius_diff = std::abs(r1 - ro)

Definition at line 185 of file GPUCACell.h.

◆ theTracks

template<typename TrackerTraits>
GPUCACellT< TrackerTraits >::theTracks = &cellTracks[0]

Definition at line 67 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__().

◆ z1

template<typename TrackerTraits>
__device__ float GPUCACellT< TrackerTraits >::z1

Definition at line 185 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::check_alignment().

◆ zi

template<typename TrackerTraits>
__device__ float float float GPUCACellT< TrackerTraits >::zi

Definition at line 185 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::check_alignment().

◆ zo

template<typename TrackerTraits>
__device__ float float float float float GPUCACellT< TrackerTraits >::zo

Definition at line 185 of file GPUCACell.h.

Referenced by GPUCACellT< TrackerTraits >::check_alignment().

Generated for CMSSW Reference Manual by doxygen 1.8.14