1 #ifndef RecoVertex_PixelVertexFinding_plugins_alpaka_clusterTracksByDensity_h 2 #define RecoVertex_PixelVertexFinding_plugins_alpaka_clusterTracksByDensity_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.totbins(),
hist.capacity(),
nt);
78 int iz =
static_cast<int>(
zt[
i] * 10.f);
79 iz = std::clamp(iz, INT8_MIN, INT8_MAX);
82 izt[
i] = iz - INT8_MIN;
87 alpaka::syncBlockThreads(acc);
90 alpaka::syncBlockThreads(acc);
96 alpaka::syncBlockThreads(acc);
101 if (
ezt2[
i] > errmax2)
114 alpaka::syncBlockThreads(acc);
133 alpaka::syncBlockThreads(acc);
141 alpaka::syncBlockThreads(acc);
153 alpaka::syncBlockThreads(acc);
184 alpaka::syncBlockThreads(acc);
187 auto&
foundClusters = alpaka::declareSharedVar<unsigned int, __COUNTER__>(acc);
189 alpaka::syncBlockThreads(acc);
194 if (
iv[
i] ==
int(
i)) {
203 alpaka::syncBlockThreads(acc);
214 alpaka::syncBlockThreads(acc);
247 #endif // RecoVertex_PixelVertexFinding_plugins_alpaka_clusterTracksByDensity_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
ALPAKA_FN_ACC ALPAKA_FN_INLINE void clusterTracksByDensity(Acc1D const &acc, VtxSoAView &data, TrkSoAView &trkdata, WsSoAView &ws, int minT, float eps, float errmax, float chi2max)
ALPAKA_FN_ACC void operator()(Acc1D const &acc, VtxSoAView data, TrkSoAView trkdata, WsSoAView ws, int minT, float eps, float errmax, float chi2max) const