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\n",
87  "tid",
88  "qual",
89  "nh",
90  "charge",
91  "pt",
92  "eta",
93  "phi",
94  "tip",
95  "zip",
96  "chi2",
97  "h1",
98  "h2",
99  "h3",
100  "h4",
101  "h5");
102 #endif
103 
104  if (m_params.onGPU_) {
105  // allocate pinned host memory only if CUDA is available
106  edm::Service<CUDAService> cs;
107  if (cs and cs->enabled()) {
108  cudaCheck(cudaMalloc(&m_counters, sizeof(Counters)));
109  cudaCheck(cudaMemset(m_counters, 0, sizeof(Counters)));
110  }
111  } else {
112  m_counters = new Counters();
113  memset(m_counters, 0, sizeof(Counters));
114  }
115 }
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 117 of file CAHitNtupletGeneratorOnGPU.cc.

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

117  {
118  if (m_params.onGPU_) {
119  // print the gpu statistics and free pinned host memory only if CUDA is available
120  edm::Service<CUDAService> cs;
121  if (cs and cs->enabled()) {
122  if (m_params.doStats_) {
123  // crash on multi-gpu processes
124  CAHitNtupletGeneratorKernelsGPU::printCounters(m_counters);
125  }
126  cudaFree(m_counters);
127  }
128  } else {
129  if (m_params.doStats_) {
130  CAHitNtupletGeneratorKernelsCPU::printCounters(m_counters);
131  }
132  delete m_counters;
133  }
134 }
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 136 of file CAHitNtupletGeneratorOnGPU.cc.

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

Referenced by CAHitNtupletCUDA::fillDescriptions().

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

216  {
217  PixelTrackHeterogeneous tracks(std::make_unique<pixelTrack::TrackSoA>());
218 
219  auto* soa = tracks.get();
220  assert(soa);
221 
222  CAHitNtupletGeneratorKernelsCPU kernels(m_params);
223  kernels.setCounters(m_counters);
224  kernels.allocateOnGPU(hits_d.nHits(), nullptr);
225 
226  kernels.buildDoublets(hits_d, nullptr);
227  kernels.launchKernels(hits_d, soa, nullptr);
228 
229  if (0 == hits_d.nHits())
230  return tracks;
231 
232  // now fit
233  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
234  fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
235 
236  if (m_params.useRiemannFit_) {
237  fitter.launchRiemannKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
238  } else {
239  fitter.launchBrokenLineKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
240  }
241 
242  kernels.classifyTuples(hits_d, soa, nullptr);
243 
244 #ifdef GPU_DEBUG
245  std::cout << "finished building pixel tracks on CPU" << std::endl;
246 #endif
247 
248  return tracks;
249 }
caConstants::maxNumberOfQuadruplets
constexpr uint32_t maxNumberOfQuadruplets
Definition: CAConstants.h:42
tracks
auto const & tracks
cannot be loose
Definition: CAHitNtupletGeneratorKernelsImpl.h:146
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 183 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().

185  {
186  PixelTrackHeterogeneous tracks(cms::cuda::make_device_unique<pixelTrack::TrackSoA>(stream));
187 
188  auto* soa = tracks.get();
189  assert(soa);
190 
191  CAHitNtupletGeneratorKernelsGPU kernels(m_params);
192  kernels.setCounters(m_counters);
193  kernels.allocateOnGPU(hits_d.nHits(), stream);
194 
195  kernels.buildDoublets(hits_d, stream);
196  kernels.launchKernels(hits_d, soa, stream);
197 
198  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
199  fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
200  if (m_params.useRiemannFit_) {
201  fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
202  } else {
203  fitter.launchBrokenLineKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
204  }
205  kernels.classifyTuples(hits_d, soa, stream);
206 
207 #ifdef GPU_DEBUG
208  cudaDeviceSynchronize();
209  cudaCheck(cudaGetLastError());
210  std::cout << "finished building pixel tracks on GPU" << std::endl;
211 #endif
212 
213  return tracks;
214 }
caConstants::maxNumberOfQuadruplets
constexpr uint32_t maxNumberOfQuadruplets
Definition: CAConstants.h:42
tracks
auto const & tracks
cannot be loose
Definition: CAHitNtupletGeneratorKernelsImpl.h:146
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