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 | 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 = TrackingRecHit2DGPU
 
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

using CAHitNtupletGeneratorOnGPU::Counters = cAHitNtupletGenerator::Counters

Definition at line 37 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::hindex_type = TrackingRecHit2DSOAView::hindex_type

Definition at line 28 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::HitContainer = pixelTrack::HitContainer

Definition at line 32 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::HitsOnCPU = TrackingRecHit2DGPU

Definition at line 27 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 26 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::OutputSoA = pixelTrack::TrackSoA

Definition at line 31 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::Params = cAHitNtupletGenerator::Params

Definition at line 36 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::Quality = pixelTrack::Quality

Definition at line 30 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::QualityCuts = cAHitNtupletGenerator::QualityCuts

Definition at line 35 of file CAHitNtupletGeneratorOnGPU.h.

using CAHitNtupletGeneratorOnGPU::Tuple = HitContainer

Definition at line 33 of file CAHitNtupletGeneratorOnGPU.h.

Constructor & Destructor Documentation

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

Definition at line 40 of file CAHitNtupletGeneratorOnGPU.h.

41  : CAHitNtupletGeneratorOnGPU(cfg, iC) {}
CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU
CAHitNtupletGeneratorOnGPU(const edm::ParameterSet &cfg, edm::ConsumesCollector &&iC)
Definition: CAHitNtupletGeneratorOnGPU.h:40
CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU ( const edm::ParameterSet &  cfg,
edm::ConsumesCollector &  iC 
)

Definition at line 60 of file CAHitNtupletGeneratorOnGPU.cc.

References fwrapper::cs, cudaCheck, CUDAService::enabled(), m_counters, m_params, cAHitNtupletGenerator::Params::onGPU_, and gpuVertexFinder::printf().

61  : m_params(cfg.getParameter<bool>("onGPU"),
62  cfg.getParameter<unsigned int>("minHitsPerNtuplet"),
63  cfg.getParameter<unsigned int>("maxNumberOfDoublets"),
64  cfg.getParameter<unsigned int>("minHitsForSharingCut"),
65  cfg.getParameter<bool>("useRiemannFit"),
66  cfg.getParameter<bool>("fitNas4"),
67  cfg.getParameter<bool>("includeJumpingForwardDoublets"),
68  cfg.getParameter<bool>("earlyFishbone"),
69  cfg.getParameter<bool>("lateFishbone"),
70  cfg.getParameter<bool>("idealConditions"),
71  cfg.getParameter<bool>("fillStatistics"),
72  cfg.getParameter<bool>("doClusterCut"),
73  cfg.getParameter<bool>("doZ0Cut"),
74  cfg.getParameter<bool>("doPtCut"),
75  cfg.getParameter<bool>("doSharedHitCut"),
76  cfg.getParameter<bool>("dupPassThrough"),
77  cfg.getParameter<bool>("useSimpleTripletCleaner"),
78  cfg.getParameter<double>("ptmin"),
79  cfg.getParameter<double>("CAThetaCutBarrel"),
80  cfg.getParameter<double>("CAThetaCutForward"),
81  cfg.getParameter<double>("hardCurvCut"),
82  cfg.getParameter<double>("dcaCutInnerTriplet"),
83  cfg.getParameter<double>("dcaCutOuterTriplet"),
84  makeQualityCuts(cfg.getParameterSet("trackQualityCuts"))) {
85 #ifdef DUMP_GPU_TK_TUPLES
86  printf("TK: %s %s %s %s %s %s %s %s %s %s %s %s %s %s %s %s %s\n",
87  "tid",
88  "qual",
89  "nh",
90  "nl",
91  "charge",
92  "pt",
93  "eta",
94  "phi",
95  "tip",
96  "zip",
97  "chi2",
98  "h1",
99  "h2",
100  "h3",
101  "h4",
102  "h5",
103  "hn");
104 #endif
105 
106  if (m_params.onGPU_) {
107  // allocate pinned host memory only if CUDA is available
108  edm::Service<CUDAService> cs;
109  if (cs and cs->enabled()) {
110  cudaCheck(cudaMalloc(&m_counters, sizeof(Counters)));
111  cudaCheck(cudaMemset(m_counters, 0, sizeof(Counters)));
112  }
113  } else {
114  m_counters = new Counters();
115  memset(m_counters, 0, sizeof(Counters));
116  }
117 }
fwrapper::cs
unique_ptr< ClusterSequence > cs
Definition: fastjetfortran_madfks.cc:47
CUDAService::enabled
bool enabled() const
Definition: CUDAService.h:22
gpuVertexFinder::printf
printf("params %d %f %f %f\n", minT, eps, errmax, chi2max)
CAHitNtupletGeneratorOnGPU::Counters
cAHitNtupletGenerator::Counters Counters
Definition: CAHitNtupletGeneratorOnGPU.h:37
edm::ParameterSet::getParameterSet
ParameterSet const & getParameterSet(std::string const &) const
Definition: ParameterSet.cc:2128
edm::ParameterSet::getParameter
T getParameter(std::string const &) const
Definition: ParameterSet.h:303
cudaCheck
#define cudaCheck(ARG,...)
Definition: cudaCheck.h:69
CAHitNtupletGeneratorOnGPU::~CAHitNtupletGeneratorOnGPU ( )

Definition at line 119 of file CAHitNtupletGeneratorOnGPU.cc.

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

119  {
120  if (m_params.onGPU_) {
121  // print the gpu statistics and free pinned host memory only if CUDA is available
122  edm::Service<CUDAService> cs;
123  if (cs and cs->enabled()) {
124  if (m_params.doStats_) {
125  // crash on multi-gpu processes
126  CAHitNtupletGeneratorKernelsGPU::printCounters(m_counters);
127  }
128  cudaFree(m_counters);
129  }
130  } else {
131  if (m_params.doStats_) {
132  CAHitNtupletGeneratorKernelsCPU::printCounters(m_counters);
133  }
134  delete m_counters;
135  }
136 }
fwrapper::cs
unique_ptr< ClusterSequence > cs
Definition: fastjetfortran_madfks.cc:47
CAHitNtupletGeneratorKernels::printCounters
static void printCounters(Counters const *counters)
Definition: CAHitNtupletGeneratorKernels.cc:12
CUDAService::enabled
bool enabled() const
Definition: CUDAService.h:22

Member Function Documentation

void CAHitNtupletGeneratorOnGPU::buildDoublets ( HitsOnCPU const &  hh,
cudaStream_t  stream 
) const
private
void CAHitNtupletGeneratorOnGPU::fillDescriptions ( edm::ParameterSetDescription &  desc)
static

Definition at line 138 of file CAHitNtupletGeneratorOnGPU.cc.

References edm::ParameterSetDescription::add(), caConstants::maxNumberOfDoublets, edm::ParameterDescriptionNode::setComment(), and HLT_FULL_cff::trackQualityCuts.

Referenced by CAHitNtupletCUDA::fillDescriptions().

138  {
139  // 87 cm/GeV = 1/(3.8T * 0.3)
140  // take less than radius given by the hardPtCut and reject everything below
141  // auto hardCurvCut = 1.f/(0.35 * 87.f);
142  desc.add<double>("ptmin", 0.9f)->setComment("Cut on minimum pt");
143  desc.add<double>("CAThetaCutBarrel", 0.002f)->setComment("Cut on RZ alignement for Barrel");
144  desc.add<double>("CAThetaCutForward", 0.003f)->setComment("Cut on RZ alignment for Forward");
145  desc.add<double>("hardCurvCut", 1.f / (0.35 * 87.f))->setComment("Cut on minimum curvature");
146  desc.add<double>("dcaCutInnerTriplet", 0.15f)->setComment("Cut on origin radius when the inner hit is on BPix1");
147  desc.add<double>("dcaCutOuterTriplet", 0.25f)->setComment("Cut on origin radius when the outer hit is on BPix1");
148  desc.add<bool>("earlyFishbone", true);
149  desc.add<bool>("lateFishbone", false);
150  desc.add<bool>("idealConditions", true);
151  desc.add<bool>("fillStatistics", false);
152  desc.add<unsigned int>("minHitsPerNtuplet", 4);
153  desc.add<unsigned int>("maxNumberOfDoublets", caConstants::maxNumberOfDoublets);
154  desc.add<unsigned int>("minHitsForSharingCut", 10)
155  ->setComment("Maximum number of hits in a tuple to clean also if the shared hit is on bpx1");
156  desc.add<bool>("includeJumpingForwardDoublets", false);
157  desc.add<bool>("fitNas4", false)->setComment("fit only 4 hits out of N");
158  desc.add<bool>("doClusterCut", true);
159  desc.add<bool>("doZ0Cut", true);
160  desc.add<bool>("doPtCut", true);
161  desc.add<bool>("useRiemannFit", false)->setComment("true for Riemann, false for BrokenLine");
162  desc.add<bool>("doSharedHitCut", true)->setComment("Sharing hit nTuples cleaning");
163  desc.add<bool>("dupPassThrough", false)->setComment("Do not reject duplicate");
164  desc.add<bool>("useSimpleTripletCleaner", true)->setComment("use alternate implementation");
165 
166  edm::ParameterSetDescription trackQualityCuts;
167  trackQualityCuts.add<double>("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut");
168  trackQualityCuts.add<std::vector<double>>("chi2Coeff", {0.9, 1.8})->setComment("chi2 at 1GeV and at ptMax above");
169  trackQualityCuts.add<double>("chi2Scale", 8.)
170  ->setComment(
171  "Factor to multiply the pT-dependent chi2 cut (currently: 8 for the broken line fit, ?? for the Riemann "
172  "fit)");
173  trackQualityCuts.add<double>("tripletMinPt", 0.5)->setComment("Min pT for triplets, in GeV");
174  trackQualityCuts.add<double>("tripletMaxTip", 0.3)->setComment("Max |Tip| for triplets, in cm");
175  trackQualityCuts.add<double>("tripletMaxZip", 12.)->setComment("Max |Zip| for triplets, in cm");
176  trackQualityCuts.add<double>("quadrupletMinPt", 0.3)->setComment("Min pT for quadruplets, in GeV");
177  trackQualityCuts.add<double>("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm");
178  trackQualityCuts.add<double>("quadrupletMaxZip", 12.)->setComment("Max |Zip| for quadruplets, in cm");
179  desc.add<edm::ParameterSetDescription>("trackQualityCuts", trackQualityCuts)
180  ->setComment(
181  "Quality cuts based on the results of the track fit:\n - apply a pT-dependent chi2 cut;\n - apply \"region "
182  "cuts\" based on the fit results (pT, Tip, Zip).");
183 }
edm::ParameterDescriptionNode::setComment
void setComment(std::string const &value)
Definition: ParameterDescriptionNode.cc:106
edm::ParameterSetDescription::add
ParameterDescriptionBase * add(U const &iLabel, T const &value)
Definition: ParameterSetDescription.h:95
caConstants::maxNumberOfDoublets
constexpr uint32_t maxNumberOfDoublets
Definition: CAConstants.h:37
HLT_FULL_cff.trackQualityCuts
tuple trackQualityCuts
Definition: HLT_FULL_cff.py:9452
static const char* CAHitNtupletGeneratorOnGPU::fillDescriptionsLabel ( )
inlinestatic

Definition at line 47 of file CAHitNtupletGeneratorOnGPU.h.

47 { return "caHitNtupletOnGPU"; }
void CAHitNtupletGeneratorOnGPU::hitNtuplets ( HitsOnCPU const &  hh,
const edm::EventSetup &  es,
bool  useRiemannFit,
cudaStream_t  cudaStream 
)
private
void CAHitNtupletGeneratorOnGPU::launchKernels ( HitsOnCPU const &  hh,
bool  useRiemannFit,
cudaStream_t  cudaStream 
) const
private
PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuples ( TrackingRecHit2DCPU const &  hits_d,
float  bfield 
) const

Definition at line 218 of file CAHitNtupletGeneratorOnGPU.cc.

References HelixFitOnGPU::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), cms::cuda::assert(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), gather_cfg::cout, cAHitNtupletGenerator::Params::fitNas4_, HeterogeneousSoA< T >::get(), 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().

218  {
219  PixelTrackHeterogeneous tracks(std::make_unique<pixelTrack::TrackSoA>());
220 
221  auto* soa = tracks.get();
222  assert(soa);
223 
224  CAHitNtupletGeneratorKernelsCPU kernels(m_params);
225  kernels.setCounters(m_counters);
226  kernels.allocateOnGPU(hits_d.nHits(), nullptr);
227 
228  kernels.buildDoublets(hits_d, nullptr);
229  kernels.launchKernels(hits_d, soa, nullptr);
230 
231  if (0 == hits_d.nHits())
232  return tracks;
233 
234  // now fit
235  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
236  fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
237 
238  if (m_params.useRiemannFit_) {
239  fitter.launchRiemannKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
240  } else {
241  fitter.launchBrokenLineKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
242  }
243 
244  kernels.classifyTuples(hits_d, soa, nullptr);
245 
246 #ifdef GPU_DEBUG
247  std::cout << "finished building pixel tracks on CPU" << std::endl;
248 #endif
249 
250  return tracks;
251 }
caConstants::maxNumberOfQuadruplets
constexpr uint32_t maxNumberOfQuadruplets
Definition: CAConstants.h:42
tracks
auto const & tracks
cannot be loose
Definition: CAHitNtupletGeneratorKernelsImpl.h:148
cms::cuda::assert
assert(be >=bs)
gpuVertexFinder::soa
ZVertexSoA * soa
Definition: gpuVertexFinder.cc:21
gather_cfg.cout
tuple cout
Definition: gather_cfg.py:144
PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuplesAsync ( TrackingRecHit2DGPU const &  hits_d,
float  bfield,
cudaStream_t  stream 
) const

Definition at line 185 of file CAHitNtupletGeneratorOnGPU.cc.

References HelixFitOnGPU::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), cms::cuda::assert(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), gather_cfg::cout, cudaCheck, cAHitNtupletGenerator::Params::fitNas4_, HeterogeneousSoA< T >::get(), 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().

187  {
188  PixelTrackHeterogeneous tracks(cms::cuda::make_device_unique<pixelTrack::TrackSoA>(stream));
189 
190  auto* soa = tracks.get();
191  assert(soa);
192 
193  CAHitNtupletGeneratorKernelsGPU kernels(m_params);
194  kernels.setCounters(m_counters);
195  kernels.allocateOnGPU(hits_d.nHits(), stream);
196 
197  kernels.buildDoublets(hits_d, stream);
198  kernels.launchKernels(hits_d, soa, stream);
199 
200  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
201  fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
202  if (m_params.useRiemannFit_) {
203  fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
204  } else {
205  fitter.launchBrokenLineKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
206  }
207  kernels.classifyTuples(hits_d, soa, stream);
208 
209 #ifdef GPU_DEBUG
210  cudaDeviceSynchronize();
211  cudaCheck(cudaGetLastError());
212  std::cout << "finished building pixel tracks on GPU" << std::endl;
213 #endif
214 
215  return tracks;
216 }
caConstants::maxNumberOfQuadruplets
constexpr uint32_t maxNumberOfQuadruplets
Definition: CAConstants.h:42
tracks
auto const & tracks
cannot be loose
Definition: CAHitNtupletGeneratorKernelsImpl.h:148
cms::cuda::stream
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t stream
Definition: HistoContainer.h:51
cms::cuda::assert
assert(be >=bs)
gpuVertexFinder::soa
ZVertexSoA * soa
Definition: gpuVertexFinder.cc:21
gather_cfg.cout
tuple cout
Definition: gather_cfg.py:144
cudaCheck
#define cudaCheck(ARG,...)
Definition: cudaCheck.h:69

Member Data Documentation

Counters* CAHitNtupletGeneratorOnGPU::m_counters = nullptr
private

Definition at line 62 of file CAHitNtupletGeneratorOnGPU.h.

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

Params CAHitNtupletGeneratorOnGPU::m_params
private

Definition at line 60 of file CAHitNtupletGeneratorOnGPU.h.

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

Generated for CMSSW Reference Manual by doxygen 1.8.5