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 = invalidHitId
	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	invalidHitId = std::numeric_limits<caConstants::hindex_type>::max()
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 28 of file GPUCACell.h.

using GPUCACell::CellNeighborsVector = caConstants::CellNeighborsVector

Definition at line 30 of file GPUCACell.h.

using GPUCACell::CellTracks = caConstants::CellTracks

Definition at line 29 of file GPUCACell.h.

using GPUCACell::CellTracksVector = caConstants::CellTracksVector

Definition at line 31 of file GPUCACell.h.

using GPUCACell::hindex_type = Hits::hindex_type

Definition at line 34 of file GPUCACell.h.

using GPUCACell::HitContainer = pixelTrack::HitContainer

Definition at line 38 of file GPUCACell.h.

using GPUCACell::Hits = TrackingRecHit2DSOAView

Definition at line 33 of file GPUCACell.h.

using GPUCACell::OuterHitOfCell = caConstants::OuterHitOfCell

Definition at line 27 of file GPUCACell.h.

using GPUCACell::OuterHitOfCellContainer = caConstants::OuterHitOfCellContainer

Definition at line 26 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 39 of file GPUCACell.h.

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

Definition at line 36 of file GPUCACell.h.

Member Enumeration Documentation

enum GPUCACell::StatusBit : uint16_t

strong

Enumerator
kUsed
kInTrack
kKilled

Definition at line 42 of file GPUCACell.h.

: 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 69 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.

                                                                                                              {
    // use smart cache
    if (outerNeighbors().empty()) {
      auto i = cellNeighbors.extend();  // maybe wasted....
      if (i > 0) {
        cellNeighbors[i].reset();
        __threadfence();
#ifdef __CUDACC__
        auto zero = (PtrAsInt)(&cellNeighbors[0]);
        atomicCAS((PtrAsInt*)(&theOuterNeighbors),
                  zero,
                  (PtrAsInt)(&cellNeighbors[i]));  // if fails we cannot give "i" back...
#else
        theOuterNeighbors = &cellNeighbors[i];
#endif
      } else
        return -1;
    }
    __threadfence();
    return outerNeighbors().push_back(t);
  }

__device__ GPUCACell::__attribute__ ( (always_inline)  )

inline

Definition at line 91 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.

                                                                                             {
    if (tracks().empty()) {
      auto i = cellTracks.extend();  // maybe wasted....
      if (i > 0) {
        cellTracks[i].reset();
        __threadfence();
#ifdef __CUDACC__
        auto zero = (PtrAsInt)(&cellTracks[0]);
        atomicCAS((PtrAsInt*)(&theTracks), zero, (PtrAsInt)(&cellTracks[i]));  // if fails we cannot give "i" back...
#else
        theTracks = &cellTracks[i];
#endif
      } else
        return -1;
    }
    __threadfence();
    return tracks().push_back(t);
  }

__device__ GPUCACell::__attribute__ ( (always_inline)  )

inline

Definition at line 110 of file GPUCACell.h.

References theTracks.

{ return *theTracks; }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 111 of file GPUCACell.h.

References theTracks.

{ return *theTracks; }

__device__ GPUCACell::__attribute__ ( (always_inline)  )

inline

Definition at line 112 of file GPUCACell.h.

References theOuterNeighbors.

{ return *theOuterNeighbors; }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 113 of file GPUCACell.h.

References theOuterNeighbors.

{ return *theOuterNeighbors; }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 114 of file GPUCACell.h.

{ return hh.xGlobal(theInnerHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 115 of file GPUCACell.h.

References theOuterHitId.

{ return hh.xGlobal(theOuterHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 116 of file GPUCACell.h.

{ return hh.yGlobal(theInnerHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 117 of file GPUCACell.h.

References theOuterHitId.

{ return hh.yGlobal(theOuterHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 118 of file GPUCACell.h.

References theInnerZ.

{ return theInnerZ; }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 120 of file GPUCACell.h.

References theOuterHitId.

{ return hh.zGlobal(theOuterHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 121 of file GPUCACell.h.

References theInnerR.

{ return theInnerR; }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 123 of file GPUCACell.h.

References theOuterHitId.

{ return hh.rGlobal(theOuterHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 125 of file GPUCACell.h.

{ return hh.iphi(theInnerHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 126 of file GPUCACell.h.

References theOuterHitId.

{ return hh.iphi(theOuterHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 128 of file GPUCACell.h.

{ return hh.detectorIndex(theInnerHitId); }

__device__ GPUCACell::__attribute__ ( (always_inline)  ) const

inline

Definition at line 129 of file GPUCACell.h.

References theOuterHitId.

{ return hh.detectorIndex(theOuterHitId); }

__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 141 of file GPUCACell.h.

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

                                                                        {
    // detIndex of the layerStart for the Phase1 Pixel Detector:
    // [BPX1, BPX2, BPX3, BPX4,  FP1,  FP2,  FP3,  FN1,  FN2,  FN3, LAST_VALID]
    // [   0,   96,  320,  672, 1184, 1296, 1408, 1520, 1632, 1744,       1856]
    auto ri = inner_r(hh);
    auto zi = inner_z(hh);

    auto ro = outer_r(hh);
    auto zo = outer_z(hh);

    auto r1 = otherCell.inner_r(hh);
    auto z1 = otherCell.inner_z(hh);
    auto isBarrel = otherCell.outer_detIndex(hh) < caConstants::last_barrel_detIndex;
    bool aligned = areAlignedRZ(r1,
                                z1,
                                ri,
                                zi,
                                ro,
                                zo,
                                ptmin,
                                isBarrel ? caThetaCutBarrel : caThetaCutForward);  // 2.f*thetaCut); // FIXME tune cuts
    return (aligned && dcaCut(hh,
                              otherCell,
                              otherCell.inner_detIndex(hh) < caConstants::last_bpix1_detIndex ? dcaCutInnerTriplet
                                                                                              : dcaCutOuterTriplet,
                              hardCurvCut));  // FIXME tune cuts
  }

constexpr unsigned int GPUCACell::inner_hit_id ( ) const

inline

Definition at line 131 of file GPUCACell.h.

{ return theInnerHitId; }

constexpr unsigned int GPUCACell::outer_hit_id ( ) const

inline

Definition at line 132 of file GPUCACell.h.

References theOuterHitId.

{ return theOuterHitId; }

__device__ void GPUCACell::print_cell ( ) const

inline

Definition at line 134 of file GPUCACell.h.

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

                                     {
    printf("printing cell: on layerPair: %d, innerHitId: %d, outerHitId: %d \n",
           theLayerPairId_,
           theInnerHitId,
           theOuterHitId);
  }

Member Data Documentation

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

static

Definition at line 40 of file GPUCACell.h.

__device__ CellTracksVector& GPUCACell::cellTracks

Definition at line 47 of file GPUCACell.h.

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

Definition at line 179 of file GPUCACell.h.

__device__ CellTracksVector Hits const& GPUCACell::hh

Definition at line 47 of file GPUCACell.h.

__device__ CellTracksVector Hits const int hindex_type GPUCACell::innerHitId

Definition at line 47 of file GPUCACell.h.

constexpr auto GPUCACell::invalidHitId = std::numeric_limits<caConstants::hindex_type>::max()

static

Definition at line 25 of file GPUCACell.h.

__device__ CellTracksVector Hits const int GPUCACell::layerPairId

Definition at line 47 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 51 of file GPUCACell.h.

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

Definition at line 181 of file GPUCACell.h.

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

Definition at line 177 of file GPUCACell.h.

__device__ float float GPUCACell::ri

Definition at line 177 of file GPUCACell.h.

Referenced by check_alignment().

__device__ float float float float GPUCACell::ro

Definition at line 177 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 184 of file GPUCACell.h.

GPUCACell::theFishboneId = invalidHitId

Definition at line 56 of file GPUCACell.h.

GPUCACell::theInnerR = hh.rGlobal(innerHitId)

Definition at line 60 of file GPUCACell.h.

Referenced by __attribute__().

GPUCACell::theInnerZ = hh.zGlobal(innerHitId)

Definition at line 59 of file GPUCACell.h.

Referenced by __attribute__().

GPUCACell::theLayerPairId_ = layerPairId

Definition at line 54 of file GPUCACell.h.

Referenced by print_cell().

GPUCACell::theOuterHitId = outerHitId

Definition at line 53 of file GPUCACell.h.

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

GPUCACell::theOuterNeighbors = &cellNeighbors[0]

Definition at line 63 of file GPUCACell.h.

Referenced by __attribute__().

GPUCACell::theStatus_ = 0

Definition at line 55 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 177 of file GPUCACell.h.

GPUCACell::theTracks = &cellTracks[0]

Definition at line 64 of file GPUCACell.h.

Referenced by __attribute__().

__device__ float GPUCACell::z1

Definition at line 177 of file GPUCACell.h.

Referenced by check_alignment().

__device__ float float float GPUCACell::zi

Definition at line 177 of file GPUCACell.h.

Referenced by check_alignment().

__device__ float float float float float GPUCACell::zo

Definition at line 177 of file GPUCACell.h.

Referenced by check_alignment().