1 #ifndef RecoTracker_PixelSeeding_plugins_alpaka_CAHitNtupletGeneratorKernelsImpl_h 2 #define RecoTracker_PixelSeeding_plugins_alpaka_CAHitNtupletGeneratorKernelsImpl_h 12 #include <type_traits> 15 #include <alpaka/alpaka.hpp> 38 template <
typename TrackerTraits>
41 template <
typename TrackerTraits>
44 template <
typename TrackerTraits>
47 template <
typename TrackerTraits>
50 template <
typename TrackerTraits>
55 template <
typename TrackerTraits>
58 template <
typename TrackerTraits>
61 template <
typename TrackerTraits>
64 template <
typename TrackerTraits>
67 template <
typename TrackerTraits>
72 template <
typename TrackerTraits>
75 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
82 uint32_t
const *__restrict__
nCells,
105 printf(
"number of found cells %d \n found tuples %d with total hits %d out of %d\n",
110 if (
apc->
get().first < TrackerTraits::maxNumberOfQuadruplets) {
117 if (
tracks_view.hitIndices().size(
idx) > TrackerTraits::maxHitsOnTrack)
126 if (
apc->
get().first >= TrackerTraits::maxNumberOfQuadruplets)
127 printf(
"Tuples overflow\n");
129 printf(
"Cells overflow\n");
133 printf(
"cellTracks overflow\n");
135 printf(
"ERROR hitToTuple overflow %d %d\n",
hitToTuple->nOnes(),
nHits);
137 printf(
"size of cellNeighbors %d \n cellTracks %d \n hitToTuple %d \n",
146 if (thisCell.hasFishbone() && !thisCell.isKilled())
148 if (thisCell.outerNeighbors().full())
149 printf(
"OuterNeighbors overflow %d in %d\n",
idx, thisCell.layerPairId());
150 if (thisCell.tracks().full())
151 printf(
"Tracks overflow %d in %d\n",
idx, thisCell.layerPairId());
152 if (thisCell.isKilled())
154 if (!thisCell.unused())
156 if ((0 ==
hitToTuple->size(thisCell.inner_hit_id())) && (0 ==
hitToTuple->size(thisCell.outer_hit_id())))
162 if ((*isOuterHitOfCell).container[
idx].full())
163 printf(
"OuterHitOfCell overflow %d\n",
idx);
167 template <
typename TrackerTraits>
170 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
173 uint32_t
const *__restrict__
nCells,
179 if (!thisCell.isKilled())
182 for (
auto it : thisCell.tracks())
190 template <
typename TrackerTraits>
193 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
196 uint32_t
const *__restrict__
nCells,
205 if (thisCell.tracks().size() < 2)
211 for (
auto it : thisCell.tracks()) {
220 for (
auto it : thisCell.tracks()) {
229 template <
typename TrackerTraits>
232 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
235 uint32_t
const *__restrict__
nCells,
243 const auto ntNCells = (*nCells);
247 if (thisCell.tracks().size() < 2)
256 int ntr = thisCell.tracks().size();
257 for (
int i = 0;
i < ntr - 1; ++
i) {
258 auto it = thisCell.tracks()[
i];
266 for (
auto j =
i + 1;
j < ntr; ++
j) {
267 auto jt = thisCell.tracks()[
j];
274 if ((cti - ctj) * (cti - ctj) > dct)
277 if ((opi - opj) * (opi - opj) > dop)
290 for (
auto it : thisCell.tracks()) {
295 if (maxQual <=
loose)
299 for (
auto it : thisCell.tracks()) {
310 for (
auto it : thisCell.tracks()) {
318 template <
typename TrackerTraits>
321 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
340 auto cellIndex =
idx;
342 auto innerHitId = thisCell.inner_hit_id();
346 uint32_t numberOfPossibleNeighbors = (*isOuterHitOfCell)[innerHitId].size();
347 auto vi = (*isOuterHitOfCell)[innerHitId].data();
348 auto ri = thisCell.inner_r(
hh);
349 auto zi = thisCell.inner_z(
hh);
350 auto ro = thisCell.outer_r(
hh);
351 auto zo = thisCell.outer_z(
hh);
352 auto isBarrel = thisCell.inner_detIndex(
hh) < TrackerTraits::last_barrel_detIndex;
356 auto otherCell = (vi[
j]);
357 auto &oc =
cells[otherCell];
358 auto r1 = oc.inner_r(
hh);
359 auto z1 = oc.inner_z(
hh);
360 bool aligned = Cell::areAlignedRZ(
372 oc.inner_detIndex(
hh) < TrackerTraits::last_bpix1_detIndex ?
params.dcaCutInnerTriplet_
373 :
params.dcaCutOuterTriplet_,
376 thisCell.setStatusBits(Cell::StatusBit::kUsed);
377 oc.setStatusBits(Cell::StatusBit::kUsed);
383 template <
typename TrackerTraits>
386 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
401 printf(
"starting producing ntuplets from %d cells \n", *
nCells);
408 if (thisCell.isKilled())
412 if (thisCell.outerNeighbors().empty())
415 auto pid = thisCell.layerPairId();
416 bool doit =
params.startingLayerPair(pid);
421 typename Cell::TmpTuple
stack;
423 bool bpix1Start =
params.startAt0(pid);
424 thisCell.template find_ntuplets<maxDepth, TAcc>(acc,
432 params.minHitsPerNtuplet_,
440 template <
typename TrackerTraits>
443 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
446 uint32_t
const *
nCells)
const {
450 if (!thisCell.tracks().empty())
451 thisCell.setStatusBits(Cell::StatusBit::kInTrack);
456 template <
typename TrackerTraits>
459 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
470 if (
nhits > TrackerTraits::maxHitsOnTrack)
471 printf(
"wrong mult %d %d\n",
it,
nhits);
478 template <
typename TrackerTraits>
481 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
492 if (
nhits > TrackerTraits::maxHitsOnTrack)
493 printf(
"wrong mult %d %d\n",
it,
nhits);
500 template <
typename TrackerTraits>
503 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
524 for (
int i = 0;
i < 5; ++
i) {
547 template <
typename TrackerTraits>
550 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
565 template <
typename TrackerTraits>
568 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
581 template <
typename TrackerTraits>
584 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
597 template <
typename TrackerTraits>
600 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
616 template <
typename TrackerTraits>
619 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
635 template <
typename TrackerTraits>
638 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
653 template <
typename TrackerTraits>
656 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
658 int *__restrict__ nshared,
693 template <
typename TrackerTraits>
695 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
697 int const *__restrict__ nshared,
713 if (nshared[
idx] > 2)
720 template <
typename TrackerTraits>
723 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
759 if ((cti - ctj) * (cti - ctj) > dct)
762 if ((opi - opj) * (opi - opj) > dop)
765 if (nlj < nli || (nlj == nli && (qj < qi || (qj == qi &&
score(
it, nli) <
score(jt, nlj)))))
777 template <
typename TrackerTraits>
780 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
794 uint32_t
l1end =
hh.hitsLayerStart()[1];
822 if (idx < l1end and nl >
nmin)
831 template <
typename TrackerTraits>
834 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
843 auto const good = Quality::strict;
853 bool onlyTriplets =
true;
891 template <
typename TrackerTraits>
894 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
937 template <
typename TrackerTraits>
940 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
957 printf(
"TK: %d %d %d %d %f %f %f %f %f %f %f %.3f %.3f %.3f %.3f %.3f %.3f %.3f\n",
982 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
986 "||Counters | nEvents | nHits | nCells | nTuples | nFitTacks | nLooseTracks | nGoodTracks | nUsedHits | " 987 "nDupHits | nFishCells | nKilledCells | nUsedCells | nZeroTrackCells ||\n");
988 printf(
"Counters Raw %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld\n",
1003 "Counters Norm %lld || %.1f| %.1f| %.1f| %.1f| %.1f| %.1f| %.1f| %.1f| %.3f| %.3f| %.3f| " 1006 c.nHits /
double(
c.nEvents),
1007 c.nCells /
double(
c.nEvents),
1008 c.nTuples /
double(
c.nEvents),
1009 c.nFitTracks /
double(
c.nEvents),
1010 c.nLooseTracks /
double(
c.nEvents),
1011 c.nGoodTracks /
double(
c.nEvents),
1012 c.nUsedHits /
double(
c.nEvents),
1013 c.nDupHits /
double(
c.nEvents),
1014 c.nFishCells /
double(
c.nCells),
1015 c.nKilledCells /
double(
c.nCells),
1016 c.nEmptyCells /
double(
c.nCells),
1017 c.nZeroTrackCells /
double(
c.nCells));
1023 #endif // RecoTracker_PixelSeeding_plugins_alpaka_CAHitNtupletGeneratorKernelsImpl_h HitContainer< TrackerTraits > const *__restrict__ tuples
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const isOuterHitOfCell
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, cms::alpakatools::AtomicPairCounter *apc) const
auto const good
min quality of good
constexpr uint32_t tkNotFound
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const * cellNeighbors
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const * cellTracks
uint32_t const *__restrict__ TkSoAView< TrackerTraits > tracks_view
ALPAKA_FN_ACC void operator()(TAcc const &acc, cms::alpakatools::AtomicPairCounter *apc1, cms::alpakatools::AtomicPairCounter *apc2, HitsConstView< TrackerTraits > hh, CACellT< TrackerTraits > *cells, uint32_t *nCells, CellNeighborsVector< TrackerTraits > *cellNeighbors, OuterHitOfCell< TrackerTraits > const *isOuterHitOfCell, CAParams< TrackerTraits > params) const
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE constexpr float zip(ConstView const &tracks, int32_t i)
caHitNtupletGenerator::Counters Counters
typename reco::TrackSoA< TrackerTraits >::HitContainer HitContainer
ALPAKA_FN_ACC void operator()(TAcc const &acc, HitToTuple< TrackerTraits > const *__restrict__ hitToTuple, Counters *counters) const
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t int32_t int iev
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ cells
TupleMultiplicity< TrackerTraits > const * tupleMultiplicity
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, TupleMultiplicity< TrackerTraits > const *tupleMultiplicity, HitToTuple< TrackerTraits > const *hitToTuple, cms::alpakatools::AtomicPairCounter *apc, CACellT< TrackerTraits > const *__restrict__ cells, uint32_t const *__restrict__ nCells, CellNeighborsVector< TrackerTraits > const *cellNeighbors, CellTracksVector< TrackerTraits > const *cellTracks, OuterHitOfCell< TrackerTraits > const *isOuterHitOfCell, int32_t nHits, uint32_t maxNumberOfDoublets, Counters *counters) const
uint32_t const *__restrict__ TkSoAView< TrackerTraits > bool dupPassThrough
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE constexpr float charge(ConstView const &tracks, int32_t i)
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, HitToTuple< TrackerTraits > *hitToTuple) const
reco::TrackSoAView< TrackerTraits > TkSoAView
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const int32_t uint32_t maxNumberOfDoublets
TkSoAView< TrackerTraits > GPUCACellT< TrackerTraits > *__restrict__ uint32_t const CellTracksVector< TrackerTraits > cms::cuda::AtomicPairCounter CAParams< TrackerTraits > params
HitsConstView< TrackerTraits > hh
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE constexpr float tip(ConstView const &tracks, int32_t i)
ALPAKA_FN_ACC void operator()(TAcc const &acc, CACellT< TrackerTraits > const *__restrict__ cells, uint32_t const *__restrict__ nCells, TkSoAView< TrackerTraits > tracks_view, bool dupPassThrough) const
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const int32_t uint32_t Counters * counters
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, uint16_t nmin, bool dupPassThrough, HitToTuple< TrackerTraits > const *__restrict__ phitToTuple) const
ALPAKA_FN_ACC void operator()(TAcc const &acc, HitsConstView< TrackerTraits > hh, TkSoAView< TrackerTraits > tracks_view, CACellT< TrackerTraits > *__restrict__ cells, uint32_t const *nCells, CellTracksVector< TrackerTraits > *cellTracks, cms::alpakatools::AtomicPairCounter *apc, CAParams< TrackerTraits > params) const
static constexpr __host__ __device__ int computeNumberOfLayers(const TrackSoAConstView &tracks, int32_t i)
ALPAKA_FN_ACC void operator()(TAcc const &acc, HitsConstView< TrackerTraits > hh, TkSoAView< TrackerTraits > tracks_view, int nmin, bool dupPassThrough, HitToTuple< TrackerTraits > const *__restrict__ phitToTuple) const
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t firstPrint
__device__ __host__ Counters get() const
QualityCuts< TrackerTraits > cuts
std::vector< Block > Blocks
Abs< T >::type abs(const T &t)
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE constexpr float phi(ConstView const &tracks, int32_t i)
typename CACellT< TrackerTraits >::HitsConstView HitsConstView
ALPAKA_FN_ACC void operator()(TAcc const &acc, CACellT< TrackerTraits > *__restrict__ cells, uint32_t const *nCells) const
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const * hitToTuple
ALPAKA_FN_ACC void operator()(TAcc const &acc, int *__restrict__ nshared, HitContainer< TrackerTraits > const *__restrict__ ptuples, Quality const *__restrict__ quality, HitToTuple< TrackerTraits > const *__restrict__ phitToTuple) const
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE constexpr bool isTriplet(ConstView const &tracks, int32_t i)
ALPAKA_FN_ACC void operator()(TAcc const &acc, int const *__restrict__ nshared, HitContainer< TrackerTraits > const *__restrict__ tuples, Quality *__restrict__ quality, bool dupPassThrough) const
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, uint16_t nmin, bool dupPassThrough, HitToTuple< TrackerTraits > const *__restrict__ phitToTuple) const
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ nCells
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, TupleMultiplicity< TrackerTraits > *tupleMultiplicity) const
auto const & foundNtuplets
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, Counters *counters) const
ALPAKA_FN_ACC void operator()(TAcc const &acc, CACellT< TrackerTraits > const *cells, uint32_t const *__restrict__ nCells, TkSoAView< TrackerTraits > tracks_view, bool dupPassThrough) const
ALPAKA_FN_ACC void operator()(TAcc const &acc, CACellT< TrackerTraits > const *cells, uint32_t const *__restrict__ nCells, TkSoAView< TrackerTraits > tracks_view) const
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, HitToTuple< TrackerTraits > *hitToTuple) const
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter * apc
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, QualityCuts< TrackerTraits > cuts) const
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t int32_t lastPrint
ALPAKA_FN_ACC void operator()(TAcc const &acc, Counters const *counters) const
ALPAKA_FN_ACC void operator()(TAcc const &acc, HitsConstView< TrackerTraits > hh, TkSoAView< TrackerTraits > tracks_view, HitToTuple< TrackerTraits > const *__restrict__ phitToTuple, int32_t firstPrint, int32_t lastPrint, int iev) const
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t nHits
The Signals That Services Can Subscribe To This is based on ActivityRegistry h
Helper function to determine trigger accepts.
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, HitsConstView< TrackerTraits > hh) const
pixelTrack::Quality Quality
ALPAKA_ASSERT_ACC(offsets)
constexpr uint32_t tkNotFound
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, uint16_t nmin, bool dupPassThrough, HitToTuple< TrackerTraits > const *__restrict__ phitToTuple) const
TrackingRecHitSoAConstView< TrackerTraits > HitsConstView
ALPAKA_FN_ACC void operator()(TAcc const &acc, TkSoAView< TrackerTraits > tracks_view, TupleMultiplicity< TrackerTraits > *tupleMultiplicity) const
HitContainer< TrackerTraits > const *__restrict__ Quality const *__restrict__ HitToTuple< TrackerTraits > const *__restrict__ phitToTuple
T1 atomicAdd(T1 *a, T2 b)
HitContainer< TrackerTraits > const *__restrict__ ptuples
typename reco::TrackSoA< TrackerTraits >::template Layout<>::View TrackSoAView