Classes
class	Producer

struct	WorkSpace

Typedefs
using	Hist = cms::cuda::HistoContainer< uint8_t, 256, 16000, 8, uint16_t >

using	TkSoA = pixelTrack::TrackSoA

using	ZVertices = ZVertexSoA

Functions
__device__	__attribute__ ((always_inline)) void clusterTracksByDensity(gpuVertexFinder

	__syncthreads ()

	assert (soa)

	assert (nvFinal<=nvIntermediate)

	assert (pdata)

	assert (zt)

	assert ((int) nt<=hist.capacity())

	assert (hist.size()==nt)

	assert (foundClusters< ZVertices::MAXVTX)

size d for d tracks hist hist	capacity ()

hist	finalize (hws)

	for (int idx=first, nt=tracks.nTracks();idx< nt;idx+=gridDim.x *blockDim.x)

	for (auto j=threadIdx.x;j< Hist::totbins();j+=blockDim.x) = float(nn[i]) / chi2[i]

	if (verbose &&0==threadIdx.x) printf("params %d %f %f %f\

	if (1==nvFinal)

	if (threadIdx.x< 32) hws[threadIdx.x]=0

pws	init ()

size d for d tracks hist	nbins ()

	while (__syncthreads_or(more))

Variables
float *__restrict__	chi2 = data.chi2

__device__ WorkSpace float	chi2Max

WorkSpace int float float float	chi2max

auto &__restrict__	data = *pdata

WorkSpace int float	eps

auto	er2mx = errmax * errmax

WorkSpace int float float	errmax

float const *__restrict__	ezt2 = ws.ezt2

auto	first = blockIdx.x * blockDim.x + threadIdx.x

auto const &	fit = tracks.stateAtBS

__shared__ unsigned int	foundClusters = 0

__shared__ Hist	hist

__shared__ Hist::Counter	hws [32]

int32_t *__restrict__	iv = ws.iv

uint8_t *__restrict__	izt = ws.izt

__device__ WorkSpace float	maxChi2

constexpr float	maxChi2ForFinalFit = 5000.f

constexpr float	maxChi2ForFirstFit = 50.f

constexpr float	maxChi2ForSplit = 9.f

WorkSpace int	minT

bool	more = true

size d for d tracks	n

__shared__ int	nloops = 0

int32_t *__restrict__	nn = data.ndof

__shared__ int	noise

auto	nt = ws.ntrks

uint32_t &	nvFinal = data.nvFinal

uint32_t &	nvIntermediate = ws.nvIntermediate

ZVertexSoA WorkSpace float float	ptMax

ZVertexSoA WorkSpace float	ptMin

float const *__restrict__	ptt2 = ws.ptt2

float *__restrict__	ptv2 = data.ptv2

WorkSpace *	pws

auto const *	quality = tracks.qualityData()

ZVertexSoA *	soa

uint16_t *__restrict__	sortInd = data.sortInd

auto const &	tracks = *ptracks

auto &__restrict__	ws = *pws

float *__restrict__	wv = data.wv

float const *__restrict__	zt = ws.zt

float *__restrict__	zv = data.zv

Typedef Documentation

◆ Hist

typedef cms::cuda::HistoContainer< uint8_t, 256, 16000, 8, uint16_t > gpuVertexFinder::Hist

Definition at line 47 of file gpuClusterTracksDBSCAN.h.

◆ TkSoA

using gpuVertexFinder::TkSoA = typedef pixelTrack::TrackSoA

Definition at line 13 of file gpuVertexFinder.h.

◆ ZVertices

using gpuVertexFinder::ZVertices = typedef ZVertexSoA

Definition at line 12 of file gpuVertexFinder.h.

Function Documentation

◆ attribute()

gpuVertexFinder::__attribute__ ( (always_inline) )

inline

Definition at line 20 of file gpuClusterTracksByDensity.h.

References __syncthreads(), funct::abs(), assert(), cms::cudacompat::atomicInc(), cms::cudacompat::blockDim, cms::cuda::OneToManyAssoc::capacity(), chi2max, data, eps, er2mx, errmax, ezt2, foundClusters, hist, hws, mps_fire::i, createfilelist::int, iv, izt, dqmiolumiharvest::j, heppy_loop::loop, visualization-live-secondInstance_cfg::m, SiStripPI::max, ZVertexSoA::MAXTRACKS, ZVertexSoA::MAXVTX, SiStripPI::min, minT, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins(), nn, nt, nvFinal, nvIntermediate, cms::cuda::OneToManyAssoc::off, pws, cms::cudacompat::threadIdx, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::totbins(), verbose, ws, and zt.

     {
     using namespace gpuVertexFinder;
     constexpr bool verbose = false;  // in principle the compiler should optmize out if false
 
     if (verbose && 0 == threadIdx.x)
       printf("params %d %f %f %f\n", minT, eps, errmax, chi2max);
 
     auto er2mx = errmax * errmax;
 
     auto& __restrict__ data = *pdata;
     auto& __restrict__ ws = *pws;
     auto nt = ws.ntrks;
     float const* __restrict__ zt = ws.zt;
     float const* __restrict__ ezt2 = ws.ezt2;
 
     uint32_t& nvFinal = data.nvFinal;
     uint32_t& nvIntermediate = ws.nvIntermediate;
 
     uint8_t* __restrict__ izt = ws.izt;
     int32_t* __restrict__ nn = data.ndof;
     int32_t* __restrict__ iv = ws.iv;
 
     assert(pdata);
     assert(zt);
 
     using Hist = cms::cuda::HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
     __shared__ Hist hist;
     __shared__ typename Hist::Counter hws[32];
     for (auto j = threadIdx.x; j < Hist::totbins(); j += blockDim.x) {
       hist.off[j] = 0;
     }
     __syncthreads();
 
     if (verbose && 0 == threadIdx.x)
       printf("booked hist with %d bins, size %d for %d tracks\n", hist.nbins(), hist.capacity(), nt);
 
     assert((int)nt <= hist.capacity());
 
     // fill hist  (bin shall be wider than "eps")
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       assert(i < ZVertices::MAXTRACKS);
       int iz = int(zt[i] * 10.);  // valid if eps<=0.1
       // iz = std::clamp(iz, INT8_MIN, INT8_MAX);  // sorry c++17 only
       iz = std::min(std::max(iz, INT8_MIN), INT8_MAX);
       izt[i] = iz - INT8_MIN;
       assert(iz - INT8_MIN >= 0);
       assert(iz - INT8_MIN < 256);
       hist.count(izt[i]);
       iv[i] = i;
       nn[i] = 0;
     }
     __syncthreads();
     if (threadIdx.x < 32)
       hws[threadIdx.x] = 0;  // used by prefix scan...
     __syncthreads();
     hist.finalize(hws);
     __syncthreads();
     assert(hist.size() == nt);
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       hist.fill(izt[i], uint16_t(i));
     }
     __syncthreads();
 
     // count neighbours
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       if (ezt2[i] > er2mx)
         continue;
       auto loop = [&](uint32_t j) {
         if (i == j)
           return;
         auto dist = std::abs(zt[i] - zt[j]);
         if (dist > eps)
           return;
         if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
           return;
         nn[i]++;
       };
 
       cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __syncthreads();
 
     // find closest above me .... (we ignore the possibility of two j at same distance from i)
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       float mdist = eps;
       auto loop = [&](uint32_t j) {
         if (nn[j] < nn[i])
           return;
         if (nn[j] == nn[i] && zt[j] >= zt[i])
           return;  // if equal use natural order...
         auto dist = std::abs(zt[i] - zt[j]);
         if (dist > mdist)
           return;
         if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
           return;  // (break natural order???)
         mdist = dist;
         iv[i] = j;  // assign to cluster (better be unique??)
       };
       cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __syncthreads();
 
 #ifdef GPU_DEBUG
     //  mini verification
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       if (iv[i] != int(i))
         assert(iv[iv[i]] != int(i));
     }
     __syncthreads();
 #endif
 
     // consolidate graph (percolate index of seed)
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       auto m = iv[i];
       while (m != iv[m])
         m = iv[m];
       iv[i] = m;
     }
 
 #ifdef GPU_DEBUG
     __syncthreads();
     //  mini verification
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       if (iv[i] != int(i))
         assert(iv[iv[i]] != int(i));
     }
 #endif
 
 #ifdef GPU_DEBUG
     // and verify that we did not spit any cluster...
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       auto minJ = i;
       auto mdist = eps;
       auto loop = [&](uint32_t j) {
         if (nn[j] < nn[i])
           return;
         if (nn[j] == nn[i] && zt[j] >= zt[i])
           return;  // if equal use natural order...
         auto dist = std::abs(zt[i] - zt[j]);
         if (dist > mdist)
           return;
         if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
           return;
         mdist = dist;
         minJ = j;
       };
       cms::cuda::forEachInBins(hist, izt[i], 1, loop);
       // should belong to the same cluster...
       assert(iv[i] == iv[minJ]);
       assert(nn[i] <= nn[iv[i]]);
     }
     __syncthreads();
 #endif
 
     __shared__ unsigned int foundClusters;
     foundClusters = 0;
     __syncthreads();
 
     // find the number of different clusters, identified by a tracks with clus[i] == i and density larger than threshold;
     // mark these tracks with a negative id.
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       if (iv[i] == int(i)) {
         if (nn[i] >= minT) {
           auto old = atomicInc(&foundClusters, 0xffffffff);
           iv[i] = -(old + 1);
         } else {  // noise
           iv[i] = -9998;
         }
       }
     }
     __syncthreads();
 
     assert(foundClusters < ZVertices::MAXVTX);
 
     // propagate the negative id to all the tracks in the cluster.
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       if (iv[i] >= 0) {
         // mark each track in a cluster with the same id as the first one
         iv[i] = iv[iv[i]];
       }
     }
     __syncthreads();
 
     // adjust the cluster id to be a positive value starting from 0
     for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
       iv[i] = -iv[i] - 1;
     }
 
     nvIntermediate = nvFinal = foundClusters;
 
     if (verbose && 0 == threadIdx.x)
       printf("found %d proto vertices\n", foundClusters);
   }

◆ __syncthreads()

gpuVertexFinder::__syncthreads ( )

inline

Definition at line 132 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), __attribute__(), cms::cuda::__attribute__(), gpuClustering::for(), for(), and CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::merge_cands_for_seed().

132 {}

◆ assert() [1/7]

gpuVertexFinder::assert ( soa )

◆ assert() [2/7]

gpuVertexFinder::assert ( nvFinal<= nvIntermediate )

◆ assert() [3/7]

gpuVertexFinder::assert ( pdata )

Referenced by __attribute__(), for(), gpuVertexFinder::Producer::make(), and while().

◆ assert() [4/7]

gpuVertexFinder::assert ( zt )

◆ assert() [5/7]

gpuVertexFinder::assert ( (int) nt<=hist.capacity() )

◆ assert() [6/7]

gpuVertexFinder::assert ( hist. size() = =nt )

◆ assert() [7/7]

gpuVertexFinder::assert ( )

◆ capacity()

size d for d tracks hist hist gpuVertexFinder::capacity ( )

Referenced by SiPixelClusterShapeCache::insert(), cms::cuda::make_SimpleVector(), btagbtvdeep::DeepBoostedJetFeatures::reserve(), and FQueue< T >::reserve().

◆ finalize()

hist gpuVertexFinder::finalize ( hws )

Referenced by reco::TaggingVariableList::insert().

◆ for() [1/2]

gpuVertexFinder::for ( int idx = first )

Definition at line 29 of file gpuVertexFinder.cc.

References assert(), cms::cudacompat::atomicAdd(), data, pixelTrack::highPurity, ZVertexSoA::idv, heavyIonCSV_trainingSettings::idx, SiStripPI::min, nHits, DiDispStaMuonMonitor_cfi::pt, ptMax, ptMin, pws, quality, soa, and tracks.

                                                                                           {
       auto nHits = tracks.nHits(idx);
       assert(nHits >= 3);
 
       // initialize soa...
       soa->idv[idx] = -1;
 
       if (tracks.isTriplet(idx))
         continue;  // no triplets
       if (quality[idx] < pixelTrack::Quality::highPurity)
         continue;
 
       auto pt = tracks.pt(idx);
 
       if (pt < ptMin)
         continue;
 
       // clamp pt
       pt = std::min(pt, ptMax);
 
       auto& data = *pws;
       auto it = atomicAdd(&data.ntrks, 1);
       data.itrk[it] = idx;
       data.zt[it] = tracks.zip(idx);
       data.ezt2[it] = fit.covariance(idx)(14);
       data.ptt2[it] = pt * pt;
     }

◆ for() [2/2]

gpuVertexFinder::for ( ) = float(nn[i]) / chi2[i]

Definition at line 50 of file gpuClusterTracksDBSCAN.h.

References hist, dqmiolumiharvest::j, and cms::cuda::OneToManyAssoc::off.

                                                                    {
       hist.off[j] = 0;
     }

◆ if() [1/3]

gpuVertexFinder::if ( verbose && 0 = = threadIdx.x )

◆ if() [2/3]

gpuVertexFinder::if ( 1 = = nvFinal )

Definition at line 53 of file gpuSortByPt2.h.

References sortInd, and cms::cudacompat::threadIdx.

                       {
       if (threadIdx.x == 0)
         sortInd[0] = 0;
       return;
     }

◆ if() [3/3]

gpuVertexFinder::if ( )

pure virtual

◆ init()

pws gpuVertexFinder::init ( )

Referenced by gpuVertexFinder::Producer::make().

◆ nbins()

size d for d tracks hist gpuVertexFinder::nbins ( )

◆ while()

gpuVertexFinder::while ( __syncthreads_or(more) )

Definition at line 109 of file gpuClusterTracksIterative.h.

References funct::abs(), assert(), cms::cudacompat::atomicMin(), cms::cuda::be, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::bin(), cms::cudacompat::blockDim, chi2max, eps, ezt2, hist, mps_fire::i, iv, izt, dqmiolumiharvest::j, dqmdumpme::k, heppy_loop::loop, visualization-live-secondInstance_cfg::m, SiStripPI::min, minT, more, cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins(), nloops, nn, nt, AlCaHLTBitMon_ParallelJobs::p, cms::cudacompat::threadIdx, and zt.

Referenced by ThreeThresholdAlgorithm::addToCandidate(), EcalSimParametersFromDD::build(), cppFunctionSkipper::delLines(), cppFunctionSkipper::filterFile(), commentSkipper.buildFileCommentSkipper::filterMultilineComment(), mkfit::MkBuilder::find_tracks_in_layers(), mkfit::MkBuilder::findTracksStandard(), dqm_interfaces.DQMcommunicator::get_runs_list(), and histoStyle::savePlots().

                                    {
       if (1 == nloops % 2) {
         for (auto i = threadIdx.x; i < nt; i += blockDim.x) {
           auto m = iv[i];
           while (m != iv[m])
             m = iv[m];
           iv[i] = m;
         }
       } else {
         more = false;
         for (auto k = threadIdx.x; k < hist.size(); k += blockDim.x) {
           auto p = hist.begin() + k;
           auto i = (*p);
           auto be = std::min(Hist::bin(izt[i]) + 1, int(hist.nbins() - 1));
           if (nn[i] < minT)
             continue;  // DBSCAN core rule
           auto loop = [&](uint32_t j) {
             assert(i != j);
             if (nn[j] < minT)
               return;  // DBSCAN core rule
             auto dist = std::abs(zt[i] - zt[j]);
             if (dist > eps)
               return;
             if (dist * dist > chi2max * (ezt2[i] + ezt2[j]))
               return;
             auto old = atomicMin(&iv[j], iv[i]);
             if (old != iv[i]) {
               // end the loop only if no changes were applied
               more = true;
             }
             atomicMin(&iv[i], old);
           };
           ++p;
           for (; p < hist.end(be); ++p)
             loop(*p);
         }  // for i
       }
       if (threadIdx.x == 0)
         ++nloops;
     }  // while