plugins/alpaka/CAPixelDoublets.h

0001 #ifndef RecoTracker_PixelSeeding_plugins_alpaka_CAPixelDoublets_h
0002 #define RecoTracker_PixelSeeding_plugins_alpaka_CAPixelDoublets_h
0003
0004 #include <type_traits>
0005
0006 #include <alpaka/alpaka.hpp>
0007
0008 #include "HeterogeneousCore/AlpakaInterface/interface/config.h"
0009 #include "HeterogeneousCore/AlpakaInterface/interface/workdivision.h"
0010
0011 #include "CAPixelDoubletsAlgos.h"
0012
0013 namespace ALPAKA_ACCELERATOR_NAMESPACE {
0014   using namespace alpaka;
0015   using namespace cms::alpakatools;
0016   namespace caPixelDoublets {
0017
0018     template <typename TrackerTraits>
0019     class InitDoublets {
0020     public:
0021       template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
0022       ALPAKA_FN_ACC void operator()(TAcc const& acc,
0023                                     OuterHitOfCell<TrackerTraits>* isOuterHitOfCell,
0024                                     int nHits,
0025                                     CellNeighborsVector<TrackerTraits>* cellNeighbors,
0026                                     CellNeighbors<TrackerTraits>* cellNeighborsContainer,
0027                                     CellTracksVector<TrackerTraits>* cellTracks,
0028                                     CellTracks<TrackerTraits>* cellTracksContainer) const {
0029         ALPAKA_ASSERT_ACC((*isOuterHitOfCell).container);
0030
0031         for (auto i : cms::alpakatools::uniform_elements(acc, nHits - isOuterHitOfCell->offset))
0032           (*isOuterHitOfCell).container[i].reset();
0033
0034         if (cms::alpakatools::once_per_grid(acc)) {
0035           cellNeighbors->construct(TrackerTraits::maxNumOfActiveDoublets, cellNeighborsContainer);
0036           cellTracks->construct(TrackerTraits::maxNumOfActiveDoublets, cellTracksContainer);
0037           [[maybe_unused]] auto i = cellNeighbors->extend(acc);
0038           ALPAKA_ASSERT_ACC(0 == i);
0039           (*cellNeighbors)[0].reset();
0040           i = cellTracks->extend(acc);
0041           ALPAKA_ASSERT_ACC(0 == i);
0042           (*cellTracks)[0].reset();
0043         }
0044       }
0045     };
0046
0047     // Not used for the moment, see below.
0048     //constexpr auto getDoubletsFromHistoMaxBlockSize = 64;  // for both x and y
0049     //constexpr auto getDoubletsFromHistoMinBlocksPerMP = 16;
0050
0051     template <typename TrackerTraits>
0052     class GetDoubletsFromHisto {
0053     public:
0054       template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
0055       // #ifdef __CUDACC__
0056       //       __launch_bounds__(getDoubletsFromHistoMaxBlockSize, getDoubletsFromHistoMinBlocksPerMP)  // TODO: Alapakify
0057       // #endif
0058       ALPAKA_FN_ACC void operator()(TAcc const& acc,
0059                                     CACellT<TrackerTraits>* cells,
0060                                     uint32_t* nCells,
0061                                     CellNeighborsVector<TrackerTraits>* cellNeighbors,
0062                                     CellTracksVector<TrackerTraits>* cellTracks,
0063                                     HitsConstView<TrackerTraits> hh,
0064                                     OuterHitOfCell<TrackerTraits>* isOuterHitOfCell,
0065                                     uint32_t nActualPairs,
0066                                     const uint32_t maxNumOfDoublets,
0067                                     CellCutsT<TrackerTraits> cuts) const {
0068         doubletsFromHisto<TrackerTraits>(
0069             acc, nActualPairs, maxNumOfDoublets, cells, nCells, cellNeighbors, cellTracks, hh, *isOuterHitOfCell, cuts);
0070       }
0071     };
0072
0073   }  // namespace caPixelDoublets
0074
0075 }  // namespace ALPAKA_ACCELERATOR_NAMESPACE
0076
0077 #endif  // RecoTracker_PixelSeeding_plugins_alpaka_CAPixelDoublets_h