CMS 3D CMS Logo

 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Properties Friends Macros Groups Pages
gpuVertexFinder.h
Go to the documentation of this file.
1 #ifndef RecoPixelVertexing_PixelVertexFinding_plugins_gpuVertexFinder_h
2 #define RecoPixelVertexing_PixelVertexFinding_plugins_gpuVertexFinder_h
3 
4 #include <cstddef>
5 #include <cstdint>
6 
8 
9 namespace gpuVertexFinder {
10 
13 
14  // workspace used in the vertex reco algos
15  struct WorkSpace {
16  static constexpr uint32_t MAXTRACKS = ZVertexSoA::MAXTRACKS;
17  static constexpr uint32_t MAXVTX = ZVertexSoA::MAXVTX;
18 
19  uint32_t ntrks; // number of "selected tracks"
20  uint16_t itrk[MAXTRACKS]; // index of original track
21  float zt[MAXTRACKS]; // input track z at bs
22  float ezt2[MAXTRACKS]; // input error^2 on the above
23  float ptt2[MAXTRACKS]; // input pt^2 on the above
24  uint8_t izt[MAXTRACKS]; // interized z-position of input tracks
25  int32_t iv[MAXTRACKS]; // vertex index for each associated track
26 
27  uint32_t nvIntermediate; // the number of vertices after splitting pruning etc.
28 
30  ntrks = 0;
31  nvIntermediate = 0;
32  }
33  };
34 
35  __global__ void init(ZVertexSoA* pdata, WorkSpace* pws) {
36  pdata->init();
37  pws->init();
38  }
39 
40  class Producer {
41  public:
45 
46  Producer(bool oneKernel,
47  bool useDensity,
48  bool useDBSCAN,
49  bool useIterative,
50  int iminT, // min number of neighbours to be "core"
51  float ieps, // max absolute distance to cluster
52  float ierrmax, // max error to be "seed"
53  float ichi2max // max normalized distance to cluster
54  )
55  : oneKernel_(oneKernel && !(useDBSCAN || useIterative)),
56  useDensity_(useDensity),
57  useDBSCAN_(useDBSCAN),
58  useIterative_(useIterative),
59  minT(iminT),
60  eps(ieps),
61  errmax(ierrmax),
62  chi2max(ichi2max) {}
63 
64  ~Producer() = default;
65 
66  ZVertexHeterogeneous makeAsync(cudaStream_t stream, TkSoA const* tksoa, float ptMin) const;
67  ZVertexHeterogeneous make(TkSoA const* tksoa, float ptMin) const;
68 
69  private:
70  const bool oneKernel_;
71  const bool useDensity_;
72  const bool useDBSCAN_;
73  const bool useIterative_;
74 
75  int minT; // min number of neighbours to be "core"
76  float eps; // max absolute distance to cluster
77  float errmax; // max error to be "seed"
78  float chi2max; // max normalized distance to cluster
79  };
80 
81 } // namespace gpuVertexFinder
82 
83 #endif // RecoPixelVertexing_PixelVertexFinding_plugins_gpuVertexFinder_h
#define __host__
#define __global__
Definition: cudaCompat.h:19
uint16_t itrk[MAXTRACKS]
static constexpr uint32_t MAXVTX
Definition: ZVertexSoA.h:12
ZVertexHeterogeneous make(TkSoA const *tksoa, float ptMin) const
__host__ __device__ void init()
static constexpr uint32_t MAXVTX
uint8_t izt[MAXTRACKS]
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t stream
Producer(bool oneKernel, bool useDensity, bool useDBSCAN, bool useIterative, int iminT, float ieps, float ierrmax, float ichi2max)
static constexpr uint32_t MAXTRACKS
TrackSoAHeterogeneousT< maxNumber()> TrackSoA
ZVertexHeterogeneous makeAsync(cudaStream_t stream, TkSoA const *tksoa, float ptMin) const
static constexpr uint32_t MAXTRACKS
Definition: ZVertexSoA.h:11
__host__ __device__ void init()
Definition: ZVertexSoA.h:23
#define __device__
ZVertexSoA WorkSpace float ptMin