diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc b/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc index f8cadf653114d..c07cb1eb2330c 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc +++ b/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc @@ -25,8 +25,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template class Kernel_BLFastFit { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, Tuples const *__restrict__ foundNtuplets, TupleMultiplicity const *__restrict__ tupleMultiplicity, TrackingRecHitSoAConstView hh, @@ -173,8 +172,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template struct Kernel_BLFit { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TupleMultiplicity const *__restrict__ tupleMultiplicity, double bField, OutputSoAView results_view, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h b/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h index 44b7bea5075fb..3908d38dab3b0 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h @@ -72,8 +72,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { } template - ALPAKA_FN_ACC ALPAKA_FN_INLINE __attribute__((always_inline)) int addOuterNeighbor( - const TAcc& acc, typename TrackerTraits::cindex_type t, CellNeighborsVector& cellNeighbors) { + ALPAKA_FN_ACC ALPAKA_FN_INLINE int addOuterNeighbor(TAcc const& acc, + typename TrackerTraits::cindex_type t, + CellNeighborsVector& cellNeighbors) { // use smart cache if (outerNeighbors().empty()) { auto i = cellNeighbors.extend(acc); // maybe wasted.... @@ -90,17 +91,16 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { (PtrAsInt)(&cellNeighbors[i]), alpaka::hierarchy::Blocks{}); // if fails we cannot give "i" back... #endif - } else + } else { return -1; + } } alpaka::mem_fence(acc, alpaka::memory_scope::Grid{}); return outerNeighbors().push_back(acc, t); } template - ALPAKA_FN_ACC ALPAKA_FN_INLINE __attribute__((always_inline)) int addTrack(TAcc const& acc, - tindex_type t, - CellTracksVector& cellTracks) { + ALPAKA_FN_ACC ALPAKA_FN_INLINE int addTrack(TAcc const& acc, tindex_type t, CellTracksVector& cellTracks) { if (tracks().empty()) { auto i = cellTracks.extend(acc); // maybe wasted.... if (i > 0) { @@ -282,8 +282,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { // trying to free the track building process from hardcoded layers, leaving // the visit of the graph based on the neighborhood connections between cells. - template - ALPAKA_FN_ACC ALPAKA_FN_INLINE void find_ntuplets(TAcc const& acc, + template + ALPAKA_FN_ACC ALPAKA_FN_INLINE void find_ntuplets(Acc1D const& acc, const HitsConstView& hh, CACellT* __restrict__ cells, CellTracksVector& cellTracks, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAFishbone.h b/RecoTracker/PixelSeeding/plugins/alpaka/CAFishbone.h index e7e1554fa7a5b..5f34641d8294a 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAFishbone.h +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAFishbone.h @@ -35,8 +35,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets { template class CAFishbone { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const& acc, + ALPAKA_FN_ACC void operator()(Acc2D const& acc, HitsConstView hh, CACellT* cells, uint32_t const* __restrict__ nCells, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc index 2d778f5e6e9de..1bd5666f2fed0 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc @@ -253,7 +253,6 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { const Vec2D blks{numberOfBlocks, 1u}; const Vec2D thrs{blockSize, stride}; const auto workDiv2D = cms::alpakatools::make_workdiv(blks, thrs); - alpaka::exec(queue, workDiv2D, CAFishbone{}, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h index e72d221f7e21c..1750a28d99666 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h @@ -73,8 +73,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_checkOverflows { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, TupleMultiplicity const *tupleMultiplicity, HitToTuple const *hitToTuple, @@ -168,8 +167,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_fishboneCleaner { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, CACellT const *cells, uint32_t const *__restrict__ nCells, TkSoAView tracks_view) const { @@ -191,8 +189,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_earlyDuplicateRemover { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, CACellT const *cells, uint32_t const *__restrict__ nCells, TkSoAView tracks_view, @@ -230,8 +227,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_fastDuplicateRemover { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, CACellT const *__restrict__ cells, uint32_t const *__restrict__ nCells, TkSoAView tracks_view, @@ -319,8 +315,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_connect { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc2D const &acc, cms::alpakatools::AtomicPairCounter *apc1, cms::alpakatools::AtomicPairCounter *apc2, // just to zero them HitsConstView hh, @@ -381,11 +376,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { } // loop on outer cells } }; + template class Kernel_find_ntuplets { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, HitsConstView hh, TkSoAView tracks_view, CACellT *__restrict__ cells, @@ -422,16 +417,16 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { typename Cell::TmpTuple stack; stack.reset(); bool bpix1Start = params.startAt0(pid); - thisCell.template find_ntuplets(acc, - hh, - cells, - *cellTracks, - tracks_view.hitIndices(), - *apc, - tracks_view.quality(), - stack, - params.minHitsPerNtuplet_, - bpix1Start); + thisCell.template find_ntuplets(acc, + hh, + cells, + *cellTracks, + tracks_view.hitIndices(), + *apc, + tracks_view.quality(), + stack, + params.minHitsPerNtuplet_, + bpix1Start); ALPAKA_ASSERT_ACC(stack.empty()); } } @@ -441,8 +436,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_mark_used { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, CACellT *__restrict__ cells, uint32_t const *nCells) const { using Cell = CACellT; @@ -457,8 +451,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_countMultiplicity { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, TupleMultiplicity *tupleMultiplicity) const { for (auto it : cms::alpakatools::uniform_elements(acc, tracks_view.hitIndices().nOnes())) { @@ -479,8 +472,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_fillMultiplicity { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, TupleMultiplicity *tupleMultiplicity) const { for (auto it : cms::alpakatools::uniform_elements(acc, tracks_view.hitIndices().nOnes())) { @@ -501,8 +493,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_classifyTracks { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, QualityCuts cuts) const { for (auto it : cms::alpakatools::uniform_elements(acc, tracks_view.hitIndices().nOnes())) { @@ -548,8 +539,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_doStatsForTracks { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView tracks_view, Counters *counters) const { + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, Counters *counters) const { for (auto idx : cms::alpakatools::uniform_elements(acc, tracks_view.hitIndices().nOnes())) { if (tracks_view.hitIndices().size(idx) == 0) break; //guard @@ -566,8 +556,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_countHitInTracks { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, HitToTuple *hitToTuple) const { for (auto idx : cms::alpakatools::uniform_elements(acc, tracks_view.hitIndices().nOnes())) { @@ -582,8 +571,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_fillHitInTracks { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, HitToTuple *hitToTuple) const { for (auto idx : cms::alpakatools::uniform_elements(acc, tracks_view.hitIndices().nOnes())) { @@ -598,8 +586,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_fillHitDetIndices { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, HitsConstView hh) const { // copy offsets @@ -617,8 +604,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_fillNLayers { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, cms::alpakatools::AtomicPairCounter *apc) const { // clamp the number of tracks to the capacity of the SoA @@ -636,8 +622,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_doStatsForHitInTracks { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, HitToTuple const *__restrict__ hitToTuple, Counters *counters) const { auto &c = *counters; @@ -654,8 +639,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_countSharedHit { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, int *__restrict__ nshared, HitContainer const *__restrict__ ptuples, Quality const *__restrict__ quality, @@ -693,8 +677,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_markSharedHit { - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, int const *__restrict__ nshared, HitContainer const *__restrict__ tuples, Quality *__restrict__ quality, @@ -721,8 +704,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_rejectDuplicate { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, uint16_t nmin, bool dupPassThrough, @@ -778,8 +760,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_sharedHitCleaner { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, HitsConstView hh, TkSoAView tracks_view, int nmin, @@ -829,11 +810,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { } } }; + template class Kernel_tripletCleaner { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, uint16_t nmin, bool dupPassThrough, @@ -892,8 +873,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_simpleTripletCleaner { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TkSoAView tracks_view, uint16_t nmin, bool dupPassThrough, @@ -938,8 +918,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { template class Kernel_print_found_ntuplets { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, HitsConstView hh, TkSoAView tracks_view, HitToTuple const *__restrict__ phitToTuple, @@ -980,8 +959,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { class Kernel_printCounters { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, Counters const *counters) const { + ALPAKA_FN_ACC void operator()(Acc1D const &acc, Counters const *counters) const { auto const &c = *counters; printf( "||Counters | nEvents | nHits | nCells | nTuples | nFitTacks | nLooseTracks | nGoodTracks | nUsedHits | " diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoublets.h b/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoublets.h index aff544bb556cb..5360ab3c877fc 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoublets.h +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoublets.h @@ -18,8 +18,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template class InitDoublets { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const& acc, + ALPAKA_FN_ACC void operator()(Acc1D const& acc, OuterHitOfCell* isOuterHitOfCell, int nHits, CellNeighborsVector* cellNeighbors, @@ -51,11 +50,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template class GetDoubletsFromHisto { public: - template >> - // #ifdef __CUDACC__ - // __launch_bounds__(getDoubletsFromHistoMaxBlockSize, getDoubletsFromHistoMinBlocksPerMP) // TODO: Alapakify - // #endif - ALPAKA_FN_ACC void operator()(TAcc const& acc, + // __launch_bounds__(getDoubletsFromHistoMaxBlockSize, getDoubletsFromHistoMinBlocksPerMP) // TODO: Alapakify + ALPAKA_FN_ACC void operator()(Acc2D const& acc, CACellT* cells, uint32_t* nCells, CellNeighborsVector* cellNeighbors, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoubletsAlgos.h b/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoubletsAlgos.h index 11363c4e2d58a..a4416ba159978 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoubletsAlgos.h +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAPixelDoubletsAlgos.h @@ -74,10 +74,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets { int phiCuts[T::nPairs]; template - ALPAKA_FN_ACC ALPAKA_FN_INLINE bool __attribute__((always_inline)) zSizeCut(const TAcc& acc, - H hh, - int i, - int o) const { + ALPAKA_FN_ACC ALPAKA_FN_INLINE bool zSizeCut(TAcc const& acc, H hh, int i, int o) const { const uint32_t mi = hh[i].detectorIndex(); bool innerB1 = mi < T::last_bpix1_detIndex; @@ -105,9 +102,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets { } template - ALPAKA_FN_ACC ALPAKA_FN_INLINE bool __attribute__((always_inline)) clusterCut(const TAcc& acc, - H hh, - uint32_t i) const { + ALPAKA_FN_ACC ALPAKA_FN_INLINE bool clusterCut(TAcc const& acc, H hh, uint32_t i) const { const uint32_t mi = hh[i].detectorIndex(); bool innerB1orB2 = mi < T::last_bpix2_detIndex; @@ -131,7 +126,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets { }; template - ALPAKA_FN_ACC ALPAKA_FN_INLINE void __attribute__((always_inline)) doubletsFromHisto( + ALPAKA_FN_ACC ALPAKA_FN_INLINE void doubletsFromHisto( const TAcc& acc, uint32_t nPairs, const uint32_t maxNumOfDoublets, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/RiemannFit.dev.cc b/RecoTracker/PixelSeeding/plugins/alpaka/RiemannFit.dev.cc index 3e6f6e9c8ed98..ab4033d6ccd54 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/RiemannFit.dev.cc +++ b/RecoTracker/PixelSeeding/plugins/alpaka/RiemannFit.dev.cc @@ -26,8 +26,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template class Kernel_FastFit { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, Tuples const *__restrict__ foundNtuplets, TupleMultiplicity const *__restrict__ tupleMultiplicity, uint32_t nHits, @@ -40,7 +39,6 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { constexpr uint32_t hitsInFit = N; ALPAKA_ASSERT_ACC(hitsInFit <= nHits); - ALPAKA_ASSERT_ACC(pfast_fit); ALPAKA_ASSERT_ACC(foundNtuplets); ALPAKA_ASSERT_ACC(tupleMultiplicity); @@ -95,8 +93,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template class Kernel_CircleFit { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TupleMultiplicity const *__restrict__ tupleMultiplicity, uint32_t nHits, double bField, @@ -141,8 +138,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { template class Kernel_LineFit { public: - template >> - ALPAKA_FN_ACC void operator()(TAcc const &acc, + ALPAKA_FN_ACC void operator()(Acc1D const &acc, TupleMultiplicity const *__restrict__ tupleMultiplicity, uint32_t nHits, double bField,