CMS 3D CMS Logo

PixelVertexProducerCUDA.cc
Go to the documentation of this file.
1 #include <cuda_runtime.h>
2 
20 
25 
26 #include "gpuVertexFinder.h"
27 
28 #undef PIXVERTEX_DEBUG_PRODUCE
29 
30 template <typename TrackerTraits>
35 
36 public:
37  explicit PixelVertexProducerCUDAT(const edm::ParameterSet& iConfig);
38  ~PixelVertexProducerCUDAT() override = default;
39 
40  static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
41 
42 private:
43  void produceOnGPU(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const;
44  void produceOnCPU(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const;
45  void produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const override;
46 
47  bool onGPU_;
48 
53 
55 
56  // Tracking cuts before sending tracks to vertex algo
57  const float ptMin_;
58  const float ptMax_;
59 };
60 
61 template <typename TrackerTraits>
63  : onGPU_(conf.getParameter<bool>("onGPU")),
64  gpuAlgo_(conf.getParameter<bool>("oneKernel"),
65  conf.getParameter<bool>("useDensity"),
66  conf.getParameter<bool>("useDBSCAN"),
67  conf.getParameter<bool>("useIterative"),
68  conf.getParameter<bool>("doSplitting"),
69  conf.getParameter<int>("minT"),
70  conf.getParameter<double>("eps"),
71  conf.getParameter<double>("errmax"),
72  conf.getParameter<double>("chi2max")),
73  ptMin_(conf.getParameter<double>("PtMin")), // 0.5 GeV
74  ptMax_(conf.getParameter<double>("PtMax")) // 75. GeV
75 {
76  if (onGPU_) {
77  tokenGPUTrack_ = consumes(conf.getParameter<edm::InputTag>("pixelTrackSrc"));
79  } else {
80  tokenCPUTrack_ = consumes(conf.getParameter<edm::InputTag>("pixelTrackSrc"));
82  }
83 }
84 
85 template <typename TrackerTraits>
88 
89  // Only one of these three algos can be used at once.
90  // Maybe this should become a Plugin Factory
91  desc.add<bool>("onGPU", true);
92  desc.add<bool>("oneKernel", true);
93  desc.add<bool>("useDensity", true);
94  desc.add<bool>("useDBSCAN", false);
95  desc.add<bool>("useIterative", false);
96  desc.add<bool>("doSplitting", true);
97 
98  desc.add<int>("minT", 2); // min number of neighbours to be "core"
99  desc.add<double>("eps", 0.07); // max absolute distance to cluster
100  desc.add<double>("errmax", 0.01); // max error to be "seed"
101  desc.add<double>("chi2max", 9.); // max normalized distance to cluster
102 
103  desc.add<double>("PtMin", 0.5);
104  desc.add<double>("PtMax", 75.);
105  desc.add<edm::InputTag>("pixelTrackSrc", edm::InputTag("pixelTracksCUDA"));
106 
107  descriptions.addWithDefaultLabel(desc);
108 }
109 
110 template <typename TrackerTraits>
113  const edm::EventSetup& iSetup) const {
115  auto hTracks = iEvent.getHandle(tokenGPUTrack_);
116 
117  cms::cuda::ScopedContextProduce ctx{*hTracks};
118  auto& tracks = ctx.get(*hTracks);
119 
120  ctx.emplace(iEvent, tokenGPUVertex_, gpuAlgo_.makeAsync(ctx.stream(), tracks.view(), ptMin_, ptMax_));
121 }
122 
123 template <typename TrackerTraits>
126  const edm::EventSetup& iSetup) const {
127  auto& tracks = iEvent.get(tokenCPUTrack_);
128 
129 #ifdef PIXVERTEX_DEBUG_PRODUCE
130  auto const& tsoa = *tracks;
131  auto maxTracks = tsoa.stride();
132  std::cout << "size of SoA " << sizeof(tsoa) << " stride " << maxTracks << std::endl;
133 
134  int32_t nt = 0;
135  for (int32_t it = 0; it < maxTracks; ++it) {
137  assert(nHits == int(tracks.view().hitIndices().size(it)));
138  if (nHits == 0)
139  break; // this is a guard: maybe we need to move to nTracks...
140  nt++;
141  }
142  std::cout << "found " << nt << " tracks in cpu SoA for Vertexing at " << tracks << std::endl;
143 #endif // PIXVERTEX_DEBUG_PRODUCE
144 
145  iEvent.emplace(tokenCPUVertex_, gpuAlgo_.make(tracks.view(), ptMin_, ptMax_));
146 }
147 
148 template <typename TrackerTraits>
151  const edm::EventSetup& iSetup) const {
152  if (onGPU_) {
153  produceOnGPU(streamID, iEvent, iSetup);
154  } else {
155  produceOnCPU(streamID, iEvent, iSetup);
156  }
157 }
158 
161 
164 
void addWithDefaultLabel(ParameterSetDescription const &psetDescription)
T getParameter(std::string const &) const
Definition: ParameterSet.h:303
edm::EDGetTokenT< TracksSoAHost > tokenCPUTrack_
void produceOnGPU(edm::StreamID streamID, edm::Event &iEvent, const edm::EventSetup &iSetup) const
~PixelVertexProducerCUDAT() override=default
assert(be >=bs)
edm::EDPutTokenT< cms::cuda::Product< ZVertexSoADevice > > tokenGPUVertex_
EDGetTokenT< ProductType > consumes(edm::InputTag const &tag)
int iEvent
Definition: GenABIO.cc:224
edm::EDPutTokenT< ZVertexSoAHost > tokenCPUVertex_
#define DEFINE_FWK_MODULE(type)
Definition: MakerMacros.h:16
void produceOnCPU(edm::StreamID streamID, edm::Event &iEvent, const edm::EventSetup &iSetup) const
edm::EDGetTokenT< cms::cuda::Product< TracksSoADevice > > tokenGPUTrack_
int nt
Definition: AMPTWrapper.h:42
static constexpr __host__ __device__ int nHits(const TrackSoAConstView &tracks, int i)
maxTracks
Definition: DMR_cfg.py:158
static void fillDescriptions(edm::ConfigurationDescriptions &descriptions)
PixelVertexProducerCUDAT(const edm::ParameterSet &iConfig)
void produce(edm::StreamID streamID, edm::Event &iEvent, const edm::EventSetup &iSetup) const override
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t nHits