#include <CAHitNtupletGeneratorOnGPU.h>

Public Types
using	Counters = cAHitNtupletGenerator::Counters

using	hindex_type = TrackingRecHit2DSOAView::hindex_type

using	HitContainer = pixelTrack::HitContainer

using	HitsOnCPU = TrackingRecHit2DGPU

using	HitsOnGPU = TrackingRecHit2DSOAView

using	OutputSoA = pixelTrack::TrackSoA

using	Params = cAHitNtupletGenerator::Params

using	Quality = pixelTrack::Quality

using	QualityCuts = cAHitNtupletGenerator::QualityCuts

using	Tuple = HitContainer

Public Member Functions
	CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &&iC)

	CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &iC)

PixelTrackHeterogeneous	makeTuples (TrackingRecHit2DCPU const &hits_d, float bfield) const

PixelTrackHeterogeneous	makeTuplesAsync (TrackingRecHit2DGPU const &hits_d, float bfield, cudaStream_t stream) const

	~CAHitNtupletGeneratorOnGPU ()

Static Public Member Functions
static void	fillDescriptions (edm::ParameterSetDescription &desc)

static const char *	fillDescriptionsLabel ()

Private Member Functions
void	buildDoublets (HitsOnCPU const &hh, cudaStream_t stream) const

void	hitNtuplets (HitsOnCPU const &hh, const edm::EventSetup &es, bool useRiemannFit, cudaStream_t cudaStream)

void	launchKernels (HitsOnCPU const &hh, bool useRiemannFit, cudaStream_t cudaStream) const

Private Attributes
Counters *	m_counters = nullptr

Params	m_params

Detailed Description

Definition at line 24 of file CAHitNtupletGeneratorOnGPU.h.

Member Typedef Documentation

◆ Counters

using CAHitNtupletGeneratorOnGPU::Counters = cAHitNtupletGenerator::Counters

Definition at line 37 of file CAHitNtupletGeneratorOnGPU.h.

◆ hindex_type

using CAHitNtupletGeneratorOnGPU::hindex_type = TrackingRecHit2DSOAView::hindex_type

Definition at line 28 of file CAHitNtupletGeneratorOnGPU.h.

◆ HitContainer

using CAHitNtupletGeneratorOnGPU::HitContainer = pixelTrack::HitContainer

Definition at line 32 of file CAHitNtupletGeneratorOnGPU.h.

◆ HitsOnCPU

using CAHitNtupletGeneratorOnGPU::HitsOnCPU = TrackingRecHit2DGPU

Definition at line 27 of file CAHitNtupletGeneratorOnGPU.h.

◆ HitsOnGPU

using CAHitNtupletGeneratorOnGPU::HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 26 of file CAHitNtupletGeneratorOnGPU.h.

◆ OutputSoA

using CAHitNtupletGeneratorOnGPU::OutputSoA = pixelTrack::TrackSoA

Definition at line 31 of file CAHitNtupletGeneratorOnGPU.h.

◆ Params

using CAHitNtupletGeneratorOnGPU::Params = cAHitNtupletGenerator::Params

Definition at line 36 of file CAHitNtupletGeneratorOnGPU.h.

◆ Quality

using CAHitNtupletGeneratorOnGPU::Quality = pixelTrack::Quality

Definition at line 30 of file CAHitNtupletGeneratorOnGPU.h.

◆ QualityCuts

using CAHitNtupletGeneratorOnGPU::QualityCuts = cAHitNtupletGenerator::QualityCuts

Definition at line 35 of file CAHitNtupletGeneratorOnGPU.h.

◆ Tuple

using CAHitNtupletGeneratorOnGPU::Tuple = HitContainer

Definition at line 33 of file CAHitNtupletGeneratorOnGPU.h.

Constructor & Destructor Documentation

◆ CAHitNtupletGeneratorOnGPU() [1/2]

CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU	(	const edm::ParameterSet &	cfg,
		edm::ConsumesCollector &&	iC
	)

inline

Definition at line 40 of file CAHitNtupletGeneratorOnGPU.h.

41 : CAHitNtupletGeneratorOnGPU(cfg, iC) {}

◆ CAHitNtupletGeneratorOnGPU() [2/2]

CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU	(	const edm::ParameterSet &	cfg,
		edm::ConsumesCollector &	iC
	)

Definition at line 60 of file CAHitNtupletGeneratorOnGPU.cc.

     : m_params(cfg.getParameter<bool>("onGPU"),
                cfg.getParameter<unsigned int>("minHitsPerNtuplet"),
                cfg.getParameter<unsigned int>("maxNumberOfDoublets"),
                cfg.getParameter<unsigned int>("minHitsForSharingCut"),
                cfg.getParameter<bool>("useRiemannFit"),
                cfg.getParameter<bool>("fit5as4"),
                cfg.getParameter<bool>("includeJumpingForwardDoublets"),
                cfg.getParameter<bool>("earlyFishbone"),
                cfg.getParameter<bool>("lateFishbone"),
                cfg.getParameter<bool>("idealConditions"),
                cfg.getParameter<bool>("fillStatistics"),
                cfg.getParameter<bool>("doClusterCut"),
                cfg.getParameter<bool>("doZ0Cut"),
                cfg.getParameter<bool>("doPtCut"),
                cfg.getParameter<bool>("doSharedHitCut"),
                cfg.getParameter<bool>("dupPassThrough"),
                cfg.getParameter<bool>("useSimpleTripletCleaner"),
                cfg.getParameter<double>("ptmin"),
                cfg.getParameter<double>("CAThetaCutBarrel"),
                cfg.getParameter<double>("CAThetaCutForward"),
                cfg.getParameter<double>("hardCurvCut"),
                cfg.getParameter<double>("dcaCutInnerTriplet"),
                cfg.getParameter<double>("dcaCutOuterTriplet"),
                makeQualityCuts(cfg.getParameterSet("trackQualityCuts"))) {
 #ifdef DUMP_GPU_TK_TUPLES
   printf("TK: %s %s % %s %s %s %s %s %s %s %s %s %s %s %s %s\n",
          "tid",
          "qual",
          "nh",
          "charge",
          "pt",
          "eta",
          "phi",
          "tip",
          "zip",
          "chi2",
          "h1",
          "h2",
          "h3",
          "h4",
          "h5");
 #endif
  
   if (m_params.onGPU_) {
     // allocate pinned host memory only if CUDA is available
     edm::Service<CUDAService> cs;
     if (cs and cs->enabled()) {
       cudaCheck(cudaMalloc(&m_counters, sizeof(Counters)));
       cudaCheck(cudaMemset(m_counters, 0, sizeof(Counters)));
     }
   } else {
     m_counters = new Counters();
     memset(m_counters, 0, sizeof(Counters));
   }
 }

References fwrapper::cs, cudaCheck, m_counters, m_params, and cAHitNtupletGenerator::Params::onGPU_.

◆ ~CAHitNtupletGeneratorOnGPU()

CAHitNtupletGeneratorOnGPU::~CAHitNtupletGeneratorOnGPU ( )

Definition at line 117 of file CAHitNtupletGeneratorOnGPU.cc.

                                                         {
   if (m_params.onGPU_) {
     // print the gpu statistics and free pinned host memory only if CUDA is available
     edm::Service<CUDAService> cs;
     if (cs and cs->enabled()) {
       if (m_params.doStats_) {
         // crash on multi-gpu processes
         CAHitNtupletGeneratorKernelsGPU::printCounters(m_counters);
       }
       cudaFree(m_counters);
     }
   } else {
     if (m_params.doStats_) {
       CAHitNtupletGeneratorKernelsCPU::printCounters(m_counters);
     }
     delete m_counters;
   }
 }

References fwrapper::cs, cAHitNtupletGenerator::Params::doStats_, m_counters, m_params, cAHitNtupletGenerator::Params::onGPU_, and CAHitNtupletGeneratorKernels< TTraits >::printCounters().

Member Function Documentation

◆ buildDoublets()

void CAHitNtupletGeneratorOnGPU::buildDoublets	(	HitsOnCPU const &	hh,
		cudaStream_t	stream
	)		const

private

◆ fillDescriptions()

void CAHitNtupletGeneratorOnGPU::fillDescriptions ( edm::ParameterSetDescription & desc )

static

Definition at line 136 of file CAHitNtupletGeneratorOnGPU.cc.

                                                                                   {
   // 87 cm/GeV = 1/(3.8T * 0.3)
   // take less than radius given by the hardPtCut and reject everything below
   // auto hardCurvCut = 1.f/(0.35 * 87.f);
   desc.add<double>("ptmin", 0.9f)->setComment("Cut on minimum pt");
   desc.add<double>("CAThetaCutBarrel", 0.002f)->setComment("Cut on RZ alignement for Barrel");
   desc.add<double>("CAThetaCutForward", 0.003f)->setComment("Cut on RZ alignment for Forward");
   desc.add<double>("hardCurvCut", 1.f / (0.35 * 87.f))->setComment("Cut on minimum curvature");
   desc.add<double>("dcaCutInnerTriplet", 0.15f)->setComment("Cut on origin radius when the inner hit is on BPix1");
   desc.add<double>("dcaCutOuterTriplet", 0.25f)->setComment("Cut on origin radius when the outer hit is on BPix1");
   desc.add<bool>("earlyFishbone", true);
   desc.add<bool>("lateFishbone", false);
   desc.add<bool>("idealConditions", true);
   desc.add<bool>("fillStatistics", false);
   desc.add<unsigned int>("minHitsPerNtuplet", 4);
   desc.add<unsigned int>("maxNumberOfDoublets", caConstants::maxNumberOfDoublets);
   desc.add<unsigned int>("minHitsForSharingCut", 5)
       ->setComment("Maximum number of hits in a tuple to clean also if the shared hit is on bpx1");
   desc.add<bool>("includeJumpingForwardDoublets", false);
   desc.add<bool>("fit5as4", true);
   desc.add<bool>("doClusterCut", true);
   desc.add<bool>("doZ0Cut", true);
   desc.add<bool>("doPtCut", true);
   desc.add<bool>("useRiemannFit", false)->setComment("true for Riemann, false for BrokenLine");
   desc.add<bool>("doSharedHitCut", true)->setComment("Sharing hit nTuples cleaning");
   desc.add<bool>("dupPassThrough", false)->setComment("Do not reject duplicate");
   desc.add<bool>("useSimpleTripletCleaner", false)->setComment("use alternate implementation");
  
   edm::ParameterSetDescription trackQualityCuts;
   trackQualityCuts.add<double>("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut");
   trackQualityCuts.add<std::vector<double>>("chi2Coeff", {0.9, 1.8})->setComment("chi2 at 1GeV and at ptMax above");
   trackQualityCuts.add<double>("chi2Scale", 8.)
       ->setComment(
           "Factor to multiply the pT-dependent chi2 cut (currently: 8 for the broken line fit, ?? for the Riemann "
           "fit)");
   trackQualityCuts.add<double>("tripletMinPt", 0.5)->setComment("Min pT for triplets, in GeV");
   trackQualityCuts.add<double>("tripletMaxTip", 0.3)->setComment("Max |Tip| for triplets, in cm");
   trackQualityCuts.add<double>("tripletMaxZip", 12.)->setComment("Max |Zip| for triplets, in cm");
   trackQualityCuts.add<double>("quadrupletMinPt", 0.3)->setComment("Min pT for quadruplets, in GeV");
   trackQualityCuts.add<double>("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm");
   trackQualityCuts.add<double>("quadrupletMaxZip", 12.)->setComment("Max |Zip| for quadruplets, in cm");
   desc.add<edm::ParameterSetDescription>("trackQualityCuts", trackQualityCuts)
       ->setComment(
           "Quality cuts based on the results of the track fit:\n  - apply a pT-dependent chi2 cut;\n  - apply \"region "
           "cuts\" based on the fit results (pT, Tip, Zip).");
 }

References edm::ParameterSetDescription::add(), submitPVResolutionJobs::desc, caConstants::maxNumberOfDoublets, and edm::ParameterDescriptionNode::setComment().

Referenced by CAHitNtupletCUDA::fillDescriptions().

◆ fillDescriptionsLabel()

static const char* CAHitNtupletGeneratorOnGPU::fillDescriptionsLabel ( )

inlinestatic

Definition at line 47 of file CAHitNtupletGeneratorOnGPU.h.

47 { return "caHitNtupletOnGPU"; }

◆ hitNtuplets()

void CAHitNtupletGeneratorOnGPU::hitNtuplets	(	HitsOnCPU const &	hh,
		const edm::EventSetup &	es,
		bool	useRiemannFit,
		cudaStream_t	cudaStream
	)

private

◆ launchKernels()

void CAHitNtupletGeneratorOnGPU::launchKernels	(	HitsOnCPU const &	hh,
		bool	useRiemannFit,
		cudaStream_t	cudaStream
	)		const

private

◆ makeTuples()

PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuples	(	TrackingRecHit2DCPU const &	hits_d,
		float	bfield
	)		const

Definition at line 217 of file CAHitNtupletGeneratorOnGPU.cc.

                                                                                                                     {
   PixelTrackHeterogeneous tracks(std::make_unique<pixelTrack::TrackSoA>());
  
   auto* soa = tracks.get();
   assert(soa);
  
   CAHitNtupletGeneratorKernelsCPU kernels(m_params);
   kernels.setCounters(m_counters);
   kernels.allocateOnGPU(hits_d.nHits(), nullptr);
  
   kernels.buildDoublets(hits_d, nullptr);
   kernels.launchKernels(hits_d, soa, nullptr);
   kernels.fillHitDetIndices(hits_d.view(), soa, nullptr);  // in principle needed only if Hits not "available"
  
   if (0 == hits_d.nHits())
     return tracks;
  
   // now fit
   HelixFitOnGPU fitter(bfield, m_params.fit5as4_);
   fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
  
   if (m_params.useRiemannFit_) {
     fitter.launchRiemannKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
   } else {
     fitter.launchBrokenLineKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
   }
  
   kernels.classifyTuples(hits_d, soa, nullptr);
  
 #ifdef GPU_DEBUG
   std::cout << "finished building pixel tracks on CPU" << std::endl;
 #endif
  
   return tracks;
 }

References HelixFitOnGPU::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), cms::cuda::assert(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), gather_cfg::cout, CAHitNtupletGeneratorKernels< TTraits >::fillHitDetIndices(), cAHitNtupletGenerator::Params::fit5as4_, HelixFitOnGPU::launchBrokenLineKernelsOnCPU(), CAHitNtupletGeneratorKernels< TTraits >::launchKernels(), HelixFitOnGPU::launchRiemannKernelsOnCPU(), m_counters, m_params, caConstants::maxNumberOfQuadruplets, TrackingRecHit2DHeterogeneous< Traits >::nHits(), CAHitNtupletGeneratorKernels< TTraits >::setCounters(), gpuVertexFinder::soa, tracks, CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity(), cAHitNtupletGenerator::Params::useRiemannFit_, and TrackingRecHit2DHeterogeneous< Traits >::view().

Referenced by CAHitNtupletCUDA::produce().

◆ makeTuplesAsync()

PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuplesAsync	(	TrackingRecHit2DGPU const &	hits_d,
		float	bfield,
		cudaStream_t	stream
	)		const

Definition at line 183 of file CAHitNtupletGeneratorOnGPU.cc.

                                                                                                {
   PixelTrackHeterogeneous tracks(cms::cuda::make_device_unique<pixelTrack::TrackSoA>(stream));
  
   auto* soa = tracks.get();
   assert(soa);
  
   CAHitNtupletGeneratorKernelsGPU kernels(m_params);
   kernels.setCounters(m_counters);
   kernels.allocateOnGPU(hits_d.nHits(), stream);
  
   kernels.buildDoublets(hits_d, stream);
   kernels.launchKernels(hits_d, soa, stream);
   kernels.fillHitDetIndices(hits_d.view(), soa, stream);  // in principle needed only if Hits not "available"
  
   HelixFitOnGPU fitter(bfield, m_params.fit5as4_);
   fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
   if (m_params.useRiemannFit_) {
     fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
   } else {
     fitter.launchBrokenLineKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
   }
   kernels.classifyTuples(hits_d, soa, stream);
  
 #ifdef GPU_DEBUG
   cudaDeviceSynchronize();
   cudaCheck(cudaGetLastError());
   std::cout << "finished building pixel tracks on GPU" << std::endl;
 #endif
  
   return tracks;
 }