gpuVertexFinder::Producer< TrackerTraits > Class Template Reference

#include <gpuVertexFinder.h>

Public Member Functions
ZVertexSoAHost	make (const TkSoAConstView &tracks_view, float ptMin, float ptMax) const
ZVertexSoADevice	makeAsync (cudaStream_t stream, const TkSoAConstView &tracks_view, float ptMin, float ptMax) const
	Producer (bool oneKernel, bool useDensity, bool useDBSCAN, bool useIterative, int iminT, float ieps, float ierrmax, float ichi2max)
	~Producer ()=default

Private Types
using	TkSoAConstView = TrackSoAConstView< TrackerTraits >

Private Attributes
float	chi2max
float	eps
float	errmax
int	minT
const bool	oneKernel_
const bool	useDBSCAN_
const bool	useDensity_
const bool	useIterative_

Detailed Description

template<typename TrackerTraits>
class gpuVertexFinder::Producer< TrackerTraits >

Definition at line 26 of file gpuVertexFinder.h.

Member Typedef Documentation

TkSoAConstView

template<typename TrackerTraits >

using gpuVertexFinder::Producer< TrackerTraits >::TkSoAConstView = TrackSoAConstView<TrackerTraits>

private

Definition at line 27 of file gpuVertexFinder.h.

Constructor & Destructor Documentation

Producer()

template<typename TrackerTraits >

gpuVertexFinder::Producer< TrackerTraits >::Producer ( bool  oneKernel, bool  useDensity, bool  useDBSCAN, bool  useIterative, int  iminT, float  ieps, float  ierrmax, float  ichi2max  )

inline

Definition at line 30 of file gpuVertexFinder.h.

        : oneKernel_(oneKernel && !(useDBSCAN || useIterative)),
          useDensity_(useDensity),
          useDBSCAN_(useDBSCAN),
          useIterative_(useIterative),
          minT(iminT),
          eps(ieps),
          errmax(ierrmax),
          chi2max(ichi2max) {}

~Producer()

template<typename TrackerTraits >

gpuVertexFinder::Producer< TrackerTraits >::~Producer ( )

default

Member Function Documentation

make()

template<typename TrackerTraits >

ZVertexSoAHost gpuVertexFinder::Producer< TrackerTraits >::make	(	const TkSoAConstView &	tracks_view,
		float	ptMin,
		float	ptMax
	)		const

Definition at line 113 of file gpuVertexFinder.cc.

References gpuVertexFinder::assert(), gpuVertexFinder::chi2max, gather_cfg::cout, cudaCheck, gpuVertexFinder::eps, gpuVertexFinder::errmax, gpuVertexFinder::fitVertices(), init, gpuVertexFinder::maxChi2ForFinalFit, gpuVertexFinder::maxChi2ForFirstFit, gpuVertexFinder::maxChi2ForSplit, gpuVertexFinder::minT, gpuVertexFinder::ptMax, gpuVertexFinder::ptMin, gpuVertexFinder::soa, gpuVertexFinder::sortByPt2(), gpuVertexFinder::splitVertices(), cms::cuda::stream, caHitNtupletGeneratorKernels::tracks_view, and AlignmentTracksFromVertexSelector_cfi::vertices.

                                                                  {
#ifdef PIXVERTEX_DEBUG_PRODUCE
    std::cout << "producing Vertices on  CPU" << std::endl;
#endif  // PIXVERTEX_DEBUG_PRODUCE
    ZVertexSoAHost vertices;
#endif
    auto soa = vertices.view();

    assert(vertices.buffer());

#ifdef __CUDACC__
    auto ws_d = gpuVertexFinder::workSpace::PixelVertexWorkSpaceSoADevice(stream);
#else
    auto ws_d = gpuVertexFinder::workSpace::PixelVertexWorkSpaceSoAHost();
#endif

#ifdef __CUDACC__
    init<<<1, 1, 0, stream>>>(soa, ws_d.view());
    auto blockSize = 128;
    auto numberOfBlocks = (tracks_view.metadata().size() + blockSize - 1) / blockSize;
    loadTracks<TrackerTraits><<<numberOfBlocks, blockSize, 0, stream>>>(tracks_view, soa, ws_d.view(), ptMin, ptMax);
    cudaCheck(cudaGetLastError());
#else
    init(soa, ws_d.view());
    loadTracks<TrackerTraits>(tracks_view, soa, ws_d.view(), ptMin, ptMax);
#endif

#ifdef __CUDACC__
    // Running too many thread lead to problems when printf is enabled.
    constexpr int maxThreadsForPrint = 1024 - 128;
    constexpr int numBlocks = 1024;
    constexpr int threadsPerBlock = 128;

    if (oneKernel_) {
      // implemented only for density clustesrs
#ifndef THREE_KERNELS
      vertexFinderOneKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
#else
      vertexFinderKernel1<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
      cudaCheck(cudaGetLastError());
      // one block per vertex...
      splitVerticesKernel<<<numBlocks, threadsPerBlock, 0, stream>>>(soa, ws_d.view(), maxChi2ForSplit);
      cudaCheck(cudaGetLastError());
      vertexFinderKernel2<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view());
#endif
    } else {  // five kernels
      if (useDensity_) {
        clusterTracksByDensityKernel<<<1, maxThreadsForPrint, 0, stream>>>(
            soa, ws_d.view(), minT, eps, errmax, chi2max);
      } else if (useDBSCAN_) {
        clusterTracksDBSCAN<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
      } else if (useIterative_) {
        clusterTracksIterative<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), minT, eps, errmax, chi2max);
      }
      cudaCheck(cudaGetLastError());
      fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFirstFit);
      cudaCheck(cudaGetLastError());
      // one block per vertex...
      splitVerticesKernel<<<numBlocks, threadsPerBlock, 0, stream>>>(soa, ws_d.view(), maxChi2ForSplit);
      cudaCheck(cudaGetLastError());
      fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFinalFit);
      cudaCheck(cudaGetLastError());
      sortByPt2Kernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view());
    }
    cudaCheck(cudaGetLastError());
#else  // __CUDACC__
    if (useDensity_) {
      clusterTracksByDensity(soa, ws_d.view(), minT, eps, errmax, chi2max);
    } else if (useDBSCAN_) {
      clusterTracksDBSCAN(soa, ws_d.view(), minT, eps, errmax, chi2max);
    } else if (useIterative_) {
      clusterTracksIterative(soa, ws_d.view(), minT, eps, errmax, chi2max);
    }
#ifdef PIXVERTEX_DEBUG_PRODUCE
    std::cout << "found " << ws_d.view().nvIntermediate() << " vertices " << std::endl;
#endif  // PIXVERTEX_DEBUG_PRODUCE
    fitVertices(soa, ws_d.view(), maxChi2ForFirstFit);
    // one block per vertex!
    splitVertices(soa, ws_d.view(), maxChi2ForSplit);
    fitVertices(soa, ws_d.view(), maxChi2ForFinalFit);
    sortByPt2(soa, ws_d.view());
#endif

    return vertices;
  }

makeAsync()

template<typename TrackerTraits >

ZVertexSoADevice gpuVertexFinder::Producer< TrackerTraits >::makeAsync	(	cudaStream_t	stream,
		const TkSoAConstView &	tracks_view,
		float	ptMin,
		float	ptMax
	)		const

Member Data Documentation

chi2max

template<typename TrackerTraits >

float gpuVertexFinder::Producer< TrackerTraits >::chi2max

private

Definition at line 62 of file gpuVertexFinder.h.

eps

template<typename TrackerTraits >

float gpuVertexFinder::Producer< TrackerTraits >::eps

private

Definition at line 60 of file gpuVertexFinder.h.

errmax

template<typename TrackerTraits >

float gpuVertexFinder::Producer< TrackerTraits >::errmax

private

Definition at line 61 of file gpuVertexFinder.h.

minT

template<typename TrackerTraits >

int gpuVertexFinder::Producer< TrackerTraits >::minT

private

Definition at line 59 of file gpuVertexFinder.h.

oneKernel_

template<typename TrackerTraits >

const bool gpuVertexFinder::Producer< TrackerTraits >::oneKernel_

private

Definition at line 54 of file gpuVertexFinder.h.

useDBSCAN_

template<typename TrackerTraits >

const bool gpuVertexFinder::Producer< TrackerTraits >::useDBSCAN_

private

Definition at line 56 of file gpuVertexFinder.h.

useDensity_

template<typename TrackerTraits >

const bool gpuVertexFinder::Producer< TrackerTraits >::useDensity_

private

Definition at line 55 of file gpuVertexFinder.h.

useIterative_

template<typename TrackerTraits >

const bool gpuVertexFinder::Producer< TrackerTraits >::useIterative_

private

Definition at line 57 of file gpuVertexFinder.h.