CMS 3D CMS Logo
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Properties Friends Macros Groups Pages
List of all members | Public Types | Public Member Functions | Public Attributes | Static Public Attributes
GPUCACell Class Reference

#include <GPUCACell.h>

Public Types

using CellNeighbors = caConstants::CellNeighbors
 
using CellNeighborsVector = caConstants::CellNeighborsVector
 
using CellTracks = caConstants::CellTracks
 
using CellTracksVector = caConstants::CellTracksVector
 
using hindex_type = Hits::hindex_type
 
using HitContainer = pixelTrack::HitContainer
 
using Hits = TrackingRecHit2DSOAView
 
using OuterHitOfCell = caConstants::OuterHitOfCell
 
using OuterHitOfCellContainer = caConstants::OuterHitOfCellContainer
 
using PtrAsInt = unsigned long long
 
using Quality = pixelTrack::Quality
 
enum  StatusBit : uint16_t { StatusBit::kUsed = 1, StatusBit::kInTrack = 2, StatusBit::kKilled = 1 << 15 }
 
using TmpTuple = cms::cuda::VecArray< uint32_t, 6 >
 

Public Member Functions

__device__ __attribute__ ((always_inline)) void init(CellNeighborsVector &cellNeighbors
 
__device__ __attribute__ ((always_inline)) int addOuterNeighbor(CellNeighbors
 
__device__ __attribute__ ((always_inline)) int addTrack(CellTracks
 
__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(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_x(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_y(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_y(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_z(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_z(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_r(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_r(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) auto inner_iphi(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) auto outer_iphi(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float inner_detIndex(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) float outer_detIndex(Hits const &hh) const
 
__device__ __attribute__ ((always_inline)) static bool areAlignedRZ(float r1
 
 assert (outerNeighbors().empty())
 
 assert (tracks().empty())
 
__device__ bool check_alignment (Hits const &hh, GPUCACell const &otherCell, const float ptmin, const float hardCurvCut, const float caThetaCutBarrel, const float caThetaCutForward, const float dcaCutInnerTriplet, const float dcaCutOuterTriplet) const
 
 GPUCACell ()=default
 
constexpr unsigned int inner_hit_id () const
 
constexpr unsigned int outer_hit_id () const
 
__device__ void print_cell () const
 

Public Attributes

__device__ CellTracksVector & cellTracks
 
float distance_13_squared = radius_diff * radius_diff + (z1 - zo) * (z1 - zo)
 
__device__ CellTracksVector
Hits const & 		hh
 
__device__ CellTracksVector
Hits const int hindex_type 		innerHitId
 
__device__ CellTracksVector
Hits const int 		layerPairId
 
__device__ CellTracksVector
Hits const 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 = std::numeric_limits<hindex_type>::max()
 
 theInnerR = hh.rGlobal(innerHitId)
 
 theInnerZ = hh.zGlobal(innerHitId)
 
 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 maxCellsPerHit = caConstants::maxCellsPerHit
 

Detailed Description

Definition at line 20 of file GPUCACell.h.

Member Typedef Documentation

using GPUCACell::CellNeighbors = caConstants::CellNeighbors

Definition at line 27 of file GPUCACell.h.

using GPUCACell::CellNeighborsVector = caConstants::CellNeighborsVector

Definition at line 29 of file GPUCACell.h.

using GPUCACell::CellTracks = caConstants::CellTracks

Definition at line 28 of file GPUCACell.h.

using GPUCACell::CellTracksVector = caConstants::CellTracksVector

Definition at line 30 of file GPUCACell.h.

using GPUCACell::hindex_type = Hits::hindex_type

Definition at line 33 of file GPUCACell.h.

using GPUCACell::HitContainer = pixelTrack::HitContainer

Definition at line 37 of file GPUCACell.h.

using GPUCACell::Hits = TrackingRecHit2DSOAView

Definition at line 32 of file GPUCACell.h.

using GPUCACell::OuterHitOfCell = caConstants::OuterHitOfCell

Definition at line 26 of file GPUCACell.h.

using GPUCACell::OuterHitOfCellContainer = caConstants::OuterHitOfCellContainer

Definition at line 25 of file GPUCACell.h.

using GPUCACell::PtrAsInt = unsigned long long

Definition at line 22 of file GPUCACell.h.

using GPUCACell::Quality = pixelTrack::Quality

Definition at line 38 of file GPUCACell.h.

using GPUCACell::TmpTuple = cms::cuda::VecArray<uint32_t, 6>

Definition at line 35 of file GPUCACell.h.

Member Enumeration Documentation

enum GPUCACell::StatusBit : uint16_t
strong
Enumerator
kUsed 
kInTrack 
kKilled 

Definition at line 41 of file GPUCACell.h.

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

Constructor & Destructor Documentation

GPUCACell::GPUCACell ( )
default

Member Function Documentation

__device__ GPUCACell::__attribute__ ( (always_inline)  )
inline
__device__ GPUCACell::__attribute__ ( (always_inline)  )
inline

Definition at line 68 of file GPUCACell.h.

References cms::cudacompat::__threadfence(), cms::cudacompat::atomicCAS(), relativeConstraints::empty, cms::cuda::SimpleVector< T >::extend(), mps_fire::i, cms::cuda::SimpleVector< T >::reset(), theOuterNeighbors, and SiPixelPI::zero.

68  {
69  // use smart cache
70  if (outerNeighbors().empty()) {
71  auto i = cellNeighbors.extend(); // maybe wasted....
72  if (i > 0) {
73  cellNeighbors[i].reset();
74  __threadfence();
75 #ifdef __CUDACC__
76  auto zero = (PtrAsInt)(&cellNeighbors[0]);
77  atomicCAS((PtrAsInt*)(&theOuterNeighbors),
78  zero,
79  (PtrAsInt)(&cellNeighbors[i])); // if fails we cannot give "i" back...
80 #else
81  theOuterNeighbors = &cellNeighbors[i];
82 #endif
83  } else
84  return -1;
85  }
86  __threadfence();
87  return outerNeighbors().push_back(t);
88  }
GPUCACell::PtrAsInt
unsigned long long PtrAsInt
Definition: GPUCACell.h:22
mps_fire.i
i
Definition: mps_fire.py:428
cms::cudacompat::atomicCAS
T1 atomicCAS(T1 *address, T1 compare, T2 val)
Definition: cudaCompat.h:36
cellNeighbors
caConstants::TupleMultiplicity const CAHitNtupletGeneratorKernelsGPU::HitToTuple const cms::cuda::AtomicPairCounter GPUCACell const *__restrict__ uint32_t const *__restrict__ gpuPixelDoublets::CellNeighborsVector const * cellNeighbors
Definition: CAHitNtupletGeneratorKernelsImpl.h:43
GPUCACell::theOuterNeighbors
theOuterNeighbors
Definition: GPUCACell.h:62
cms::cudacompat::__threadfence
void __threadfence()
Definition: cudaCompat.h:109
__device__ GPUCACell::__attribute__ ( (always_inline)  )
inline

Definition at line 90 of file GPUCACell.h.

References cms::cudacompat::__threadfence(), cms::cudacompat::atomicCAS(), relativeConstraints::empty, cms::cuda::SimpleVector< T >::extend(), mps_fire::i, cms::cuda::SimpleVector< T >::reset(), theTracks, tracks, and SiPixelPI::zero.

90  {
91  if (tracks().empty()) {
92  auto i = cellTracks.extend(); // maybe wasted....
93  if (i > 0) {
94  cellTracks[i].reset();
95  __threadfence();
96 #ifdef __CUDACC__
97  auto zero = (PtrAsInt)(&cellTracks[0]);
98  atomicCAS((PtrAsInt*)(&theTracks), zero, (PtrAsInt)(&cellTracks[i])); // if fails we cannot give "i" back...
99 #else
100  theTracks = &cellTracks[i];
101 #endif
102  } else
103  return -1;
104  }
105  __threadfence();
106  return tracks().push_back(t);
107  }
GPUCACell::PtrAsInt
unsigned long long PtrAsInt
Definition: GPUCACell.h:22
mps_fire.i
i
Definition: mps_fire.py:428
cms::cudacompat::atomicCAS
T1 atomicCAS(T1 *address, T1 compare, T2 val)
Definition: cudaCompat.h:36
tracks
auto const & tracks
cannot be loose
Definition: CAHitNtupletGeneratorKernelsImpl.h:146
cms::cuda::SimpleVector::extend
__device__ int extend(int size=1)
Definition: SimpleVector.h:84
cms::cuda::SimpleVector::reset
constexpr void reset()
Definition: SimpleVector.h:108
GPUCACell::cellTracks
__device__ CellTracksVector & cellTracks
Definition: GPUCACell.h:46
cms::cudacompat::__threadfence
void __threadfence()
Definition: cudaCompat.h:109
__device__ GPUCACell::__attribute__ ( (always_inline)  )
inline

Definition at line 109 of file GPUCACell.h.

References theTracks.

109 { return *theTracks; }
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 110 of file GPUCACell.h.

References theTracks.

110 { return *theTracks; }
__device__ GPUCACell::__attribute__ ( (always_inline)  )
inline

Definition at line 111 of file GPUCACell.h.

References theOuterNeighbors.

111 { return *theOuterNeighbors; }
GPUCACell::theOuterNeighbors
theOuterNeighbors
Definition: GPUCACell.h:62
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 112 of file GPUCACell.h.

References theOuterNeighbors.

112 { return *theOuterNeighbors; }
GPUCACell::theOuterNeighbors
theOuterNeighbors
Definition: GPUCACell.h:62
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 113 of file GPUCACell.h.

113 { return hh.xGlobal(theInnerHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 114 of file GPUCACell.h.

References theOuterHitId.

114 { return hh.xGlobal(theOuterHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 115 of file GPUCACell.h.

115 { return hh.yGlobal(theInnerHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 116 of file GPUCACell.h.

References theOuterHitId.

116 { return hh.yGlobal(theOuterHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 117 of file GPUCACell.h.

References theInnerZ.

117 { return theInnerZ; }
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 119 of file GPUCACell.h.

References theOuterHitId.

119 { return hh.zGlobal(theOuterHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 120 of file GPUCACell.h.

References theInnerR.

120 { return theInnerR; }
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 122 of file GPUCACell.h.

References theOuterHitId.

122 { return hh.rGlobal(theOuterHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 124 of file GPUCACell.h.

124 { return hh.iphi(theInnerHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 125 of file GPUCACell.h.

References theOuterHitId.

125 { return hh.iphi(theOuterHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 127 of file GPUCACell.h.

127 { return hh.detectorIndex(theInnerHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  ) const
inline

Definition at line 128 of file GPUCACell.h.

References theOuterHitId.

128 { return hh.detectorIndex(theOuterHitId); }
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
__device__ GPUCACell::__attribute__ ( (always_inline)  )
inline
GPUCACell::assert ( outerNeighbors().empty()  )
GPUCACell::assert ( tracks().empty()  )
__device__ bool GPUCACell::check_alignment ( Hits const &  hh,
GPUCACell 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 140 of file GPUCACell.h.

References PixelPluginsPhase0_cfi::isBarrel, caConstants::last_barrel_detIndex, caConstants::last_bpix1_detIndex, diffTwoXMLs::r1, ri, ro, z1, zi, and zo.

147  {
148  // detIndex of the layerStart for the Phase1 Pixel Detector:
149  // [BPX1, BPX2, BPX3, BPX4, FP1, FP2, FP3, FN1, FN2, FN3, LAST_VALID]
150  // [ 0, 96, 320, 672, 1184, 1296, 1408, 1520, 1632, 1744, 1856]
151  auto ri = inner_r(hh);
152  auto zi = inner_z(hh);
153 
154  auto ro = outer_r(hh);
155  auto zo = outer_z(hh);
156 
157  auto r1 = otherCell.inner_r(hh);
158  auto z1 = otherCell.inner_z(hh);
159  auto isBarrel = otherCell.outer_detIndex(hh) < caConstants::last_barrel_detIndex;
160  bool aligned = areAlignedRZ(r1,
161  z1,
162  ri,
163  zi,
164  ro,
165  zo,
166  ptmin,
167  isBarrel ? caThetaCutBarrel : caThetaCutForward); // 2.f*thetaCut); // FIXME tune cuts
168  return (aligned && dcaCut(hh,
169  otherCell,
170  otherCell.inner_detIndex(hh) < caConstants::last_bpix1_detIndex ? dcaCutInnerTriplet
171  : dcaCutOuterTriplet,
172  hardCurvCut)); // FIXME tune cuts
173  }
GPUCACell::ro
__device__ float float float float ro
Definition: GPUCACell.h:176
PixelPluginsPhase0_cfi.isBarrel
tuple isBarrel
Definition: PixelPluginsPhase0_cfi.py:17
GPUCACell::zo
__device__ float float float float float zo
Definition: GPUCACell.h:176
GPUCACell::zi
__device__ float float float zi
Definition: GPUCACell.h:176
GPUCACell::ptmin
__device__ float float float float float const float ptmin
Definition: GPUCACell.h:176
GPUCACell::z1
__device__ float z1
Definition: GPUCACell.h:176
GPUCACell::hh
__device__ CellTracksVector Hits const & hh
Definition: GPUCACell.h:46
caConstants::last_bpix1_detIndex
constexpr uint32_t last_bpix1_detIndex
Definition: CAConstants.h:64
caConstants::last_barrel_detIndex
constexpr uint32_t last_barrel_detIndex
Definition: CAConstants.h:65
GPUCACell::ri
__device__ float float ri
Definition: GPUCACell.h:176
constexpr unsigned int GPUCACell::inner_hit_id ( ) const
inline

Definition at line 130 of file GPUCACell.h.

130 { return theInnerHitId; }
constexpr unsigned int GPUCACell::outer_hit_id ( ) const
inline

Definition at line 131 of file GPUCACell.h.

References theOuterHitId.

131 { return theOuterHitId; }
__device__ void GPUCACell::print_cell ( ) const
inline

Definition at line 133 of file GPUCACell.h.

References gpuVertexFinder::printf(), theLayerPairId_, and theOuterHitId.

133  {
134  printf("printing cell: on layerPair: %d, innerHitId: %d, outerHitId: %d \n",
135  theLayerPairId_,
136  theInnerHitId,
137  theOuterHitId);
138  }
GPUCACell::theLayerPairId_
theLayerPairId_
Definition: GPUCACell.h:53
gpuVertexFinder::printf
printf("params %d %f %f %f\n", minT, eps, errmax, chi2max)

Member Data Documentation

constexpr auto GPUCACell::bad = pixelTrack::Quality::bad
static

Definition at line 39 of file GPUCACell.h.

__device__ CellTracksVector& GPUCACell::cellTracks

Definition at line 46 of file GPUCACell.h.

float GPUCACell::distance_13_squared = radius_diff * radius_diff + (z1 - zo) * (z1 - zo)

Definition at line 178 of file GPUCACell.h.

__device__ CellTracksVector Hits const& GPUCACell::hh

Definition at line 46 of file GPUCACell.h.

__device__ CellTracksVector Hits const int hindex_type GPUCACell::innerHitId

Definition at line 46 of file GPUCACell.h.

__device__ CellTracksVector Hits const int GPUCACell::layerPairId

Definition at line 46 of file GPUCACell.h.

Referenced by gpuPixelDoublets::__attribute__().

constexpr auto GPUCACell::maxCellsPerHit = caConstants::maxCellsPerHit
static

Definition at line 24 of file GPUCACell.h.

Referenced by gpuPixelDoublets::__attribute__().

__device__ CellTracksVector Hits const int hindex_type hindex_type GPUCACell::outerHitId
Initial value:
{
theInnerHitId = innerHitId

Definition at line 50 of file GPUCACell.h.

float GPUCACell::pMin = ptmin * std::sqrt(distance_13_squared)

Definition at line 180 of file GPUCACell.h.

__device__ float float float float float const float GPUCACell::ptmin

Definition at line 176 of file GPUCACell.h.

__device__ float float GPUCACell::ri

Definition at line 176 of file GPUCACell.h.

Referenced by check_alignment().

__device__ float float float float GPUCACell::ro

Definition at line 176 of file GPUCACell.h.

Referenced by check_alignment().

float GPUCACell::tan_12_13_half_mul_distance_13_squared = fabs(z1 * (ri - ro) + zi * (ro - r1) + zo * (r1 - ri))

Definition at line 183 of file GPUCACell.h.

GPUCACell::theFishboneId = std::numeric_limits<hindex_type>::max()

Definition at line 55 of file GPUCACell.h.

GPUCACell::theInnerR = hh.rGlobal(innerHitId)

Definition at line 59 of file GPUCACell.h.

Referenced by __attribute__().

GPUCACell::theInnerZ = hh.zGlobal(innerHitId)

Definition at line 58 of file GPUCACell.h.

Referenced by __attribute__().

GPUCACell::theLayerPairId_ = layerPairId

Definition at line 53 of file GPUCACell.h.

Referenced by print_cell().

GPUCACell::theOuterHitId = outerHitId

Definition at line 52 of file GPUCACell.h.

Referenced by __attribute__(), outer_hit_id(), and print_cell().

GPUCACell::theOuterNeighbors = &cellNeighbors[0]

Definition at line 62 of file GPUCACell.h.

Referenced by __attribute__().

GPUCACell::theStatus_ = 0

Definition at line 54 of file GPUCACell.h.

__device__ float float float float float const float const float GPUCACell::thetaCut
Initial value:
{
float radius_diff = std::abs(r1 - ro)

Definition at line 176 of file GPUCACell.h.

GPUCACell::theTracks = &cellTracks[0]

Definition at line 63 of file GPUCACell.h.

Referenced by __attribute__().

__device__ float GPUCACell::z1

Definition at line 176 of file GPUCACell.h.

Referenced by check_alignment().

__device__ float float float GPUCACell::zi

Definition at line 176 of file GPUCACell.h.

Referenced by check_alignment().

__device__ float float float float float GPUCACell::zo

Definition at line 176 of file GPUCACell.h.

Referenced by check_alignment().

Generated for CMSSW Reference Manual by doxygen 1.8.5