PixelSeeding/plugins/gpuPixelDoublets.h

0001 #ifndef RecoTracker_PixelSeeding_plugins_gpuPixelDoublets_h
0002 #define RecoTracker_PixelSeeding_plugins_gpuPixelDoublets_h
0003
0004 #include "RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h"
0005
0006 #define CONSTANT_VAR __constant__
0007
0008 namespace gpuPixelDoublets {
0009
0010   template <typename TrackerTraits>
0011   __global__ void initDoublets(OuterHitOfCell<TrackerTraits> isOuterHitOfCell,
0012                                int nHits,
0013                                CellNeighborsVector<TrackerTraits>* cellNeighbors,
0014                                CellNeighbors<TrackerTraits>* cellNeighborsContainer,
0015                                CellTracksVector<TrackerTraits>* cellTracks,
0016                                CellTracks<TrackerTraits>* cellTracksContainer) {
0017     assert(isOuterHitOfCell.container);
0018     int first = blockIdx.x * blockDim.x + threadIdx.x;
0019     for (int i = first; i < nHits - isOuterHitOfCell.offset; i += gridDim.x * blockDim.x)
0020       isOuterHitOfCell.container[i].reset();
0021
0022     if (0 == first) {
0023       cellNeighbors->construct(TrackerTraits::maxNumOfActiveDoublets, cellNeighborsContainer);
0024       cellTracks->construct(TrackerTraits::maxNumOfActiveDoublets, cellTracksContainer);
0025       auto i = cellNeighbors->extend();
0026       assert(0 == i);
0027       (*cellNeighbors)[0].reset();
0028       i = cellTracks->extend();
0029       assert(0 == i);
0030       (*cellTracks)[0].reset();
0031     }
0032   }
0033
0034   constexpr auto getDoubletsFromHistoMaxBlockSize = 64;  // for both x and y
0035   constexpr auto getDoubletsFromHistoMinBlocksPerMP = 16;
0036
0037   template <typename TrackerTraits>
0038   __global__
0039 #ifdef __CUDACC__
0040   __launch_bounds__(getDoubletsFromHistoMaxBlockSize, getDoubletsFromHistoMinBlocksPerMP)
0041 #endif
0042       void getDoubletsFromHisto(GPUCACellT<TrackerTraits>* cells,
0043                                 uint32_t* nCells,
0044                                 CellNeighborsVector<TrackerTraits>* cellNeighbors,
0045                                 CellTracksVector<TrackerTraits>* cellTracks,
0046                                 HitsConstView<TrackerTraits> hh,
0047                                 OuterHitOfCell<TrackerTraits> isOuterHitOfCell,
0048                                 int nActualPairs,
0049                                 const int maxNumOfDoublets,
0050                                 CellCutsT<TrackerTraits>* const cuts) {
0051
0052     doubletsFromHisto<TrackerTraits>(
0053         nActualPairs, maxNumOfDoublets, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, *cuts);
0054   }
0055
0056 }  // namespace gpuPixelDoublets
0057
0058 #endif  // RecoTracker_PixelSeeding_plugins_gpuPixelDoublets_h