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> 39 template <
typename TrackerTraits>
42 template <
typename TrackerTraits>
45 template <
typename TrackerTraits>
48 template <
typename TrackerTraits>
51 template <
typename TrackerTraits>
56 template <
typename TrackerTraits>
59 template <
typename TrackerTraits>
62 template <
typename TrackerTraits>
65 template <
typename TrackerTraits>
68 template <
typename TrackerTraits>
73 template <
typename TrackerTraits>
76 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
83 uint32_t
const *__restrict__
nCells,
106 printf(
"number of found cells %d \n found tuples %d with total hits %d out of %d\n",
111 if (
apc->
get().first < TrackerTraits::maxNumberOfQuadruplets) {
118 if (
tracks_view.hitIndices().size(
idx) > TrackerTraits::maxHitsOnTrack)
127 if (
apc->
get().first >= TrackerTraits::maxNumberOfQuadruplets)
128 printf(
"Tuples overflow\n");
130 printf(
"Cells overflow\n");
134 printf(
"cellTracks overflow\n");
136 printf(
"ERROR hitToTuple overflow %d %d\n",
hitToTuple->nOnes(),
nHits);
138 printf(
"size of cellNeighbors %d \n cellTracks %d \n hitToTuple %d \n",
147 if (thisCell.hasFishbone() && !thisCell.isKilled())
149 if (thisCell.outerNeighbors().full())
150 printf(
"OuterNeighbors overflow %d in %d\n",
idx, thisCell.layerPairId());
151 if (thisCell.tracks().full())
152 printf(
"Tracks overflow %d in %d\n",
idx, thisCell.layerPairId());
153 if (thisCell.isKilled())
155 if (!thisCell.unused())
157 if ((0 ==
hitToTuple->size(thisCell.inner_hit_id())) && (0 ==
hitToTuple->size(thisCell.outer_hit_id())))
163 if ((*isOuterHitOfCell).container[
idx].full())
164 printf(
"OuterHitOfCell overflow %d\n",
idx);
168 template <
typename TrackerTraits>
171 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
174 uint32_t
const *__restrict__
nCells,
180 if (!thisCell.isKilled())
183 for (
auto it : thisCell.tracks())
191 template <
typename TrackerTraits>
194 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
197 uint32_t
const *__restrict__
nCells,
206 if (thisCell.tracks().size() < 2)
212 for (
auto it : thisCell.tracks()) {
221 for (
auto it : thisCell.tracks()) {
230 template <
typename TrackerTraits>
233 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
236 uint32_t
const *__restrict__
nCells,
244 const auto ntNCells = (*nCells);
248 if (thisCell.tracks().size() < 2)
257 int ntr = thisCell.tracks().size();
258 for (
int i = 0;
i < ntr - 1; ++
i) {
259 auto it = thisCell.tracks()[
i];
267 for (
auto j =
i + 1;
j < ntr; ++
j) {
268 auto jt = thisCell.tracks()[
j];
275 if ((cti - ctj) * (cti - ctj) > dct)
278 if ((opi - opj) * (opi - opj) > dop)
291 for (
auto it : thisCell.tracks()) {
296 if (maxQual <=
loose)
300 for (
auto it : thisCell.tracks()) {
311 for (
auto it : thisCell.tracks()) {
319 template <
typename TrackerTraits>
322 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
341 auto cellIndex =
idx;
343 auto innerHitId = thisCell.inner_hit_id();
347 uint32_t numberOfPossibleNeighbors = (*isOuterHitOfCell)[innerHitId].size();
348 auto vi = (*isOuterHitOfCell)[innerHitId].data();
349 auto ri = thisCell.inner_r(
hh);
350 auto zi = thisCell.inner_z(
hh);
351 auto ro = thisCell.outer_r(
hh);
352 auto zo = thisCell.outer_z(
hh);
353 auto isBarrel = thisCell.inner_detIndex(
hh) < TrackerTraits::last_barrel_detIndex;
357 auto otherCell = (vi[
j]);
358 auto &oc =
cells[otherCell];
359 auto r1 = oc.inner_r(
hh);
360 auto z1 = oc.inner_z(
hh);
361 bool aligned = Cell::areAlignedRZ(
373 oc.inner_detIndex(
hh) < TrackerTraits::last_bpix1_detIndex ?
params.dcaCutInnerTriplet_
374 :
params.dcaCutOuterTriplet_,
377 thisCell.setStatusBits(Cell::StatusBit::kUsed);
378 oc.setStatusBits(Cell::StatusBit::kUsed);
384 template <
typename TrackerTraits>
387 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
402 printf(
"starting producing ntuplets from %d cells \n", *
nCells);
409 if (thisCell.isKilled())
413 if (thisCell.outerNeighbors().empty())
416 auto pid = thisCell.layerPairId();
417 bool doit =
params.startingLayerPair(pid);
422 typename Cell::TmpTuple
stack;
424 bool bpix1Start =
params.startAt0(pid);
425 thisCell.template find_ntuplets<maxDepth, TAcc>(acc,
433 params.minHitsPerNtuplet_,
441 template <
typename TrackerTraits>
444 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
447 uint32_t
const *
nCells)
const {
451 if (!thisCell.tracks().empty())
452 thisCell.setStatusBits(Cell::StatusBit::kInTrack);
457 template <
typename TrackerTraits>
460 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
471 if (
nhits > TrackerTraits::maxHitsOnTrack)
472 printf(
"wrong mult %d %d\n",
it,
nhits);
479 template <
typename TrackerTraits>
482 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
493 if (
nhits > TrackerTraits::maxHitsOnTrack)
494 printf(
"wrong mult %d %d\n",
it,
nhits);
501 template <
typename TrackerTraits>
504 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
525 for (
int i = 0;
i < 5; ++
i) {
548 template <
typename TrackerTraits>
551 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
566 template <
typename TrackerTraits>
569 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
582 template <
typename TrackerTraits>
585 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
598 template <
typename TrackerTraits>
601 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
617 template <
typename TrackerTraits>
620 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
636 template <
typename TrackerTraits>
639 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
654 template <
typename TrackerTraits>
657 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
659 int *__restrict__ nshared,
694 template <
typename TrackerTraits>
696 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
698 int const *__restrict__ nshared,
714 if (nshared[
idx] > 2)
721 template <
typename TrackerTraits>
724 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
760 if ((cti - ctj) * (cti - ctj) > dct)
763 if ((opi - opj) * (opi - opj) > dop)
766 if (nlj < nli || (nlj == nli && (qj < qi || (qj == qi &&
score(
it, nli) <
score(jt, nlj)))))
778 template <
typename TrackerTraits>
781 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
795 uint32_t
l1end =
hh.hitsLayerStart()[1];
823 if (idx < l1end and nl >
nmin)
832 template <
typename TrackerTraits>
835 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
844 auto const good = Quality::strict;
854 bool onlyTriplets =
true;
892 template <
typename TrackerTraits>
895 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
938 template <
typename TrackerTraits>
941 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
958 printf(
"TK: %d %d %d %d %f %f %f %f %f %f %f %.3f %.3f %.3f %.3f %.3f %.3f %.3f\n",
983 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
987 "||Counters | nEvents | nHits | nCells | nTuples | nFitTacks | nLooseTracks | nGoodTracks | nUsedHits | " 988 "nDupHits | nFishCells | nKilledCells | nUsedCells | nZeroTrackCells ||\n");
989 printf(
"Counters Raw %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld %lld\n",
1004 "Counters Norm %lld || %.1f| %.1f| %.1f| %.1f| %.1f| %.1f| %.1f| %.1f| %.3f| %.3f| %.3f| " 1007 c.nHits /
double(
c.nEvents),
1008 c.nCells /
double(
c.nEvents),
1009 c.nTuples /
double(
c.nEvents),
1010 c.nFitTracks /
double(
c.nEvents),
1011 c.nLooseTracks /
double(
c.nEvents),
1012 c.nGoodTracks /
double(
c.nEvents),
1013 c.nUsedHits /
double(
c.nEvents),
1014 c.nDupHits /
double(
c.nEvents),
1015 c.nFishCells /
double(
c.nCells),
1016 c.nKilledCells /
double(
c.nCells),
1017 c.nEmptyCells /
double(
c.nCells),
1018 c.nZeroTrackCells /
double(
c.nCells));
1024 #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
constexpr bool isNotFinite(T x)
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