CMS 3D CMS Logo
List of all members | Public Types | Public Member Functions | Static Public Member Functions | Private Member Functions | Private Attributes
CAHitNtupletGeneratorOnGPU Class Reference

#include <CAHitNtupletGeneratorOnGPU.h>

Public Types

using Counters = cAHitNtupletGenerator::Counters
 
using hindex_type = TrackingRecHit2DSOAView::hindex_type
 
using HitContainer = pixelTrack::HitContainer
 
using HitsOnCPU = TrackingRecHit2DCUDA
 
using HitsOnGPU = TrackingRecHit2DSOAView
 
using OutputSoA = pixelTrack::TrackSoA
 
using Params = cAHitNtupletGenerator::Params
 
using Quality = pixelTrack::Quality
 
using QualityCuts = cAHitNtupletGenerator::QualityCuts
 
using Tuple = HitContainer
 

Public Member Functions

 CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &&iC)
 
 CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &iC)
 
PixelTrackHeterogeneous makeTuples (TrackingRecHit2DCPU const &hits_d, float bfield) const
 
PixelTrackHeterogeneous makeTuplesAsync (TrackingRecHit2DGPU const &hits_d, float bfield, cudaStream_t stream) const
 
 ~CAHitNtupletGeneratorOnGPU ()
 

Static Public Member Functions

static void fillDescriptions (edm::ParameterSetDescription &desc)
 
static const char * fillDescriptionsLabel ()
 

Private Member Functions

void buildDoublets (HitsOnCPU const &hh, cudaStream_t stream) const
 
void hitNtuplets (HitsOnCPU const &hh, const edm::EventSetup &es, bool useRiemannFit, cudaStream_t cudaStream)
 
void launchKernels (HitsOnCPU const &hh, bool useRiemannFit, cudaStream_t cudaStream) const
 

Private Attributes

Counters * m_counters = nullptr
 
Params m_params
 

Detailed Description

Definition at line 24 of file CAHitNtupletGeneratorOnGPU.h.

Member Typedef Documentation

◆ Counters

using CAHitNtupletGeneratorOnGPU::Counters = cAHitNtupletGenerator::Counters

Definition at line 37 of file CAHitNtupletGeneratorOnGPU.h.

◆ hindex_type

using CAHitNtupletGeneratorOnGPU::hindex_type = TrackingRecHit2DSOAView::hindex_type

Definition at line 28 of file CAHitNtupletGeneratorOnGPU.h.

◆ HitContainer

using CAHitNtupletGeneratorOnGPU::HitContainer = pixelTrack::HitContainer

Definition at line 32 of file CAHitNtupletGeneratorOnGPU.h.

◆ HitsOnCPU

using CAHitNtupletGeneratorOnGPU::HitsOnCPU = TrackingRecHit2DCUDA

Definition at line 27 of file CAHitNtupletGeneratorOnGPU.h.

◆ HitsOnGPU

using CAHitNtupletGeneratorOnGPU::HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 26 of file CAHitNtupletGeneratorOnGPU.h.

◆ OutputSoA

using CAHitNtupletGeneratorOnGPU::OutputSoA = pixelTrack::TrackSoA

Definition at line 31 of file CAHitNtupletGeneratorOnGPU.h.

◆ Params

using CAHitNtupletGeneratorOnGPU::Params = cAHitNtupletGenerator::Params

Definition at line 36 of file CAHitNtupletGeneratorOnGPU.h.

◆ Quality

using CAHitNtupletGeneratorOnGPU::Quality = pixelTrack::Quality

Definition at line 30 of file CAHitNtupletGeneratorOnGPU.h.

◆ QualityCuts

using CAHitNtupletGeneratorOnGPU::QualityCuts = cAHitNtupletGenerator::QualityCuts

Definition at line 35 of file CAHitNtupletGeneratorOnGPU.h.

◆ Tuple

using CAHitNtupletGeneratorOnGPU::Tuple = HitContainer

Definition at line 33 of file CAHitNtupletGeneratorOnGPU.h.

Constructor & Destructor Documentation

◆ CAHitNtupletGeneratorOnGPU() [1/2]

CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU ( const edm::ParameterSet &  cfg,
edm::ConsumesCollector &&  iC 
)
inline

Definition at line 40 of file CAHitNtupletGeneratorOnGPU.h.

41  : CAHitNtupletGeneratorOnGPU(cfg, iC) {}

◆ CAHitNtupletGeneratorOnGPU() [2/2]

CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU ( const edm::ParameterSet &  cfg,
edm::ConsumesCollector &  iC 
)

Definition at line 56 of file CAHitNtupletGeneratorOnGPU.cc.

57  : m_params(cfg.getParameter<bool>("onGPU"),
58  cfg.getParameter<unsigned int>("minHitsPerNtuplet"),
59  cfg.getParameter<unsigned int>("maxNumberOfDoublets"),
60  cfg.getParameter<unsigned int>("minHitsForSharingCut"),
61  cfg.getParameter<bool>("useRiemannFit"),
62  cfg.getParameter<bool>("fit5as4"),
63  cfg.getParameter<bool>("includeJumpingForwardDoublets"),
64  cfg.getParameter<bool>("earlyFishbone"),
65  cfg.getParameter<bool>("lateFishbone"),
66  cfg.getParameter<bool>("idealConditions"),
67  cfg.getParameter<bool>("fillStatistics"),
68  cfg.getParameter<bool>("doClusterCut"),
69  cfg.getParameter<bool>("doZ0Cut"),
70  cfg.getParameter<bool>("doPtCut"),
71  cfg.getParameter<bool>("doSharedHitCut"),
72  cfg.getParameter<double>("ptmin"),
73  cfg.getParameter<double>("CAThetaCutBarrel"),
74  cfg.getParameter<double>("CAThetaCutForward"),
75  cfg.getParameter<double>("hardCurvCut"),
76  cfg.getParameter<double>("dcaCutInnerTriplet"),
77  cfg.getParameter<double>("dcaCutOuterTriplet"),
78  makeQualityCuts(cfg.getParameterSet("trackQualityCuts"))) {
79 #ifdef DUMP_GPU_TK_TUPLES
80  printf("TK: %s %s % %s %s %s %s %s %s %s %s %s %s %s %s %s\n",
81  "tid",
82  "qual",
83  "nh",
84  "charge",
85  "pt",
86  "eta",
87  "phi",
88  "tip",
89  "zip",
90  "chi2",
91  "h1",
92  "h2",
93  "h3",
94  "h4",
95  "h5");
96 #endif
97 
98  if (m_params.onGPU_) {
99  // allocate pinned host memory only if CUDA is available
100  edm::Service<CUDAService> cs;
101  if (cs and cs->enabled()) {
102  cudaCheck(cudaMalloc(&m_counters, sizeof(Counters)));
103  cudaCheck(cudaMemset(m_counters, 0, sizeof(Counters)));
104  }
105  } else {
106  m_counters = new Counters();
107  memset(m_counters, 0, sizeof(Counters));
108  }
109 }

References fwrapper::cs, cudaCheck, m_counters, m_params, and cAHitNtupletGenerator::Params::onGPU_.

◆ ~CAHitNtupletGeneratorOnGPU()

CAHitNtupletGeneratorOnGPU::~CAHitNtupletGeneratorOnGPU ( )

Definition at line 111 of file CAHitNtupletGeneratorOnGPU.cc.

111  {
112  if (m_params.onGPU_) {
113  // print the gpu statistics and free pinned host memory only if CUDA is available
114  edm::Service<CUDAService> cs;
115  if (cs and cs->enabled()) {
116  if (m_params.doStats_) {
117  // crash on multi-gpu processes
118  CAHitNtupletGeneratorKernelsGPU::printCounters(m_counters);
119  }
120  cudaFree(m_counters);
121  }
122  } else {
123  if (m_params.doStats_) {
124  CAHitNtupletGeneratorKernelsCPU::printCounters(m_counters);
125  }
126  delete m_counters;
127  }
128 }

References fwrapper::cs, cAHitNtupletGenerator::Params::doStats_, m_counters, m_params, cAHitNtupletGenerator::Params::onGPU_, and CAHitNtupletGeneratorKernels< TTraits >::printCounters().

Member Function Documentation

◆ buildDoublets()

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

◆ fillDescriptions()

void CAHitNtupletGeneratorOnGPU::fillDescriptions ( edm::ParameterSetDescription &  desc)
static

Definition at line 130 of file CAHitNtupletGeneratorOnGPU.cc.

130  {
131  // 87 cm/GeV = 1/(3.8T * 0.3)
132  // take less than radius given by the hardPtCut and reject everything below
133  // auto hardCurvCut = 1.f/(0.35 * 87.f);
134  desc.add<double>("ptmin", 0.9f)->setComment("Cut on minimum pt");
135  desc.add<double>("CAThetaCutBarrel", 0.002f)->setComment("Cut on RZ alignement for Barrel");
136  desc.add<double>("CAThetaCutForward", 0.003f)->setComment("Cut on RZ alignment for Forward");
137  desc.add<double>("hardCurvCut", 1.f / (0.35 * 87.f))->setComment("Cut on minimum curvature");
138  desc.add<double>("dcaCutInnerTriplet", 0.15f)->setComment("Cut on origin radius when the inner hit is on BPix1");
139  desc.add<double>("dcaCutOuterTriplet", 0.25f)->setComment("Cut on origin radius when the outer hit is on BPix1");
140  desc.add<bool>("earlyFishbone", true);
141  desc.add<bool>("lateFishbone", false);
142  desc.add<bool>("idealConditions", true);
143  desc.add<bool>("fillStatistics", false);
144  desc.add<unsigned int>("minHitsPerNtuplet", 4);
145  desc.add<unsigned int>("maxNumberOfDoublets", caConstants::maxNumberOfDoublets);
146  desc.add<unsigned int>("minHitsForSharingCut", 5)
147  ->setComment("Maximum number of hits in a tuple to clean also if the shared hit is on bpx1");
148  desc.add<bool>("includeJumpingForwardDoublets", false);
149  desc.add<bool>("fit5as4", true);
150  desc.add<bool>("doClusterCut", true);
151  desc.add<bool>("doZ0Cut", true);
152  desc.add<bool>("doPtCut", true);
153  desc.add<bool>("useRiemannFit", false)->setComment("true for Riemann, false for BrokenLine");
154  desc.add<bool>("doSharedHitCut", true)->setComment("Sharing hit nTuples cleaning");
155 
156  edm::ParameterSetDescription trackQualityCuts;
157  trackQualityCuts.add<double>("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut");
158  trackQualityCuts.add<std::vector<double>>("chi2Coeff", {0.68177776, 0.74609577, -0.08035491, 0.00315399})
159  ->setComment("Polynomial coefficients to derive the pT-dependent chi2 cut");
160  trackQualityCuts.add<double>("chi2Scale", 30.)
161  ->setComment(
162  "Factor to multiply the pT-dependent chi2 cut (currently: 30 for the broken line fit, 45 for the Riemann "
163  "fit)");
164  trackQualityCuts.add<double>("tripletMinPt", 0.5)->setComment("Min pT for triplets, in GeV");
165  trackQualityCuts.add<double>("tripletMaxTip", 0.3)->setComment("Max |Tip| for triplets, in cm");
166  trackQualityCuts.add<double>("tripletMaxZip", 12.)->setComment("Max |Zip| for triplets, in cm");
167  trackQualityCuts.add<double>("quadrupletMinPt", 0.3)->setComment("Min pT for quadruplets, in GeV");
168  trackQualityCuts.add<double>("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm");
169  trackQualityCuts.add<double>("quadrupletMaxZip", 12.)->setComment("Max |Zip| for quadruplets, in cm");
170  desc.add<edm::ParameterSetDescription>("trackQualityCuts", trackQualityCuts)
171  ->setComment(
172  "Quality cuts based on the results of the track fit:\n - apply a pT-dependent chi2 cut;\n - apply \"region "
173  "cuts\" based on the fit results (pT, Tip, Zip).");
174 }

References edm::ParameterSetDescription::add(), submitPVResolutionJobs::desc, caConstants::maxNumberOfDoublets, and edm::ParameterDescriptionNode::setComment().

Referenced by CAHitNtupletCUDA::fillDescriptions().

◆ fillDescriptionsLabel()

static const char* CAHitNtupletGeneratorOnGPU::fillDescriptionsLabel ( )
inlinestatic

Definition at line 47 of file CAHitNtupletGeneratorOnGPU.h.

47 { return "caHitNtupletOnGPU"; }

◆ hitNtuplets()

void CAHitNtupletGeneratorOnGPU::hitNtuplets ( HitsOnCPU const &  hh,
const edm::EventSetup &  es,
bool  useRiemannFit,
cudaStream_t  cudaStream 
)
private

◆ launchKernels()

void CAHitNtupletGeneratorOnGPU::launchKernels ( HitsOnCPU const &  hh,
bool  useRiemannFit,
cudaStream_t  cudaStream 
) const
private

◆ makeTuples()

PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuples ( TrackingRecHit2DCPU const &  hits_d,
float  bfield 
) const

Definition at line 204 of file CAHitNtupletGeneratorOnGPU.cc.

204  {
205  PixelTrackHeterogeneous tracks(std::make_unique<pixelTrack::TrackSoA>());
206 
207  auto* soa = tracks.get();
208  assert(soa);
209 
210  CAHitNtupletGeneratorKernelsCPU kernels(m_params);
211  kernels.setCounters(m_counters);
212  kernels.allocateOnGPU(nullptr);
213 
214  kernels.buildDoublets(hits_d, nullptr);
215  kernels.launchKernels(hits_d, soa, nullptr);
216  kernels.fillHitDetIndices(hits_d.view(), soa, nullptr); // in principle needed only if Hits not "available"
217 
218  if (0 == hits_d.nHits())
219  return tracks;
220 
221  // now fit
222  HelixFitOnGPU fitter(bfield, m_params.fit5as4_);
223  fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
224 
225  if (m_params.useRiemannFit_) {
226  fitter.launchRiemannKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
227  } else {
228  fitter.launchBrokenLineKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
229  }
230 
231  kernels.classifyTuples(hits_d, soa, nullptr);
232 
233  return tracks;
234 }

References HelixFitOnGPU::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), cms::cuda::assert(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), CAHitNtupletGeneratorKernels< TTraits >::fillHitDetIndices(), cAHitNtupletGenerator::Params::fit5as4_, HelixFitOnGPU::launchBrokenLineKernelsOnCPU(), CAHitNtupletGeneratorKernels< TTraits >::launchKernels(), HelixFitOnGPU::launchRiemannKernelsOnCPU(), m_counters, m_params, caConstants::maxNumberOfQuadruplets, TrackingRecHit2DHeterogeneous< Traits >::nHits(), CAHitNtupletGeneratorKernels< TTraits >::setCounters(), gpuVertexFinder::soa, tracks, CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity(), cAHitNtupletGenerator::Params::useRiemannFit_, and TrackingRecHit2DHeterogeneous< Traits >::view().

Referenced by CAHitNtupletCUDA::produce().

◆ makeTuplesAsync()

PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuplesAsync ( TrackingRecHit2DGPU const &  hits_d,
float  bfield,
cudaStream_t  stream 
) const

Definition at line 176 of file CAHitNtupletGeneratorOnGPU.cc.

178  {
179  PixelTrackHeterogeneous tracks(cms::cuda::make_device_unique<pixelTrack::TrackSoA>(stream));
180 
181  auto* soa = tracks.get();
182 
183  CAHitNtupletGeneratorKernelsGPU kernels(m_params);
184  kernels.setCounters(m_counters);
185 
186  kernels.allocateOnGPU(stream);
187 
188  kernels.buildDoublets(hits_d, stream);
189  kernels.launchKernels(hits_d, soa, stream);
190  kernels.fillHitDetIndices(hits_d.view(), soa, stream); // in principle needed only if Hits not "available"
191 
192  HelixFitOnGPU fitter(bfield, m_params.fit5as4_);
193  fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
194  if (m_params.useRiemannFit_) {
195  fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
196  } else {
197  fitter.launchBrokenLineKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
198  }
199  kernels.classifyTuples(hits_d, soa, stream);
200 
201  return tracks;
202 }

References HelixFitOnGPU::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), CAHitNtupletGeneratorKernels< TTraits >::fillHitDetIndices(), cAHitNtupletGenerator::Params::fit5as4_, HelixFitOnGPU::launchBrokenLineKernels(), CAHitNtupletGeneratorKernels< TTraits >::launchKernels(), HelixFitOnGPU::launchRiemannKernels(), m_counters, m_params, caConstants::maxNumberOfQuadruplets, TrackingRecHit2DHeterogeneous< Traits >::nHits(), CAHitNtupletGeneratorKernels< TTraits >::setCounters(), gpuVertexFinder::soa, cms::cuda::stream, tracks, CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity(), cAHitNtupletGenerator::Params::useRiemannFit_, and TrackingRecHit2DHeterogeneous< Traits >::view().

Referenced by CAHitNtupletCUDA::produce().

Member Data Documentation

◆ m_counters

Counters* CAHitNtupletGeneratorOnGPU::m_counters = nullptr
private

Definition at line 62 of file CAHitNtupletGeneratorOnGPU.h.

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

◆ m_params

Params CAHitNtupletGeneratorOnGPU::m_params
private

Definition at line 60 of file CAHitNtupletGeneratorOnGPU.h.

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

edm::ParameterSetDescription::add
ParameterDescriptionBase * add(U const &iLabel, T const &value)
Definition: ParameterSetDescription.h:95
fwrapper::cs
unique_ptr< ClusterSequence > cs
Definition: fastjetfortran_madfks.cc:47
caConstants::maxNumberOfDoublets
constexpr uint32_t maxNumberOfDoublets
Definition: CAConstants.h:36
cms::cuda::stream
cudaStream_t stream
Definition: HistoContainer.h:57
edm::ParameterSetDescription
Definition: ParameterSetDescription.h:52
cms::cuda::assert
assert(be >=bs)
gpuVertexFinder::soa
ZVertexSoA * soa
Definition: gpuVertexFinderImpl.h:24
HeterogeneousSoA
Definition: HeterogeneousSoA.h:13
caConstants::maxNumberOfQuadruplets
constexpr uint32_t maxNumberOfQuadruplets
Definition: CAConstants.h:41
cAHitNtupletGenerator::Params::onGPU_
const bool onGPU_
Definition: CAHitNtupletGeneratorKernels.h:100
tracks
const uint32_t *__restrict__ const HitContainer *__restrict__ TkSoA *__restrict__ tracks
Definition: CAHitNtupletGeneratorKernelsImpl.h:159
edm::Service
Definition: Service.h:30
CAHitNtupletGeneratorOnGPU::Counters
cAHitNtupletGenerator::Counters Counters
Definition: CAHitNtupletGeneratorOnGPU.h:37
CAHitNtupletGeneratorKernels::printCounters
static void printCounters(Counters const *counters)
Definition: CAHitNtupletGeneratorKernels.cc:4
looper.cfg
cfg
Definition: looper.py:297
CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU
CAHitNtupletGeneratorOnGPU(const edm::ParameterSet &cfg, edm::ConsumesCollector &&iC)
Definition: CAHitNtupletGeneratorOnGPU.h:40
CAHitNtupletGeneratorKernels
Definition: CAHitNtupletGeneratorKernels.h:147
submitPVResolutionJobs.desc
string desc
Definition: submitPVResolutionJobs.py:251
cudaCheck
#define cudaCheck(ARG,...)
Definition: cudaCheck.h:62
cAHitNtupletGenerator::Params::useRiemannFit_
const bool useRiemannFit_
Definition: CAHitNtupletGeneratorKernels.h:104
CAHitNtupletGeneratorOnGPU::m_counters
Counters * m_counters
Definition: CAHitNtupletGeneratorOnGPU.h:62
CAHitNtupletGeneratorOnGPU::m_params
Params m_params
Definition: CAHitNtupletGeneratorOnGPU.h:60
edm::ParameterDescriptionNode::setComment
void setComment(std::string const &value)
Definition: ParameterDescriptionNode.cc:106
cAHitNtupletGenerator::Params::doStats_
const bool doStats_
Definition: CAHitNtupletGeneratorKernels.h:110
HelixFitOnGPU
Definition: HelixFitOnGPU.h:34
cAHitNtupletGenerator::Params::fit5as4_
const bool fit5as4_
Definition: CAHitNtupletGeneratorKernels.h:105
Generated for CMSSW Reference Manual by doxygen 1.8.16