CAHitNtupletGeneratorOnGPU Class Reference

#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
void	beginJob ()
	CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &&iC)
	CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &iC)
void	endJob ()
PixelTrackHeterogeneous	makeTuples (TrackingRecHit2DCPU const &hits_d, float bfield) const
PixelTrackHeterogeneous	makeTuplesAsync (TrackingRecHit2DGPU const &hits_d, float bfield, cudaStream_t stream) const

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.

      : 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>("fitNas4"),
               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 %s %s\n",
         "tid",
         "qual",
         "nh",
         "nl",
         "charge",
         "pt",
         "eta",
         "phi",
         "tip",
         "zip",
         "chi2",
         "h1",
         "h2",
         "h3",
         "h4",
         "h5",
         "hn");
#endif
}

Member Function Documentation

beginJob()

void CAHitNtupletGeneratorOnGPU::beginJob

(

void

)

Definition at line 154 of file CAHitNtupletGeneratorOnGPU.cc.

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

Referenced by CAHitNtupletCUDA::beginJob().

                                          {
  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));
  }
}

buildDoublets()

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

private

endJob()

void CAHitNtupletGeneratorOnGPU::endJob

(

void

)

Definition at line 168 of file CAHitNtupletGeneratorOnGPU.cc.

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

Referenced by CAHitNtupletCUDA::endJob().

                                        {
  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;
  }
}

fillDescriptions()

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

static

Definition at line 107 of file CAHitNtupletGeneratorOnGPU.cc.

References submitPVResolutionJobs::desc, caConstants::maxNumberOfDoublets, and HLT_2022v15_cff::trackQualityCuts.

Referenced by CAHitNtupletCUDA::fillDescriptions().

                                                                                  {
  // 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", 10)
      ->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>("fitNas4", false)->setComment("fit only 4 hits out of N");
  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", true)->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).");
}

fillDescriptionsLabel()

static const char* CAHitNtupletGeneratorOnGPU::fillDescriptionsLabel ( )

inlinestatic

Definition at line 45 of file CAHitNtupletGeneratorOnGPU.h.

{ 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 220 of file CAHitNtupletGeneratorOnGPU.cc.

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

Referenced by CAHitNtupletCUDA::produce().

                                                                                                                    {
  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);

  if (0 == hits_d.nHits())
    return tracks;

  // now fit
  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
  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

  // check that the fixed-size SoA does not overflow
  auto const& tsoa = *soa;
  auto maxTracks = tsoa.stride();
  auto nTracks = tsoa.nTracks();
  assert(nTracks < maxTracks);
  if (nTracks == maxTracks - 1) {
    edm::LogWarning("PixelTracks") << "Unsorted reconstructed pixel tracks truncated to " << maxTracks - 1
                                   << " candidates";
  }

  return tracks;
}

makeTuplesAsync()

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

Definition at line 187 of file CAHitNtupletGeneratorOnGPU.cc.

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

Referenced by CAHitNtupletCUDA::produce().

                                                                                               {
  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);

  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
  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;
}

Member Data Documentation

m_counters

Counters* CAHitNtupletGeneratorOnGPU::m_counters = nullptr

private

Definition at line 63 of file CAHitNtupletGeneratorOnGPU.h.

Referenced by beginJob(), endJob(), makeTuples(), and makeTuplesAsync().

m_params

Params CAHitNtupletGeneratorOnGPU::m_params

private

Definition at line 61 of file CAHitNtupletGeneratorOnGPU.h.

Referenced by beginJob(), endJob(), makeTuples(), and makeTuplesAsync().