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)

hist	finalize (hws)

	for (int idx=first, nt=TkSoA::stride();idx< nt;idx+=gridDim.x *blockDim.x)

	for (auto j=threadIdx.x;j< Hist::totbins();j+=blockDim.x)

	if (1==nvFinal)

pws	init ()

	printf ("params %d %f %f %f\n", minT, eps, errmax, chi2max)

	printf ("booked hist with %d bins, size %d for %d tracks\n", hist.nbins(), hist.capacity(), nt)

wv *[i]	printf ("found %d proto clusters ", foundClusters)

	printf ("and %d noise\n", noise)

	printf ("found %d proto vertices\n", foundClusters)

	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] = 0

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

__shared__ int	nloops = 0

int32_t *__restrict__	nn = data.ndof

__shared__ int	noise = 0

auto	nt = ws.ntrks

uint32_t &	nvFinal = data.nvFinal

uint32_t &	nvIntermediate = ws.nvIntermediate

ZVertexSoA WorkSpace float	ptMin

float const *__restrict__	ptt2 = ws.ptt2

float *__restrict__	ptv2 = data.ptv2

WorkSpace *	pws

auto const *	quality = tracks.qualityData()

	return

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

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

Definition at line 47 of file gpuClusterTracksDBSCAN.h.

using gpuVertexFinder::TkSoA = typedef pixelTrack::TrackSoA

Definition at line 12 of file gpuVertexFinder.h.

using gpuVertexFinder::ZVertices = typedef ZVertexSoA

Definition at line 11 of file gpuVertexFinder.h.

Function Documentation

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(), data, eps, er2mx, errmax, ezt2, foundClusters, hist, hws, mps_fire::i, iv, izt, dqmiolumiharvest::j, heppy_loop::loop, visualization-live-secondInstance_cfg::m, SiStripPI::max, ZVertexSoA::MAXTRACKS, ZVertexSoA::MAXVTX, min(), cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins(), nn, nt, nvFinal, nvIntermediate, cms::cuda::OneToManyAssoc::off, printf(), 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);
   }

gpuVertexFinder::__syncthreads ( )

inline

Definition at line 108 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), __attribute__(), cms::cuda::__attribute__(), gpuClustering::for(), and for().

108 {}

gpuVertexFinder::assert ( soa )

gpuVertexFinder::assert ( nvFinal<= nvIntermediate )

gpuVertexFinder::assert ( pdata )

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

gpuVertexFinder::assert ( zt )

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

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

gpuVertexFinder::assert ( )

hist gpuVertexFinder::finalize ( hws )

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

gpuVertexFinder::for ( int idx = first )

Definition at line 29 of file gpuVertexFinder.cc.

References cms::cudacompat::atomicAdd(), data, fit, pixelTrack::highPurity, ZVertexSoA::idv, nHits, DiDispStaMuonMonitor_cfi::pt, pws, quality, and tracks.

                                                                                        {
       auto nHits = tracks.nHits(idx);
       if (nHits == 0)
         break;  // this is a guard: maybe we need to move to nTracks...
 
       // 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;
 
       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;
     }

gpuVertexFinder::for ( )

Definition at line 50 of file gpuClusterTracksDBSCAN.h.

References dqmiolumiharvest::j.

                                                                    {
       hist.off[j] = 0;
     }

gpuVertexFinder::if ( 1 = = nvFinal )

Definition at line 53 of file gpuSortByPt2.h.

References cms::cudacompat::threadIdx.

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

pws gpuVertexFinder::init ( )

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

gpuVertexFinder::printf	(	"params %d %f %f %f\n"	,
		minT	,
		eps	,
		errmax	,
		chi2max
	)

gpuVertexFinder::printf	(	"booked hist with %d	bins,
		size%d for%d tracks\n"	,
		hist.	nbins(),
		hist.	capacity(),
		nt
	)

wv* [i] gpuVertexFinder::printf	(	"found %d proto clusters "	,
		foundClusters
	)

gpuVertexFinder::printf	(	"and %d noise\n"	,
		noise
	)

gpuVertexFinder::printf	(	"found %d proto vertices\n"	,
		foundClusters
	)

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, mps_fire::i, dqmiolumiharvest::j, isotrackApplyRegressor::k, heppy_loop::loop, visualization-live-secondInstance_cfg::m, min(), nloops, nt, AlCaHLTBitMon_ParallelJobs::p, and cms::cudacompat::threadIdx.

Referenced by ThreeThresholdAlgorithm::addToCandidate(), EcalSimParametersFromDD::build(), commentSkipper.buildFileCommentSkipper::filterMultilineComment(), 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