1 #ifndef RecoPixelVertexing_PixelVertexFinding_plugins_gpuClusterTracksDBSCAN_h
2 #define RecoPixelVertexing_PixelVertexFinding_plugins_gpuClusterTracksDBSCAN_h
13 namespace gpuVertexFinder {
31 auto& __restrict__
data = *pdata;
34 float const* __restrict__
zt = ws.zt;
35 float const* __restrict__
ezt2 = ws.ezt2;
40 uint8_t* __restrict__
izt = ws.izt;
41 int32_t* __restrict__
nn = data.ndof;
42 int32_t* __restrict__
iv = ws.iv;
56 printf(
"booked hist with %d bins, size %d for %d tracks\n", hist.nbins(), hist.capacity(),
nt);
63 int iz = int(zt[
i] * 10.);
64 iz = std::clamp(iz, INT8_MIN, INT8_MAX);
65 izt[
i] = iz - INT8_MIN;
66 assert(iz - INT8_MIN >= 0);
67 assert(iz - INT8_MIN < 256);
80 hist.fill(izt[
i], uint16_t(i));
88 auto loop = [&](uint32_t
j) {
98 cms::cuda::forEachInBins(hist, izt[i], 1,
loop);
108 auto loop = [&](uint32_t
j) {
120 cms::cuda::forEachInBins(hist, izt[i], 1,
loop);
129 assert(iv[iv[i]] !=
int(i));
148 assert(iv[iv[i]] !=
int(i));
158 assert(zt[iv[i]] <= zt[i]);
159 auto loop = [&](uint32_t
j) {
167 if (iv[i] != iv[
j]) {
168 printf(
"ERROR %d %d %f %f %d\n", i, iv[i], zt[i], zt[iv[i]], iv[iv[i]]);
169 printf(
" %d %d %f %f %d\n", j, iv[j], zt[j], zt[iv[j]], iv[iv[j]]);
174 cms::cuda::forEachInBins(hist, izt[i], 1,
loop);
185 auto loop = [&](uint32_t
j) {
191 if (dist * dist >
chi2max * (ezt2[i] + ezt2[
j]))
196 cms::cuda::forEachInBins(hist, izt[i], 1,
loop);
206 if (iv[i] ==
int(i)) {
208 auto old =
atomicInc(&foundClusters, 0xffffffff);
236 printf(
"found %d proto vertices\n", foundClusters);
241 #endif // RecoPixelVertexing_PixelVertexFinding_plugins_gpuClusterTracksDBSCAN_h
float const *__restrict__ ezt2
float const *__restrict__ zt
static constexpr uint32_t MAXVTX
WorkSpace int float float float chi2max
T1 atomicInc(T1 *a, T2 b)
printf("params %d %f %f %f\n", minT, eps, errmax, chi2max)
Abs< T >::type abs(const T &t)
typename Base::Counter Counter
__shared__ Hist::Counter hws[32]
uint8_t *__restrict__ izt
WorkSpace int float float errmax
static constexpr uint32_t totbins()
static constexpr uint32_t MAXTRACKS
__shared__ unsigned int foundClusters
uint32_t & nvIntermediate