1 #ifndef RecoVertex_PixelVertexFinding_clusterTracksIterativeAlpaka_h 2 #define RecoVertex_PixelVertexFinding_clusterTracksIterativeAlpaka_h 8 #include <alpaka/alpaka.hpp> 44 float const* __restrict__
zt =
ws.zt();
45 float const* __restrict__
ezt2 =
ws.ezt2();
46 uint8_t* __restrict__
izt =
ws.izt();
47 int32_t* __restrict__
iv =
ws.iv();
48 int32_t* __restrict__
nn =
trkdata.ndof();
56 auto&
hist = alpaka::declareSharedVar<Hist, __COUNTER__>(acc);
57 auto&
hws = alpaka::declareSharedVar<Hist::Counter[32], __COUNTER__>(acc);
65 alpaka::syncBlockThreads(acc);
69 printf(
"booked hist with %d bins, size %d for %d tracks\n",
hist.nbins(),
hist.capacity(),
nt);
78 int iz =
static_cast<int>(
zt[
i] * 10.f);
79 iz = std::clamp(iz, INT8_MIN, INT8_MAX);
80 izt[
i] = iz - INT8_MIN;
87 alpaka::syncBlockThreads(acc);
90 alpaka::syncBlockThreads(acc);
96 alpaka::syncBlockThreads(acc);
101 if (
ezt2[
i] > errmax2)
115 auto&
nloops = alpaka::declareSharedVar<int, __COUNTER__>(acc);
117 alpaka::syncBlockThreads(acc);
121 while (alpaka::syncBlockThreadsPredicate<alpaka::BlockOr>(acc,
more)) {
132 auto p =
hist.begin() +
k;
180 auto&
foundClusters = alpaka::declareSharedVar<unsigned int, __COUNTER__>(acc);
182 alpaka::syncBlockThreads(acc);
187 if (
iv[
i] ==
int(
i)) {
196 alpaka::syncBlockThreads(acc);
207 alpaka::syncBlockThreads(acc);
226 #endif // RecoTracker_PixelVertexFinding_plugins_clusterTracksIterativeAlpaka_h
ALPAKA_ASSERT_ACC(nvFinal<=nvIntermediate)
auto &__restrict__ trkdata
float const *__restrict__ zt
cms::cuda::HistoContainer< uint8_t, 256, 16000, 8, uint16_t > Hist
T1 atomicInc(T1 *a, T2 b)
Abs< T >::type abs(const T &t)
float const *__restrict__ ezt2
__shared__ Hist::Counter hws[32]
::vertexFinder::PixelVertexWorkSpaceSoAView WsSoAView
T1 atomicMin(T1 *a, T2 b)
ALPAKA_FN_ACC void operator()(Acc1D const &acc, VtxSoAView data, TrkSoAView trkdata, WsSoAView ws, int minT, float eps, float errmax, float chi2max) const