gpuVertexFinder::Producer Class Reference

#include <gpuVertexFinder.h>

Public Types
using	TkSoA = pixelTrack::TrackSoA
using	WorkSpace = gpuVertexFinder::WorkSpace
using	ZVertices = ZVertexSoA

Public Member Functions
ZVertexHeterogeneous	make (TkSoA const *tksoa, float ptMin) const
ZVertexHeterogeneous	makeAsync (cudaStream_t stream, TkSoA const *tksoa, float ptMin) const
	Producer (bool oneKernel, bool useDensity, bool useDBSCAN, bool useIterative, int iminT, float ieps, float ierrmax, float ichi2max)
	~Producer ()=default

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

Detailed Description

Definition at line 40 of file gpuVertexFinder.h.

Member Typedef Documentation

using gpuVertexFinder::Producer::TkSoA = pixelTrack::TrackSoA

Definition at line 44 of file gpuVertexFinder.h.

using gpuVertexFinder::Producer::WorkSpace = gpuVertexFinder::WorkSpace

Definition at line 43 of file gpuVertexFinder.h.

using gpuVertexFinder::Producer::ZVertices = ZVertexSoA

Definition at line 42 of file gpuVertexFinder.h.

Constructor & Destructor Documentation

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

inline

Definition at line 46 of file gpuVertexFinder.h.

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

gpuVertexFinder::Producer::~Producer ( )

default

Member Function Documentation

ZVertexHeterogeneous gpuVertexFinder::Producer::make	(	TkSoA const *	tksoa,
		float	ptMin
	)		const

Definition at line 102 of file gpuVertexFinder.cc.

References gpuVertexFinder::assert(), chi2max, gather_cfg::cout, cudaCheck, eps, errmax, HeterogeneousSoA< T >::get(), gpuVertexFinder::init(), gpuVertexFinder::maxChi2ForFinalFit, gpuVertexFinder::maxChi2ForFirstFit, gpuVertexFinder::maxChi2ForSplit, minT, oneKernel_, gpuVertexFinder::ptMin, gpuVertexFinder::soa, cms::cuda::stream, TrackSoAHeterogeneousT< S >::stride(), useDBSCAN_, useDensity_, useIterative_, and beam_dqm_sourceclient-live_cfg::vertices.

Referenced by PixelVertexProducerCUDA::produceOnCPU().

                                                                           {
#ifdef PIXVERTEX_DEBUG_PRODUCE
    std::cout << "producing Vertices on  CPU" << std::endl;
#endif  // PIXVERTEX_DEBUG_PRODUCE
    ZVertexHeterogeneous vertices(std::make_unique<ZVertexSoA>());
#endif
    assert(tksoa);
    auto* soa = vertices.get();
    assert(soa);

#ifdef __CUDACC__
    auto ws_d = cms::cuda::make_device_unique<WorkSpace>(stream);
#else
    auto ws_d = std::make_unique<WorkSpace>();
#endif

#ifdef __CUDACC__
    init<<<1, 1, 0, stream>>>(soa, ws_d.get());
    auto blockSize = 128;
    auto numberOfBlocks = (TkSoA::stride() + blockSize - 1) / blockSize;
    loadTracks<<<numberOfBlocks, blockSize, 0, stream>>>(tksoa, soa, ws_d.get(), ptMin);
    cudaCheck(cudaGetLastError());
#else
    init(soa, ws_d.get());
    loadTracks(tksoa, soa, ws_d.get(), ptMin);
#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.get(), minT, eps, errmax, chi2max);
#else
      vertexFinderKernel1<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get(), minT, eps, errmax, chi2max);
      cudaCheck(cudaGetLastError());
      // one block per vertex...
      splitVerticesKernel<<<numBlocks, threadsPerBlock, 0, stream>>>(soa, ws_d.get(), maxChi2ForSplit);
      cudaCheck(cudaGetLastError());
      vertexFinderKernel2<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get());
#endif
    } else {  // five kernels
      if (useDensity_) {
        clusterTracksByDensityKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get(), minT, eps, errmax, chi2max);
      } else if (useDBSCAN_) {
        clusterTracksDBSCAN<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get(), minT, eps, errmax, chi2max);
      } else if (useIterative_) {
        clusterTracksIterative<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get(), minT, eps, errmax, chi2max);
      }
      cudaCheck(cudaGetLastError());
      fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get(), maxChi2ForFirstFit);
      cudaCheck(cudaGetLastError());
      // one block per vertex...
      splitVerticesKernel<<<numBlocks, threadsPerBlock, 0, stream>>>(soa, ws_d.get(), maxChi2ForSplit);
      cudaCheck(cudaGetLastError());
      fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get(), maxChi2ForFinalFit);
      cudaCheck(cudaGetLastError());
      sortByPt2Kernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.get());
    }
    cudaCheck(cudaGetLastError());
#else  // __CUDACC__
    if (useDensity_) {
      clusterTracksByDensity(soa, ws_d.get(), minT, eps, errmax, chi2max);
    } else if (useDBSCAN_) {
      clusterTracksDBSCAN(soa, ws_d.get(), minT, eps, errmax, chi2max);
    } else if (useIterative_) {
      clusterTracksIterative(soa, ws_d.get(), minT, eps, errmax, chi2max);
    }
#ifdef PIXVERTEX_DEBUG_PRODUCE
    std::cout << "found " << (*ws_d).nvIntermediate << " vertices " << std::endl;
#endif  // PIXVERTEX_DEBUG_PRODUCE
    fitVertices(soa, ws_d.get(), maxChi2ForFirstFit);
    // one block per vertex!
    splitVertices(soa, ws_d.get(), maxChi2ForSplit);
    fitVertices(soa, ws_d.get(), maxChi2ForFinalFit);
    sortByPt2(soa, ws_d.get());
#endif

    return vertices;
  }

ZVertexHeterogeneous gpuVertexFinder::Producer::makeAsync	(	cudaStream_t	stream,
		TkSoA const *	tksoa,
		float	ptMin
	)		const

Referenced by PixelVertexProducerCUDA::produceOnGPU().

Member Data Documentation

float gpuVertexFinder::Producer::chi2max

private

Definition at line 78 of file gpuVertexFinder.h.

Referenced by make().

float gpuVertexFinder::Producer::eps

private

Definition at line 76 of file gpuVertexFinder.h.

Referenced by make().

float gpuVertexFinder::Producer::errmax

private

Definition at line 77 of file gpuVertexFinder.h.

Referenced by make().

int gpuVertexFinder::Producer::minT

private

Definition at line 75 of file gpuVertexFinder.h.

Referenced by make().

const bool gpuVertexFinder::Producer::oneKernel_

private

Definition at line 70 of file gpuVertexFinder.h.

Referenced by make().

const bool gpuVertexFinder::Producer::useDBSCAN_

private

Definition at line 72 of file gpuVertexFinder.h.

Referenced by make().

const bool gpuVertexFinder::Producer::useDensity_

private

Definition at line 71 of file gpuVertexFinder.h.

Referenced by make().

const bool gpuVertexFinder::Producer::useIterative_

private

Definition at line 73 of file gpuVertexFinder.h.

Referenced by make().