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

#include <CAHitNtupletGeneratorKernels.h>

Public Types

using Counters = cAHitNtupletGenerator::Counters
 
using HitContainer = pixelTrack::HitContainer
 
using HitsOnCPU = TrackingRecHit2DHeterogeneous< Traits >
 
using HitsOnGPU = TrackingRecHit2DSOAView
 
using HitsView = TrackingRecHit2DSOAView
 
using HitToTuple = caConstants::HitToTuple
 
using Params = cAHitNtupletGenerator::Params
 
using Quality = pixelTrack::Quality
 
using QualityCuts = cAHitNtupletGenerator::QualityCuts
 
using TkSoA = pixelTrack::TrackSoA
 
using Traits = TTraits
 
using TupleMultiplicity = caConstants::TupleMultiplicity
 
template<typename T >
using unique_ptr = typename Traits::template unique_ptr< T >
 

Public Member Functions

void allocateOnGPU (int32_t nHits, cudaStream_t stream)
 
void buildDoublets (HitsOnCPU const &hh, cudaStream_t stream)
 
 CAHitNtupletGeneratorKernels (Params const &params)
 
void classifyTuples (HitsOnCPU const &hh, TkSoA *tuples_d, cudaStream_t cudaStream)
 
void cleanup (cudaStream_t cudaStream)
 
void fillHitDetIndices (HitsView const *hv, TkSoA *tuples_d, cudaStream_t cudaStream)
 
void launchKernels (HitsOnCPU const &hh, TkSoA *tuples_d, cudaStream_t cudaStream)
 
void setCounters (Counters *counters)
 
TupleMultiplicity const * tupleMultiplicity () const
 
 ~CAHitNtupletGeneratorKernels ()=default
 

Static Public Member Functions

static void printCounters (Counters const *counters)
 

Private Member Functions

uint32_t nDoubletBlocks (uint32_t blockSize)
 Compute the number of doublet blocks for block size. More...
 
uint32_t nQuadrupletBlocks (uint32_t blockSize)
 Compute the number of quadruplet blocks for block size. More...
 

Private Attributes

unique_ptr< unsigned char[]> cellStorage_
 
Counters * counters_ = nullptr
 
unique_ptr< HitToTuple > device_hitToTuple_
 
cms::cuda::AtomicPairCounter * device_hitToTuple_apc_ = nullptr
 
unique_ptr< HitToTuple::Counter[]> device_hitToTupleStorage_
 
cms::cuda::AtomicPairCounter * device_hitTuple_apc_ = nullptr
 
unique_ptr< GPUCACell::OuterHitOfCell[]> device_isOuterHitOfCell_
 
uint32_t * device_nCells_ = nullptr
 
unique_ptr< cms::cuda::AtomicPairCounter::c_type[]> device_storage_
 
unique_ptr< caConstants::CellNeighborsVector > device_theCellNeighbors_
 
caConstants::CellNeighbors * device_theCellNeighborsContainer_
 
unique_ptr< GPUCACell[]> device_theCells_
 
unique_ptr< caConstants::CellTracksVector > device_theCellTracks_
 
caConstants::CellTracks * device_theCellTracksContainer_
 
unique_ptr< TupleMultiplicity > device_tupleMultiplicity_
 
HitToTuple::View hitToTupleView_
 
Params const & params_
 
const uint32_t paramsMaxDoubletes3Quarters_
 Intermediate result avoiding repeated computations. More...
 

Detailed Description

template<typename TTraits>
class CAHitNtupletGeneratorKernels< TTraits >

Definition at line 156 of file CAHitNtupletGeneratorKernels.h.

Member Typedef Documentation

◆ Counters

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::Counters = cAHitNtupletGenerator::Counters

Definition at line 162 of file CAHitNtupletGeneratorKernels.h.

◆ HitContainer

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::HitContainer = pixelTrack::HitContainer

Definition at line 176 of file CAHitNtupletGeneratorKernels.h.

◆ HitsOnCPU

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::HitsOnCPU = TrackingRecHit2DHeterogeneous<Traits>

Definition at line 169 of file CAHitNtupletGeneratorKernels.h.

◆ HitsOnGPU

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 168 of file CAHitNtupletGeneratorKernels.h.

◆ HitsView

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::HitsView = TrackingRecHit2DSOAView

Definition at line 167 of file CAHitNtupletGeneratorKernels.h.

◆ HitToTuple

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::HitToTuple = caConstants::HitToTuple

Definition at line 171 of file CAHitNtupletGeneratorKernels.h.

◆ Params

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::Params = cAHitNtupletGenerator::Params

Definition at line 161 of file CAHitNtupletGeneratorKernels.h.

◆ Quality

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::Quality = pixelTrack::Quality

Definition at line 174 of file CAHitNtupletGeneratorKernels.h.

◆ QualityCuts

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::QualityCuts = cAHitNtupletGenerator::QualityCuts

Definition at line 160 of file CAHitNtupletGeneratorKernels.h.

◆ TkSoA

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::TkSoA = pixelTrack::TrackSoA

Definition at line 175 of file CAHitNtupletGeneratorKernels.h.

◆ Traits

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::Traits = TTraits

Definition at line 158 of file CAHitNtupletGeneratorKernels.h.

◆ TupleMultiplicity

template<typename TTraits >
using CAHitNtupletGeneratorKernels< TTraits >::TupleMultiplicity = caConstants::TupleMultiplicity

Definition at line 172 of file CAHitNtupletGeneratorKernels.h.

◆ unique_ptr

template<typename TTraits >
template<typename T >
using CAHitNtupletGeneratorKernels< TTraits >::unique_ptr = typename Traits::template unique_ptr<T>

Definition at line 165 of file CAHitNtupletGeneratorKernels.h.

Constructor & Destructor Documentation

◆ CAHitNtupletGeneratorKernels()

template<typename TTraits >
CAHitNtupletGeneratorKernels< TTraits >::CAHitNtupletGeneratorKernels ( Params const &  params)
inline

Definition at line 178 of file CAHitNtupletGeneratorKernels.h.

179  : params_(params), paramsMaxDoubletes3Quarters_(3 * params.maxNumberOfDoublets_ / 4) {}

◆ ~CAHitNtupletGeneratorKernels()

template<typename TTraits >
CAHitNtupletGeneratorKernels< TTraits >::~CAHitNtupletGeneratorKernels ( )
default

Member Function Documentation

◆ allocateOnGPU()

void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU ( int32_t  nHits,
cudaStream_t  stream 
)

Definition at line 9 of file CAHitNtupletGeneratorKernelsAlloc.cc.

9  {
10 #endif
11  // ALLOCATIONS FOR THE INTERMEDIATE RESULTS (STAYS ON WORKER)
14 
15  device_theCellNeighbors_ = Traits::template make_unique<caConstants::CellNeighborsVector>(stream);
16  device_theCellTracks_ = Traits::template make_unique<caConstants::CellTracksVector>(stream);
17 
18 #ifdef GPU_DEBUG
19  std::cout << "Allocation for tuple building. N hits " << nHits << std::endl;
20 #endif
21 
22  nHits++; // storage requires one more counter;
23  assert(nHits > 0);
24  device_hitToTuple_ = Traits::template make_unique<HitToTuple>(stream);
25  device_hitToTupleStorage_ = Traits::template make_unique<HitToTuple::Counter[]>(nHits, stream);
26  hitToTupleView_.assoc = device_hitToTuple_.get();
27  hitToTupleView_.offStorage = device_hitToTupleStorage_.get();
28  hitToTupleView_.offSize = nHits;
29 
30  device_tupleMultiplicity_ = Traits::template make_unique<TupleMultiplicity>(stream);
31 
32  device_storage_ = Traits::template make_unique<cms::cuda::AtomicPairCounter::c_type[]>(3, stream);
33 
34  device_hitTuple_apc_ = (cms::cuda::AtomicPairCounter*)device_storage_.get();
35  device_hitToTuple_apc_ = (cms::cuda::AtomicPairCounter*)device_storage_.get() + 1;
36  device_nCells_ = (uint32_t*)(device_storage_.get() + 2);
37 
38  // FIXME: consider collapsing these 3 in one adhoc kernel
39  if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
40  cudaCheck(cudaMemsetAsync(device_nCells_, 0, sizeof(uint32_t), stream));
41  } else {
42  *device_nCells_ = 0;
43  }
44  cms::cuda::launchZero(device_tupleMultiplicity_.get(), stream);
45  cms::cuda::launchZero(hitToTupleView_, stream); // we may wish to keep it in the edm
46 #ifdef GPU_DEBUG
47  cudaDeviceSynchronize();
48  cudaCheck(cudaGetLastError());
49 #endif
50 }

References cms::cuda::assert(), cms::cuda::OneToManyAssocView< Assoc >::assoc, gather_cfg::cout, cudaCheck, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTupleStorage_, CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_storage_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_, CAHitNtupletGeneratorKernels< TTraits >::hitToTupleView_, nHits, cms::cuda::OneToManyAssocView< Assoc >::offSize, cms::cuda::OneToManyAssocView< Assoc >::offStorage, cms::cuda::stream, svgfig::template(), and relativeConstraints::value.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

◆ buildDoublets()

void CAHitNtupletGeneratorKernelsCPU::buildDoublets ( HitsOnCPU const &  hh,
cudaStream_t  stream 
)

Definition at line 14 of file CAHitNtupletGeneratorKernels.cc.

14  {
15  auto nhits = hh.nHits();
16 
17 #ifdef NTUPLE_DEBUG
18  std::cout << "building Doublets out of " << nhits << " Hits" << std::endl;
19 #endif
20 
21  // use "nhits" to heuristically dimension the workspace
22 
23  // no need to use the Traits allocations, since we know this is being compiled for the CPU
24  //device_isOuterHitOfCell_ = Traits::template make_unique<GPUCACell::OuterHitOfCell[]>(std::max(1U, nhits), stream);
25  device_isOuterHitOfCell_ = std::make_unique<GPUCACell::OuterHitOfCell[]>(std::max(1U, nhits));
26  assert(device_isOuterHitOfCell_.get());
27 
28  auto cellStorageSize = caConstants::maxNumOfActiveDoublets * sizeof(GPUCACell::CellNeighbors) +
29  caConstants::maxNumOfActiveDoublets * sizeof(GPUCACell::CellTracks);
30  // no need to use the Traits allocations, since we know this is being compiled for the CPU
31  //cellStorage_ = Traits::template make_unique<unsigned char[]>(cellStorageSize, stream);
32  cellStorage_ = std::make_unique<unsigned char[]>(cellStorageSize);
33  device_theCellNeighborsContainer_ = (GPUCACell::CellNeighbors *)cellStorage_.get();
34  device_theCellTracksContainer_ = (GPUCACell::CellTracks *)(cellStorage_.get() + caConstants::maxNumOfActiveDoublets *
35  sizeof(GPUCACell::CellNeighbors));
36 
37  gpuPixelDoublets::initDoublets(device_isOuterHitOfCell_.get(),
38  nhits,
39  device_theCellNeighbors_.get(),
40  device_theCellNeighborsContainer_,
41  device_theCellTracks_.get(),
42  device_theCellTracksContainer_);
43 
44  // no need to use the Traits allocations, since we know this is being compiled for the CPU
45  //device_theCells_ = Traits::template make_unique<GPUCACell[]>(params_.maxNumberOfDoublets_, stream);
46  device_theCells_ = std::make_unique<GPUCACell[]>(params_.maxNumberOfDoublets_);
47  if (0 == nhits)
48  return; // protect against empty events
49 
50  // take all layer pairs into account
51  auto nActualPairs = gpuPixelDoublets::nPairs;
52  if (not params_.includeJumpingForwardDoublets_) {
53  // exclude forward "jumping" layer pairs
54  nActualPairs = gpuPixelDoublets::nPairsForTriplets;
55  }
56  if (params_.minHitsPerNtuplet_ > 3) {
57  // for quadruplets, exclude all "jumping" layer pairs
58  nActualPairs = gpuPixelDoublets::nPairsForQuadruplets;
59  }
60 
61  assert(nActualPairs <= gpuPixelDoublets::nPairs);
62  gpuPixelDoublets::getDoubletsFromHisto(device_theCells_.get(),
63  device_nCells_,
64  device_theCellNeighbors_.get(),
65  device_theCellTracks_.get(),
66  hh.view(),
67  device_isOuterHitOfCell_.get(),
68  nActualPairs,
69  params_.idealConditions_,
70  params_.doClusterCut_,
71  params_.doZ0Cut_,
72  params_.doPtCut_,
73  params_.maxNumberOfDoublets_);
74 }

References cms::cuda::assert(), CAHitNtupletGeneratorKernels< TTraits >::cellStorage_, gather_cfg::cout, CAHitNtupletGeneratorKernels< TTraits >::device_isOuterHitOfCell_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighborsContainer_, CAHitNtupletGeneratorKernels< TTraits >::device_theCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracksContainer_, cAHitNtupletGenerator::Params::doClusterCut_, cAHitNtupletGenerator::Params::doPtCut_, cAHitNtupletGenerator::Params::doZ0Cut_, hh, cAHitNtupletGenerator::Params::idealConditions_, cAHitNtupletGenerator::Params::includeJumpingForwardDoublets_, SiStripPI::max, cAHitNtupletGenerator::Params::maxNumberOfDoublets_, caConstants::maxNumOfActiveDoublets, cAHitNtupletGenerator::Params::minHitsPerNtuplet_, gpuPixelDoublets::nActualPairs, nhits, gpuPixelDoublets::nPairs, gpuPixelDoublets::nPairsForQuadruplets, gpuPixelDoublets::nPairsForTriplets, CAHitNtupletGeneratorKernels< TTraits >::params_, and mitigatedMETSequence_cff::U.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

◆ classifyTuples()

void CAHitNtupletGeneratorKernelsCPU::classifyTuples ( HitsOnCPU const &  hh,
TkSoA *  tuples_d,
cudaStream_t  cudaStream 
)

Definition at line 156 of file CAHitNtupletGeneratorKernels.cc.

156  {
157  auto const *tuples_d = &tracks_d->hitIndices;
158  auto *quality_d = tracks_d->qualityData();
159 
160  // classify tracks based on kinematics
161  kernel_classifyTracks(tuples_d, tracks_d, params_.cuts_, quality_d);
162 
163  if (params_.lateFishbone_) {
164  // apply fishbone cleaning to good tracks
165  kernel_fishboneCleaner(device_theCells_.get(), device_nCells_, quality_d);
166  }
167 
168  // remove duplicates (tracks that share a doublet)
169  kernel_fastDuplicateRemover(device_theCells_.get(), device_nCells_, tuples_d, tracks_d, params_.dupPassThrough_);
170 
171  // fill hit->track "map"
172  if (params_.doSharedHitCut_ || params_.doStats_) {
173  kernel_countHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
174  cms::cuda::launchFinalize(hitToTupleView_, cudaStream);
175  kernel_fillHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
176  }
177 
178  // remove duplicates (tracks that share at least one hit)
179  if (params_.doSharedHitCut_) {
180  kernel_rejectDuplicate(hh.view(),
181  tuples_d,
182  tracks_d,
183  quality_d,
184  params_.minHitsForSharingCut_,
185  params_.dupPassThrough_,
186  device_hitToTuple_.get());
187 
188  kernel_sharedHitCleaner(hh.view(),
189  tuples_d,
190  tracks_d,
191  quality_d,
192  params_.minHitsForSharingCut_,
193  params_.dupPassThrough_,
194  device_hitToTuple_.get());
195  if (params_.useSimpleTripletCleaner_) {
196  kernel_simpleTripletCleaner(hh.view(),
197  tuples_d,
198  tracks_d,
199  quality_d,
200  params_.minHitsForSharingCut_,
201  params_.dupPassThrough_,
202  device_hitToTuple_.get());
203  } else {
204  kernel_tripletCleaner(hh.view(),
205  tuples_d,
206  tracks_d,
207  quality_d,
208  params_.minHitsForSharingCut_,
209  params_.dupPassThrough_,
210  device_hitToTuple_.get());
211  }
212  }
213  if (params_.doStats_) {
214  // counters (add flag???)
215  kernel_doStatsForHitInTracks(device_hitToTuple_.get(), counters_);
216  kernel_doStatsForTracks(tuples_d, quality_d, counters_);
217  }
218 
219 #ifdef DUMP_GPU_TK_TUPLES
220  static std::atomic<int> iev(0);
221  ++iev;
222  kernel_print_found_ntuplets(hh.view(), tuples_d, tracks_d, quality_d, device_hitToTuple_.get(), 100, iev);
223 #endif
224 }

References CAHitNtupletGeneratorKernels< TTraits >::counters_, cAHitNtupletGenerator::Params::cuts_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCells_, cAHitNtupletGenerator::Params::doSharedHitCut_, cAHitNtupletGenerator::Params::doStats_, cAHitNtupletGenerator::Params::dupPassThrough_, hh, TrackSoAHeterogeneousT< S >::hitIndices, CAHitNtupletGeneratorKernels< TTraits >::hitToTupleView_, iev, cAHitNtupletGenerator::Params::lateFishbone_, cAHitNtupletGenerator::Params::minHitsForSharingCut_, CAHitNtupletGeneratorKernels< TTraits >::params_, TrackSoAHeterogeneousT< S >::qualityData(), and cAHitNtupletGenerator::Params::useSimpleTripletCleaner_.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

◆ cleanup()

template<typename TTraits >
void CAHitNtupletGeneratorKernels< TTraits >::cleanup ( cudaStream_t  cudaStream)

◆ fillHitDetIndices()

void CAHitNtupletGeneratorKernelsCPU::fillHitDetIndices ( HitsView const *  hv,
TkSoA *  tuples_d,
cudaStream_t  cudaStream 
)

Definition at line 9 of file CAHitNtupletGeneratorKernels.cc.

9  {
10  kernel_fillHitDetIndices(&tracks_d->hitIndices, hv, &tracks_d->detIndices);
11 }

References TrackSoAHeterogeneousT< S >::detIndices, and TrackSoAHeterogeneousT< S >::hitIndices.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

◆ launchKernels()

void CAHitNtupletGeneratorKernelsCPU::launchKernels ( HitsOnCPU const &  hh,
TkSoA *  tuples_d,
cudaStream_t  cudaStream 
)

Definition at line 77 of file CAHitNtupletGeneratorKernels.cc.

77  {
78  auto *tuples_d = &tracks_d->hitIndices;
79  auto *quality_d = tracks_d->qualityData();
80 
81  assert(tuples_d && quality_d);
82 
83  // zero tuples
84  cms::cuda::launchZero(tuples_d, cudaStream);
85 
86  auto nhits = hh.nHits();
87 
88  // std::cout << "N hits " << nhits << std::endl;
89  // if (nhits<2) std::cout << "too few hits " << nhits << std::endl;
90 
91  //
92  // applying conbinatoric cleaning such as fishbone at this stage is too expensive
93  //
94 
95  kernel_connect(device_hitTuple_apc_,
96  device_hitToTuple_apc_, // needed only to be reset, ready for next kernel
97  hh.view(),
98  device_theCells_.get(),
99  device_nCells_,
100  device_theCellNeighbors_.get(),
101  device_isOuterHitOfCell_.get(),
102  params_.hardCurvCut_,
103  params_.ptmin_,
104  params_.CAThetaCutBarrel_,
105  params_.CAThetaCutForward_,
106  params_.dcaCutInnerTriplet_,
107  params_.dcaCutOuterTriplet_);
108 
109  if (nhits > 1 && params_.earlyFishbone_) {
110  gpuPixelDoublets::fishbone(
111  hh.view(), device_theCells_.get(), device_nCells_, device_isOuterHitOfCell_.get(), nhits, false);
112  }
113 
114  kernel_find_ntuplets(hh.view(),
115  device_theCells_.get(),
116  device_nCells_,
117  device_theCellTracks_.get(),
118  tuples_d,
119  device_hitTuple_apc_,
120  quality_d,
121  params_.minHitsPerNtuplet_);
122  if (params_.doStats_)
123  kernel_mark_used(hh.view(), device_theCells_.get(), device_nCells_);
124 
125  cms::cuda::finalizeBulk(device_hitTuple_apc_, tuples_d);
126 
127  // remove duplicates (tracks that share a doublet)
128  kernel_earlyDuplicateRemover(device_theCells_.get(), device_nCells_, tuples_d, quality_d, params_.dupPassThrough_);
129 
130  kernel_countMultiplicity(tuples_d, quality_d, device_tupleMultiplicity_.get());
131  cms::cuda::launchFinalize(device_tupleMultiplicity_.get(), cudaStream);
132  kernel_fillMultiplicity(tuples_d, quality_d, device_tupleMultiplicity_.get());
133 
134  if (nhits > 1 && params_.lateFishbone_) {
135  gpuPixelDoublets::fishbone(
136  hh.view(), device_theCells_.get(), device_nCells_, device_isOuterHitOfCell_.get(), nhits, true);
137  }
138 
139  if (params_.doStats_) {
140  kernel_checkOverflows(tuples_d,
141  device_tupleMultiplicity_.get(),
142  device_hitToTuple_.get(),
143  device_hitTuple_apc_,
144  device_theCells_.get(),
145  device_nCells_,
146  device_theCellNeighbors_.get(),
147  device_theCellTracks_.get(),
148  device_isOuterHitOfCell_.get(),
149  nhits,
150  params_.maxNumberOfDoublets_,
151  counters_);
152  }
153 }

References cms::cuda::assert(), cAHitNtupletGenerator::Params::CAThetaCutBarrel_, cAHitNtupletGenerator::Params::CAThetaCutForward_, CAHitNtupletGeneratorKernels< TTraits >::counters_, cAHitNtupletGenerator::Params::dcaCutInnerTriplet_, cAHitNtupletGenerator::Params::dcaCutOuterTriplet_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_isOuterHitOfCell_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_, cAHitNtupletGenerator::Params::doStats_, cAHitNtupletGenerator::Params::dupPassThrough_, cAHitNtupletGenerator::Params::earlyFishbone_, cAHitNtupletGenerator::Params::hardCurvCut_, hh, TrackSoAHeterogeneousT< S >::hitIndices, cAHitNtupletGenerator::Params::lateFishbone_, cAHitNtupletGenerator::Params::maxNumberOfDoublets_, cAHitNtupletGenerator::Params::minHitsPerNtuplet_, nhits, CAHitNtupletGeneratorKernels< TTraits >::params_, cAHitNtupletGenerator::Params::ptmin_, and TrackSoAHeterogeneousT< S >::qualityData().

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

◆ nDoubletBlocks()

template<typename TTraits >
uint32_t CAHitNtupletGeneratorKernels< TTraits >::nDoubletBlocks ( uint32_t  blockSize)
inlineprivate

Compute the number of doublet blocks for block size.

Definition at line 227 of file CAHitNtupletGeneratorKernels.h.

227  {
228  // We want (3 * params_.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize, but first part is pre-computed.
229  return (paramsMaxDoubletes3Quarters_ + blockSize - 1) / blockSize;
230  }

References CAHitNtupletGeneratorKernels< TTraits >::paramsMaxDoubletes3Quarters_.

◆ nQuadrupletBlocks()

template<typename TTraits >
uint32_t CAHitNtupletGeneratorKernels< TTraits >::nQuadrupletBlocks ( uint32_t  blockSize)
inlineprivate

Compute the number of quadruplet blocks for block size.

Definition at line 233 of file CAHitNtupletGeneratorKernels.h.

233  {
234  // caConstants::maxNumberOfQuadruplets is a constexpr, so the compiler will pre compute the 3*max/4
235  return (3 * caConstants::maxNumberOfQuadruplets / 4 + blockSize - 1) / blockSize;
236  }

References caConstants::maxNumberOfQuadruplets.

◆ printCounters()

void CAHitNtupletGeneratorKernelsCPU::printCounters ( Counters const *  counters)
static

Definition at line 4 of file CAHitNtupletGeneratorKernels.cc.

4  {
5  kernel_printCounters(counters);
6 }

References counters.

Referenced by CAHitNtupletGeneratorOnGPU::~CAHitNtupletGeneratorOnGPU().

◆ setCounters()

template<typename TTraits >
void CAHitNtupletGeneratorKernels< TTraits >::setCounters ( Counters *  counters)
inline

Definition at line 195 of file CAHitNtupletGeneratorKernels.h.

195 { counters_ = counters; }

References counters, and CAHitNtupletGeneratorKernels< TTraits >::counters_.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

◆ tupleMultiplicity()

template<typename TTraits >
TupleMultiplicity const* CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity ( ) const
inline

Definition at line 182 of file CAHitNtupletGeneratorKernels.h.

182 { return device_tupleMultiplicity_.get(); }

References CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

Member Data Documentation

◆ cellStorage_

template<typename TTraits >
unique_ptr<unsigned char[]> CAHitNtupletGeneratorKernels< TTraits >::cellStorage_
private

Definition at line 201 of file CAHitNtupletGeneratorKernels.h.

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

◆ counters_

template<typename TTraits >
Counters* CAHitNtupletGeneratorKernels< TTraits >::counters_ = nullptr
private

Definition at line 198 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), CAHitNtupletGeneratorKernels< TTraits >::launchKernels(), and CAHitNtupletGeneratorKernels< TTraits >::setCounters().

◆ device_hitToTuple_

template<typename TTraits >
unique_ptr<HitToTuple> CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_
private

Definition at line 211 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_hitToTuple_apc_

template<typename TTraits >
cms::cuda::AtomicPairCounter* CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_apc_ = nullptr
private

Definition at line 215 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_hitToTupleStorage_

template<typename TTraits >
unique_ptr<HitToTuple::Counter[]> CAHitNtupletGeneratorKernels< TTraits >::device_hitToTupleStorage_
private

Definition at line 212 of file CAHitNtupletGeneratorKernels.h.

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

◆ device_hitTuple_apc_

template<typename TTraits >
cms::cuda::AtomicPairCounter* CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_ = nullptr
private

Definition at line 217 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_isOuterHitOfCell_

template<typename TTraits >
unique_ptr<GPUCACell::OuterHitOfCell[]> CAHitNtupletGeneratorKernels< TTraits >::device_isOuterHitOfCell_
private

Definition at line 208 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_nCells_

template<typename TTraits >
uint32_t* CAHitNtupletGeneratorKernels< TTraits >::device_nCells_ = nullptr
private

Definition at line 209 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_storage_

template<typename TTraits >
unique_ptr<cms::cuda::AtomicPairCounter::c_type[]> CAHitNtupletGeneratorKernels< TTraits >::device_storage_
private

Definition at line 221 of file CAHitNtupletGeneratorKernels.h.

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

◆ device_theCellNeighbors_

template<typename TTraits >
unique_ptr<caConstants::CellNeighborsVector> CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_
private

Definition at line 202 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_theCellNeighborsContainer_

template<typename TTraits >
caConstants::CellNeighbors* CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighborsContainer_
private

Definition at line 203 of file CAHitNtupletGeneratorKernels.h.

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

◆ device_theCells_

template<typename TTraits >
unique_ptr<GPUCACell[]> CAHitNtupletGeneratorKernels< TTraits >::device_theCells_
private

Definition at line 207 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_theCellTracks_

template<typename TTraits >
unique_ptr<caConstants::CellTracksVector> CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_
private

Definition at line 204 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ device_theCellTracksContainer_

template<typename TTraits >
caConstants::CellTracks* CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracksContainer_
private

Definition at line 205 of file CAHitNtupletGeneratorKernels.h.

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

◆ device_tupleMultiplicity_

template<typename TTraits >
unique_ptr<TupleMultiplicity> CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_
private

Definition at line 219 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::launchKernels(), and CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity().

◆ hitToTupleView_

template<typename TTraits >
HitToTuple::View CAHitNtupletGeneratorKernels< TTraits >::hitToTupleView_
private

Definition at line 213 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), and CAHitNtupletGeneratorKernels< TTraits >::classifyTuples().

◆ params_

template<typename TTraits >
Params const& CAHitNtupletGeneratorKernels< TTraits >::params_
private

Definition at line 223 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

◆ paramsMaxDoubletes3Quarters_

template<typename TTraits >
const uint32_t CAHitNtupletGeneratorKernels< TTraits >::paramsMaxDoubletes3Quarters_
private

Intermediate result avoiding repeated computations.

Definition at line 225 of file CAHitNtupletGeneratorKernels.h.

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

cAHitNtupletGenerator::Params::earlyFishbone_
const bool earlyFishbone_
Definition: CAHitNtupletGeneratorKernels.h:114
cms::cuda::OneToManyAssocView::assoc
Assoc * assoc
Definition: OneToManyAssoc.h:27
CAHitNtupletGeneratorKernels::device_theCellTracksContainer_
caConstants::CellTracks * device_theCellTracksContainer_
Definition: CAHitNtupletGeneratorKernels.h:205
cms::cuda::OneToManyAssocView::offSize
int32_t offSize
Definition: OneToManyAssoc.h:30
cms::cuda::OneToManyAssocView::offStorage
Counter * offStorage
Definition: OneToManyAssoc.h:28
CAHitNtupletGeneratorKernels::device_hitTuple_apc_
cms::cuda::AtomicPairCounter * device_hitTuple_apc_
Definition: CAHitNtupletGeneratorKernels.h:217
iev
const HitContainer *__restrict__ const TkSoA *__restrict__ const Quality *__restrict__ const CAHitNtupletGeneratorKernelsGPU::HitToTuple *__restrict__ int32_t int iev
Definition: CAHitNtupletGeneratorKernelsImpl.h:862
CalibrationSummaryClient_cfi.params
params
Definition: CalibrationSummaryClient_cfi.py:14
cAHitNtupletGenerator::Params::minHitsForSharingCut_
const uint16_t minHitsForSharingCut_
Definition: CAHitNtupletGeneratorKernels.h:110
GPUCACell::CellNeighbors
caConstants::CellNeighbors CellNeighbors
Definition: GPUCACell.h:26
gather_cfg.cout
cout
Definition: gather_cfg.py:144
CAHitNtupletGeneratorKernels::params_
Params const & params_
Definition: CAHitNtupletGeneratorKernels.h:223
cms::cuda::stream
uint32_t const T *__restrict__ const uint32_t *__restrict__ int32_t int Histo::index_type cudaStream_t stream
Definition: HistoContainer.h:51
cms::cuda::assert
assert(be >=bs)
CAHitNtupletGeneratorKernels::device_theCellTracks_
unique_ptr< caConstants::CellTracksVector > device_theCellTracks_
Definition: CAHitNtupletGeneratorKernels.h:204
cAHitNtupletGenerator::Params::dcaCutInnerTriplet_
const float dcaCutInnerTriplet_
Definition: CAHitNtupletGeneratorKernels.h:128
CAHitNtupletGeneratorKernels::device_theCellNeighbors_
unique_ptr< caConstants::CellNeighborsVector > device_theCellNeighbors_
Definition: CAHitNtupletGeneratorKernels.h:202
CAHitNtupletGeneratorKernels::device_theCells_
unique_ptr< GPUCACell[]> device_theCells_
Definition: CAHitNtupletGeneratorKernels.h:207
caConstants::maxNumberOfQuadruplets
constexpr uint32_t maxNumberOfQuadruplets
Definition: CAConstants.h:41
CAHitNtupletGeneratorKernels::device_nCells_
uint32_t * device_nCells_
Definition: CAHitNtupletGeneratorKernels.h:209
CAHitNtupletGeneratorKernels::device_tupleMultiplicity_
unique_ptr< TupleMultiplicity > device_tupleMultiplicity_
Definition: CAHitNtupletGeneratorKernels.h:219
CAHitNtupletGeneratorKernels::device_hitToTuple_apc_
cms::cuda::AtomicPairCounter * device_hitToTuple_apc_
Definition: CAHitNtupletGeneratorKernels.h:215
cAHitNtupletGenerator::Params::idealConditions_
const bool idealConditions_
Definition: CAHitNtupletGeneratorKernels.h:116
cAHitNtupletGenerator::Params::lateFishbone_
const bool lateFishbone_
Definition: CAHitNtupletGeneratorKernels.h:115
cAHitNtupletGenerator::Params::dupPassThrough_
const bool dupPassThrough_
Definition: CAHitNtupletGeneratorKernels.h:122
cAHitNtupletGenerator::Params::doSharedHitCut_
const bool doSharedHitCut_
Definition: CAHitNtupletGeneratorKernels.h:121
counters
const caConstants::TupleMultiplicity const CAHitNtupletGeneratorKernelsGPU::HitToTuple cms::cuda::AtomicPairCounter const GPUCACell *__restrict__ const uint32_t *__restrict__ const gpuPixelDoublets::CellNeighborsVector const gpuPixelDoublets::CellTracksVector const GPUCACell::OuterHitOfCell *__restrict__ int32_t uint32_t CAHitNtupletGeneratorKernelsGPU::Counters * counters
Definition: CAHitNtupletGeneratorKernelsImpl.h:53
cAHitNtupletGenerator::Params::maxNumberOfDoublets_
const uint32_t maxNumberOfDoublets_
Definition: CAHitNtupletGeneratorKernels.h:109
CAHitNtupletGeneratorKernels::device_theCellNeighborsContainer_
caConstants::CellNeighbors * device_theCellNeighborsContainer_
Definition: CAHitNtupletGeneratorKernels.h:203
cAHitNtupletGenerator::Params::minHitsPerNtuplet_
const uint32_t minHitsPerNtuplet_
Definition: CAHitNtupletGeneratorKernels.h:108
nHits
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t nHits
Definition: BrokenLineFitOnGPU.h:27
cAHitNtupletGenerator::Params::useSimpleTripletCleaner_
const bool useSimpleTripletCleaner_
Definition: CAHitNtupletGeneratorKernels.h:123
CAHitNtupletGeneratorKernels::device_hitToTuple_
unique_ptr< HitToTuple > device_hitToTuple_
Definition: CAHitNtupletGeneratorKernels.h:211
cAHitNtupletGenerator::Params::hardCurvCut_
const float hardCurvCut_
Definition: CAHitNtupletGeneratorKernels.h:127
cAHitNtupletGenerator::Params::dcaCutOuterTriplet_
const float dcaCutOuterTriplet_
Definition: CAHitNtupletGeneratorKernels.h:129
caConstants::maxNumOfActiveDoublets
constexpr uint32_t maxNumOfActiveDoublets
Definition: CAConstants.h:40
nhits
Definition: HIMultiTrackSelector.h:42
gpuPixelDoublets::nPairsForQuadruplets
constexpr int nPairsForQuadruplets
Definition: gpuPixelDoublets.h:10
mitigatedMETSequence_cff.U
U
Definition: mitigatedMETSequence_cff.py:36
gpuPixelDoublets::nActualPairs
uint32_t CellNeighborsVector CellTracksVector TrackingRecHit2DSOAView const *__restrict__ GPUCACell::OuterHitOfCell int nActualPairs
Definition: gpuPixelDoublets.h:99
CAHitNtupletGeneratorKernels::device_isOuterHitOfCell_
unique_ptr< GPUCACell::OuterHitOfCell[]> device_isOuterHitOfCell_
Definition: CAHitNtupletGeneratorKernels.h:208
CAHitNtupletGeneratorKernels::device_storage_
unique_ptr< cms::cuda::AtomicPairCounter::c_type[]> device_storage_
Definition: CAHitNtupletGeneratorKernels.h:221
hh
const auto & hh
Definition: CAHitNtupletGeneratorKernelsImpl.h:552
cms::cuda::AtomicPairCounter
Definition: AtomicPairCounter.h:11
SiStripPI::max
Definition: SiStripPayloadInspectorHelper.h:169
cAHitNtupletGenerator::Params::cuts_
QualityCuts cuts_
Definition: CAHitNtupletGeneratorKernels.h:132
cAHitNtupletGenerator::Params::doPtCut_
const bool doPtCut_
Definition: CAHitNtupletGeneratorKernels.h:120
cAHitNtupletGenerator::Params::includeJumpingForwardDoublets_
const bool includeJumpingForwardDoublets_
Definition: CAHitNtupletGeneratorKernels.h:113
cAHitNtupletGenerator::Params::CAThetaCutForward_
const float CAThetaCutForward_
Definition: CAHitNtupletGeneratorKernels.h:126
cms::cuda::VecArray
Definition: VecArray.h:14
svgfig.template
def template(fileName, svg, replaceme="REPLACEME")
Definition: svgfig.py:521
cAHitNtupletGenerator::Params::doZ0Cut_
const bool doZ0Cut_
Definition: CAHitNtupletGeneratorKernels.h:119
CAHitNtupletGeneratorKernels::paramsMaxDoubletes3Quarters_
const uint32_t paramsMaxDoubletes3Quarters_
Intermediate result avoiding repeated computations.
Definition: CAHitNtupletGeneratorKernels.h:225
gpuPixelDoublets::nPairsForTriplets
constexpr int nPairsForTriplets
Definition: gpuPixelDoublets.h:11
gpuPixelDoublets::nPairs
constexpr int nPairs
Definition: gpuPixelDoublets.h:12
cudaCheck
#define cudaCheck(ARG,...)
Definition: cudaCheck.h:69
CAHitNtupletGeneratorKernels::hitToTupleView_
HitToTuple::View hitToTupleView_
Definition: CAHitNtupletGeneratorKernels.h:213
relativeConstraints.value
value
Definition: relativeConstraints.py:53
CAHitNtupletGeneratorKernels::device_hitToTupleStorage_
unique_ptr< HitToTuple::Counter[]> device_hitToTupleStorage_
Definition: CAHitNtupletGeneratorKernels.h:212
TrackSoAHeterogeneousT::hitIndices
HitContainer hitIndices
Definition: TrackSoAHeterogeneousT.h:62
CAHitNtupletGeneratorKernels::cellStorage_
unique_ptr< unsigned char[]> cellStorage_
Definition: CAHitNtupletGeneratorKernels.h:201
cAHitNtupletGenerator::Params::doClusterCut_
const bool doClusterCut_
Definition: CAHitNtupletGeneratorKernels.h:118
CAHitNtupletGeneratorKernels::counters_
Counters * counters_
Definition: CAHitNtupletGeneratorKernels.h:198
cAHitNtupletGenerator::Params::doStats_
const bool doStats_
Definition: CAHitNtupletGeneratorKernels.h:117
cAHitNtupletGenerator::Params::ptmin_
const float ptmin_
Definition: CAHitNtupletGeneratorKernels.h:124
cAHitNtupletGenerator::Params::CAThetaCutBarrel_
const float CAThetaCutBarrel_
Definition: CAHitNtupletGeneratorKernels.h:125
Generated for CMSSW Reference Manual by doxygen 1.8.16