PixelVertexFinding/plugins/gpuVertexFinder.h

0001 #ifndef RecoPixelVertexing_PixelVertexFinding_plugins_gpuVertexFinder_h
0002 #define RecoPixelVertexing_PixelVertexFinding_plugins_gpuVertexFinder_h
0003
0004 #include <cstddef>
0005 #include <cstdint>
0006
0007 #include "CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h"
0008 #include "CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h"
0009
0010 namespace gpuVertexFinder {
0011
0012   using ZVertices = ZVertexSoA;
0013   using TkSoA = pixelTrack::TrackSoA;
0014
0015   // workspace used in the vertex reco algos
0016   struct WorkSpace {
0017     static constexpr uint32_t MAXTRACKS = ZVertexSoA::MAXTRACKS;
0018     static constexpr uint32_t MAXVTX = ZVertexSoA::MAXVTX;
0019
0020     uint32_t ntrks;            // number of "selected tracks"
0021     uint16_t itrk[MAXTRACKS];  // index of original track
0022     float zt[MAXTRACKS];       // input track z at bs
0023     float ezt2[MAXTRACKS];     // input error^2 on the above
0024     float ptt2[MAXTRACKS];     // input pt^2 on the above
0025     uint8_t izt[MAXTRACKS];    // interized z-position of input tracks
0026     int32_t iv[MAXTRACKS];     // vertex index for each associated track
0027
0028     uint32_t nvIntermediate;  // the number of vertices after splitting pruning etc.
0029
0030     __host__ __device__ void init() {
0031       ntrks = 0;
0032       nvIntermediate = 0;
0033     }
0034   };
0035
0036   __global__ void init(ZVertexSoA* pdata, WorkSpace* pws) {
0037     pdata->init();
0038     pws->init();
0039   }
0040
0041   class Producer {
0042   public:
0043     using ZVertices = ZVertexSoA;
0044     using WorkSpace = gpuVertexFinder::WorkSpace;
0045     using TkSoA = pixelTrack::TrackSoA;
0046
0047     Producer(bool oneKernel,
0048              bool useDensity,
0049              bool useDBSCAN,
0050              bool useIterative,
0051              int iminT,      // min number of neighbours to be "core"
0052              float ieps,     // max absolute distance to cluster
0053              float ierrmax,  // max error to be "seed"
0054              float ichi2max  // max normalized distance to cluster
0055              )
0056         : oneKernel_(oneKernel && !(useDBSCAN || useIterative)),
0057           useDensity_(useDensity),
0058           useDBSCAN_(useDBSCAN),
0059           useIterative_(useIterative),
0060           minT(iminT),
0061           eps(ieps),
0062           errmax(ierrmax),
0063           chi2max(ichi2max) {}
0064
0065     ~Producer() = default;
0066
0067     ZVertexHeterogeneous makeAsync(cudaStream_t stream, TkSoA const* tksoa, float ptMin, float ptMax) const;
0068     ZVertexHeterogeneous make(TkSoA const* tksoa, float ptMin, float ptMax) const;
0069
0070   private:
0071     const bool oneKernel_;
0072     const bool useDensity_;
0073     const bool useDBSCAN_;
0074     const bool useIterative_;
0075
0076     int minT;       // min number of neighbours to be "core"
0077     float eps;      // max absolute distance to cluster
0078     float errmax;   // max error to be "seed"
0079     float chi2max;  // max normalized distance to cluster
0080   };
0081
0082 }  // namespace gpuVertexFinder
0083
0084 #endif  // RecoPixelVertexing_PixelVertexFinding_plugins_gpuVertexFinder_h