CMS 3D CMS Logo

All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Properties Friends Macros Modules Pages
List of all members | Public Member Functions | Private Types | Private Attributes
gpuVertexFinder::Producer< TrackerTraits > Class Template Reference

#include <gpuVertexFinder.h>

Public Member Functions

ZVertexSoAHost make (const TkSoAConstView &tracks_view, float ptMin, float ptMax) const
 
ZVertexSoADevice makeAsync (cudaStream_t stream, const TkSoAConstView &tracks_view, float ptMin, float ptMax) const
 
 Producer (bool oneKernel, bool useDensity, bool useDBSCAN, bool useIterative, bool doSplitting, int iminT, float ieps, float ierrmax, float ichi2max)
 
 ~Producer ()=default
 

Private Types

using TkSoAConstView = TrackSoAConstView< TrackerTraits >
 

Private Attributes

float chi2max
 
const bool doSplitting_
 
float eps
 
float errmax
 
int minT
 
const bool oneKernel_
 
const bool useDBSCAN_
 
const bool useDensity_
 
const bool useIterative_
 

Detailed Description

template<typename TrackerTraits>
class gpuVertexFinder::Producer< TrackerTraits >

Definition at line 26 of file gpuVertexFinder.h.

Member Typedef Documentation

◆ TkSoAConstView

template<typename TrackerTraits >
using gpuVertexFinder::Producer< TrackerTraits >::TkSoAConstView = TrackSoAConstView<TrackerTraits>
private

Definition at line 27 of file gpuVertexFinder.h.

Constructor & Destructor Documentation

◆ Producer()

template<typename TrackerTraits >
gpuVertexFinder::Producer< TrackerTraits >::Producer ( bool  oneKernel,
bool  useDensity,
bool  useDBSCAN,
bool  useIterative,
bool  doSplitting,
int  iminT,
float  ieps,
float  ierrmax,
float  ichi2max 
)
inline

◆ ~Producer()

template<typename TrackerTraits >
gpuVertexFinder::Producer< TrackerTraits >::~Producer ( )
default

Member Function Documentation

◆ make()

template<typename TrackerTraits >
ZVertexSoAHost gpuVertexFinder::Producer< TrackerTraits >::make ( const TkSoAConstView tracks_view,
float  ptMin,
float  ptMax 
) const

Definition at line 113 of file gpuVertexFinder.cc.

References gpuVertexFinder::assert(), gpuVertexFinder::chi2max, gather_cfg::cout, cudaCheck, gpuVertexFinder::eps, gpuVertexFinder::errmax, gpuVertexFinder::fitVertices(), init, gpuVertexFinder::maxChi2ForFinalFit, gpuVertexFinder::maxChi2ForFirstFit, gpuVertexFinder::maxChi2ForSplit, gpuVertexFinder::minT, gpuVertexFinder::ptMax, gpuVertexFinder::ptMin, gpuVertexFinder::soa, gpuVertexFinder::sortByPt2(), gpuVertexFinder::splitVertices(), cms::cuda::stream, caHitNtupletGeneratorKernels::tracks_view, and AlignmentTracksFromVertexSelector_cfi::vertices.

115  {
116 #ifdef PIXVERTEX_DEBUG_PRODUCE
117  std::cout << "producing Vertices on CPU" << std::endl;
118 #endif // PIXVERTEX_DEBUG_PRODUCE
120 #endif
121  auto soa = vertices.view();
122 
123  assert(vertices.buffer());
124 
125 #ifdef __CUDACC__
127 #else
129 #endif
130 
131 #ifdef __CUDACC__
132  init<<<1, 1, 0, stream>>>(soa, ws_d.view());
133  auto blockSize = 128;
134  auto numberOfBlocks = (tracks_view.metadata().size() + blockSize - 1) / blockSize;
135  loadTracks<TrackerTraits><<<numberOfBlocks, blockSize, 0, stream>>>(tracks_view, soa, ws_d.view(), ptMin, ptMax);
136  cudaCheck(cudaGetLastError());
137 #else
138  init(soa, ws_d.view());
139  loadTracks<TrackerTraits>(tracks_view, soa, ws_d.view(), ptMin, ptMax);
140 #endif
141 
142 #ifdef __CUDACC__
143  // Running too many thread lead to problems when printf is enabled.
144  constexpr int maxThreadsForPrint = 1024 - 128;
145  constexpr int numBlocks = 1024;
146  constexpr int threadsPerBlock = 128;
147 
148  if (oneKernel_) {
149  // implemented only for density clustesrs
150 #ifndef THREE_KERNELS
151  vertexFinderOneKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
152 #else
153  vertexFinderKernel1<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
154  cudaCheck(cudaGetLastError());
155  // one block per vertex...
156  splitVerticesKernel<<<numBlocks, threadsPerBlock, 0, stream>>>(soa, ws_d.view(), maxChi2ForSplit);
157  cudaCheck(cudaGetLastError());
158  vertexFinderKernel2<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view());
159 #endif
160  } else { // five kernels
161  if (useDensity_) {
162  clusterTracksByDensityKernel<<<1, maxThreadsForPrint, 0, stream>>>(
163  soa, ws_d.view(), minT, eps, errmax, chi2max);
164  } else if (useDBSCAN_) {
165  clusterTracksDBSCAN<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
166  } else if (useIterative_) {
167  clusterTracksIterative<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
168  }
169  cudaCheck(cudaGetLastError());
170  fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFirstFit);
171  cudaCheck(cudaGetLastError());
172  if (doSplitting_) {
173  // one block per vertex...
174  splitVerticesKernel<<<numBlocks, threadsPerBlock, 0, stream>>>(soa, ws_d.view(), maxChi2ForSplit);
175  cudaCheck(cudaGetLastError());
176  fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFinalFit);
177  cudaCheck(cudaGetLastError());
178  }
179  sortByPt2Kernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view());
180  }
181  cudaCheck(cudaGetLastError());
182 #else // __CUDACC__
183  if (useDensity_) {
184  clusterTracksByDensity(soa, ws_d.view(), minT, eps, errmax, chi2max);
185  } else if (useDBSCAN_) {
186  clusterTracksDBSCAN(soa, ws_d.view(), minT, eps, errmax, chi2max);
187  } else if (useIterative_) {
188  clusterTracksIterative(soa, ws_d.view(), minT, eps, errmax, chi2max);
189  }
190 #ifdef PIXVERTEX_DEBUG_PRODUCE
191  std::cout << "found " << ws_d.view().nvIntermediate() << " vertices " << std::endl;
192 #endif // PIXVERTEX_DEBUG_PRODUCE
193  fitVertices(soa, ws_d.view(), maxChi2ForFirstFit);
194  // one block per vertex!
195  if (doSplitting_) {
196  splitVertices(soa, ws_d.view(), maxChi2ForSplit);
197  fitVertices(soa, ws_d.view(), maxChi2ForFinalFit);
198  }
199  sortByPt2(soa, ws_d.view());
200 #endif
201 
202  return vertices;
203  }
uint32_t const *__restrict__ TkSoAView< TrackerTraits > tracks_view
fitVertices(pdata, pws, maxChi2ForFirstFit)
PixelVertexWorkSpaceSoAHost< zVertex::utilities::MAXTRACKS > PixelVertexWorkSpaceSoAHost
int init
Definition: HydjetWrapper.h:66
VtxSoAView WsSoAView float ptMin
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t stream
constexpr float maxChi2ForFirstFit
sortByPt2(pdata, pws)
VtxSoAView WsSoAView float float ptMax
splitVertices(pdata, pws, maxChi2ForSplit)
constexpr float maxChi2ForFinalFit
#define cudaCheck(ARG,...)
Definition: cudaCheck.h:69
constexpr float maxChi2ForSplit
PixelVertexWorkSpaceSoADevice< zVertex::utilities::MAXTRACKS > PixelVertexWorkSpaceSoADevice

◆ makeAsync()

template<typename TrackerTraits >
ZVertexSoADevice gpuVertexFinder::Producer< TrackerTraits >::makeAsync ( cudaStream_t  stream,
const TkSoAConstView tracks_view,
float  ptMin,
float  ptMax 
) const

Member Data Documentation

◆ chi2max

template<typename TrackerTraits >
float gpuVertexFinder::Producer< TrackerTraits >::chi2max
private

Definition at line 65 of file gpuVertexFinder.h.

◆ doSplitting_

template<typename TrackerTraits >
const bool gpuVertexFinder::Producer< TrackerTraits >::doSplitting_
private

Definition at line 60 of file gpuVertexFinder.h.

◆ eps

template<typename TrackerTraits >
float gpuVertexFinder::Producer< TrackerTraits >::eps
private

Definition at line 63 of file gpuVertexFinder.h.

◆ errmax

template<typename TrackerTraits >
float gpuVertexFinder::Producer< TrackerTraits >::errmax
private

Definition at line 64 of file gpuVertexFinder.h.

◆ minT

template<typename TrackerTraits >
int gpuVertexFinder::Producer< TrackerTraits >::minT
private

Definition at line 62 of file gpuVertexFinder.h.

◆ oneKernel_

template<typename TrackerTraits >
const bool gpuVertexFinder::Producer< TrackerTraits >::oneKernel_
private

Definition at line 56 of file gpuVertexFinder.h.

◆ useDBSCAN_

template<typename TrackerTraits >
const bool gpuVertexFinder::Producer< TrackerTraits >::useDBSCAN_
private

Definition at line 58 of file gpuVertexFinder.h.

◆ useDensity_

template<typename TrackerTraits >
const bool gpuVertexFinder::Producer< TrackerTraits >::useDensity_
private

Definition at line 57 of file gpuVertexFinder.h.

◆ useIterative_

template<typename TrackerTraits >
const bool gpuVertexFinder::Producer< TrackerTraits >::useIterative_
private

Definition at line 59 of file gpuVertexFinder.h.