Namespaces
	workSpace

Classes
class	Producer

Typedefs
using	helper = TracksUtilities< TrackerTraits >

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

using	VtxSoAView = zVertex::ZVertexSoAView

using	WsSoAView = gpuVertexFinder::workSpace::PixelVertexWorkSpaceSoAView

Functions
__device__	__attribute__ ((always_inline)) void clusterTracksByDensity(VtxSoAView &pdata

	__attribute__ ((always_inline)) void clusterTracksByDensityKernel(VtxSoAView pdata

	__syncthreads ()

	assert (ptv2)

	assert (sortInd)

	assert (wv)

	assert (chi2)

	assert (nvFinal<=nvIntermediate)

	assert (zt)

	assert (ezt2)

	assert (izt)

	assert (nn)

	assert (iv)

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

	assert (hist.size()==nt)

	assert (foundClusters< zVertex::utilities::MAXVTX)

size d for d tracks hist hist	capacity ()

hist	finalize (hws)

	fitVertices (pdata, pws, maxChi2ForFirstFit)

	fitVertices (pdata, pws, maxChi2ForFinalFit)

	for (int idx=first, nt=tracks_view.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

size d for d tracks hist	nbins ()

	sortByPt2 (pdata, pws)

	splitVertices (pdata, pws, maxChi2ForSplit)

	while (__syncthreads_or(more))

Variables
float *__restrict__	chi2 = data.chi2()

__device__ WsSoAView float	chi2Max

__device__ WsSoAView int float float float	chi2max

auto &__restrict__	data = pdata

__device__ WsSoAView int float	eps

auto	er2mx = errmax * errmax

__device__ WsSoAView int float float	errmax

float const *__restrict__	ezt2 = ws.ezt2()

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

__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__ WsSoAView float	maxChi2

constexpr float	maxChi2ForFinalFit = 5000.f

constexpr float	maxChi2ForFirstFit = 50.f

constexpr float	maxChi2ForSplit = 9.f

__device__ WsSoAView 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()

VtxSoAView WsSoAView float float	ptMax

VtxSoAView WsSoAView float	ptMin

float const *__restrict__	ptt2 = ws.ptt2()

float *__restrict__	ptv2 = data.ptv2()

__device__ WsSoAView &	pws

VtxSoAView	soa

uint16_t *__restrict__	sortInd = data.sortInd()

constexpr bool	verbose = false

auto &__restrict__	ws = pws

float *__restrict__	wv = data.wv()

float const *__restrict__	zt = ws.zt()

float *__restrict__	zv = data.zv()

Typedef Documentation

◆ helper

using gpuVertexFinder::helper = typedef TracksUtilities<TrackerTraits>

Definition at line 32 of file gpuVertexFinder.cc.

◆ Hist

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

Definition at line 54 of file gpuClusterTracksByDensity.h.

◆ VtxSoAView

using gpuVertexFinder::VtxSoAView = typedef zVertex::ZVertexSoAView

Definition at line 17 of file gpuVertexFinder.h.

◆ WsSoAView

using gpuVertexFinder::WsSoAView = typedef gpuVertexFinder::workSpace::PixelVertexWorkSpaceSoAView

Definition at line 18 of file gpuVertexFinder.h.

Function Documentation

◆ attribute() [1/2]

__device__ gpuVertexFinder::__attribute__ ( (always_inline) ) &

inline

◆ attribute() [2/2]

gpuVertexFinder::__attribute__ ( (always_inline) )

inline

◆ __syncthreads()

gpuVertexFinder::__syncthreads ( )

inline

Definition at line 132 of file cudaCompat.h.

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

132 {}

◆ assert() [1/13]

gpuVertexFinder::assert ( ptv2 )

◆ assert() [2/13]

gpuVertexFinder::assert ( sortInd )

◆ assert() [3/13]

gpuVertexFinder::assert ( wv )

◆ assert() [4/13]

gpuVertexFinder::assert ( chi2 )

◆ assert() [5/13]

gpuVertexFinder::assert ( nvFinal<= nvIntermediate )

◆ assert() [6/13]

gpuVertexFinder::assert ( zt )

Referenced by for(), gpuVertexFinder::Producer< TrackerTraits >::make(), and while().

◆ assert() [7/13]

gpuVertexFinder::assert ( ezt2 )

◆ assert() [8/13]

gpuVertexFinder::assert ( izt )

◆ assert() [9/13]

gpuVertexFinder::assert ( nn )

◆ assert() [10/13]

gpuVertexFinder::assert ( iv )

◆ assert() [11/13]

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

◆ assert() [12/13]

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

◆ assert() [13/13]

gpuVertexFinder::assert ( )

◆ capacity()

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

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

◆ finalize()

hist gpuVertexFinder::finalize ( hws )

◆ fitVertices() [1/2]

gpuVertexFinder::fitVertices	(	pdata	,
		pws	,
		maxChi2ForFirstFit
	)

Referenced by gpuVertexFinder::Producer< TrackerTraits >::make(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::VertexFinderOneKernel::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::FitVerticesKernel::operator()().

◆ fitVertices() [2/2]

gpuVertexFinder::fitVertices	(	pdata	,
		pws	,
		maxChi2ForFinalFit
	)

◆ for() [1/2]

gpuVertexFinder::for ( int idx = first )

Definition at line 34 of file gpuVertexFinder.cc.

References assert(), cms::cudacompat::atomicAdd(), data, pixelTrack::highPurity, heavyIonCSV_trainingSettings::idx, TracksUtilities< TrackerTraits >::isTriplet(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::it, SiStripPI::min, nHits, TracksUtilities< TrackerTraits >::nHits(), DiDispStaMuonMonitor_cfi::pt, ptMax, ptMin, pws, quality, soa, caHitNtupletGeneratorKernels::tracks_view, and TracksUtilities< TrackerTraits >::zip().

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

◆ for() [2/2]

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

Definition at line 57 of file gpuClusterTracksByDensity.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

◆ nbins()

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

◆ sortByPt2()

gpuVertexFinder::sortByPt2	(	pdata	,
		pws
	)

Referenced by gpuVertexFinder::Producer< TrackerTraits >::make().

◆ splitVertices()

gpuVertexFinder::splitVertices	(	pdata	,
		pws	,
		maxChi2ForSplit
	)

Referenced by gpuVertexFinder::Producer< TrackerTraits >::make(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::VertexFinderOneKernel::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::SplitVerticesKernel::operator()().

◆ while()

gpuVertexFinder::while ( __syncthreads_or(more) )

Definition at line 111 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