CAHitNtupletGeneratorOnGPU< TrackerTraits > Class Template Reference

#include <CAHitNtupletGeneratorOnGPU.h>

Public Types
using	CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
using	CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
using	Counters = caHitNtupletGenerator::Counters
using	GPUCACell = GPUCACellT< TrackerTraits >
using	hindex_type = typename TrackingRecHitSoA< TrackerTraits >::hindex_type
using	HitContainer = typename TrackSoA< TrackerTraits >::HitContainer
using	HitsConstView = TrackingRecHitSoAConstView< TrackerTraits >
using	HitsOnDevice = TrackingRecHitSoADevice< TrackerTraits >
using	HitsOnHost = TrackingRecHitSoAHost< TrackerTraits >
using	HitsView = TrackingRecHitSoAView< TrackerTraits >
using	HitToTuple = caStructures::HitToTupleT< TrackerTraits >
using	OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
using	Params = caHitNtupletGenerator::ParamsT< TrackerTraits >
using	Quality = pixelTrack::Quality
using	QualityCuts = pixelTrack::QualityCutsT< TrackerTraits >
using	TrackSoADevice = TrackSoAHeterogeneousDevice< TrackerTraits >
using	TrackSoAHost = TrackSoAHeterogeneousHost< TrackerTraits >
using	Tuple = HitContainer
using	TupleMultiplicity = caStructures::TupleMultiplicityT< TrackerTraits >

Public Member Functions
void	beginJob ()
	CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &&iC)
	CAHitNtupletGeneratorOnGPU (const edm::ParameterSet &cfg, edm::ConsumesCollector &iC)
void	endJob ()
template<>
void	fillDescriptions (edm::ParameterSetDescription &desc)
template<>
void	fillDescriptions (edm::ParameterSetDescription &desc)
template<>
void	fillDescriptions (edm::ParameterSetDescription &desc)
TrackSoAHost	makeTuples (HitsOnHost const &hits_d, float bfield) const
TrackSoADevice	makeTuplesAsync (HitsOnDevice const &hits_d, float bfield, cudaStream_t stream) const

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

Private Member Functions
void	buildDoublets (const HitsConstView &hh, cudaStream_t stream) const
void	hitNtuplets (const HitsConstView &hh, const edm::EventSetup &es, bool useRiemannFit, cudaStream_t cudaStream)
void	launchKernels (const HitsConstView &hh, bool useRiemannFit, cudaStream_t cudaStream) const

Private Attributes
Counters *	m_counters = nullptr
Params	m_params

Detailed Description

template<typename TrackerTraits>
class CAHitNtupletGeneratorOnGPU< TrackerTraits >

Definition at line 28 of file CAHitNtupletGeneratorOnGPU.h.

Member Typedef Documentation

CellNeighborsVector

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::CellNeighborsVector = caStructures::CellNeighborsVectorT<TrackerTraits>

Definition at line 46 of file CAHitNtupletGeneratorOnGPU.h.

CellTracksVector

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::CellTracksVector = caStructures::CellTracksVectorT<TrackerTraits>

Definition at line 47 of file CAHitNtupletGeneratorOnGPU.h.

Counters

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::Counters = caHitNtupletGenerator::Counters

Definition at line 53 of file CAHitNtupletGeneratorOnGPU.h.

GPUCACell

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::GPUCACell = GPUCACellT<TrackerTraits>

Definition at line 40 of file CAHitNtupletGeneratorOnGPU.h.

hindex_type

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::hindex_type = typename TrackingRecHitSoA<TrackerTraits>::hindex_type

Definition at line 34 of file CAHitNtupletGeneratorOnGPU.h.

HitContainer

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::HitContainer = typename TrackSoA<TrackerTraits>::HitContainer

Definition at line 43 of file CAHitNtupletGeneratorOnGPU.h.

HitsConstView

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::HitsConstView = TrackingRecHitSoAConstView<TrackerTraits>

Definition at line 31 of file CAHitNtupletGeneratorOnGPU.h.

HitsOnDevice

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::HitsOnDevice = TrackingRecHitSoADevice<TrackerTraits>

Definition at line 32 of file CAHitNtupletGeneratorOnGPU.h.

HitsOnHost

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::HitsOnHost = TrackingRecHitSoAHost<TrackerTraits>

Definition at line 33 of file CAHitNtupletGeneratorOnGPU.h.

HitsView

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::HitsView = TrackingRecHitSoAView<TrackerTraits>

Definition at line 30 of file CAHitNtupletGeneratorOnGPU.h.

HitToTuple

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::HitToTuple = caStructures::HitToTupleT<TrackerTraits>

Definition at line 36 of file CAHitNtupletGeneratorOnGPU.h.

OuterHitOfCell

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::OuterHitOfCell = caStructures::OuterHitOfCellT<TrackerTraits>

Definition at line 38 of file CAHitNtupletGeneratorOnGPU.h.

Params

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::Params = caHitNtupletGenerator::ParamsT<TrackerTraits>

Definition at line 52 of file CAHitNtupletGeneratorOnGPU.h.

Quality

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::Quality = pixelTrack::Quality

Definition at line 49 of file CAHitNtupletGeneratorOnGPU.h.

QualityCuts

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::QualityCuts = pixelTrack::QualityCutsT<TrackerTraits>

Definition at line 51 of file CAHitNtupletGeneratorOnGPU.h.

TrackSoADevice

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::TrackSoADevice = TrackSoAHeterogeneousDevice<TrackerTraits>

Definition at line 42 of file CAHitNtupletGeneratorOnGPU.h.

TrackSoAHost

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::TrackSoAHost = TrackSoAHeterogeneousHost<TrackerTraits>

Definition at line 41 of file CAHitNtupletGeneratorOnGPU.h.

Tuple

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::Tuple = HitContainer

Definition at line 44 of file CAHitNtupletGeneratorOnGPU.h.

TupleMultiplicity

template<typename TrackerTraits >

using CAHitNtupletGeneratorOnGPU< TrackerTraits >::TupleMultiplicity = caStructures::TupleMultiplicityT<TrackerTraits>

Definition at line 37 of file CAHitNtupletGeneratorOnGPU.h.

Constructor & Destructor Documentation

CAHitNtupletGeneratorOnGPU() [1/2]

template<typename TrackerTraits >

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

inline

Definition at line 56 of file CAHitNtupletGeneratorOnGPU.h.

      : CAHitNtupletGeneratorOnGPU(cfg, iC) {}

CAHitNtupletGeneratorOnGPU() [2/2]

template<typename TrackerTraits >

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

Definition at line 141 of file CAHitNtupletGeneratorOnGPU.cc.

    : m_params(makeCommonParams(cfg),
               makeCellCuts<TrackerTraits>(cfg),
               topologyCuts<TrackerTraits>::makeQualityCuts(cfg.getParameterSet("trackQualityCuts")),
               topologyCuts<TrackerTraits>::makeCACuts(cfg)) {
#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()

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::beginJob

(

void

)

Definition at line 299 of file CAHitNtupletGeneratorOnGPU.cc.

References ecalDigis_cff::cuda, and cudaCheck.

                                                         {
  if (m_params.onGPU_) {
    // allocate pinned host memory only if CUDA is available
    edm::Service<CUDAInterface> cuda;
    if (cuda and cuda->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()

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::buildDoublets ( const HitsConstView &  hh, cudaStream_t  stream  ) const

private

endJob()

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::endJob

(

void

)

Definition at line 314 of file CAHitNtupletGeneratorOnGPU.cc.

References ecalDigis_cff::cuda, CAHitNtupletGeneratorKernelsGPU< TrackerTraits >::printCounters(), and CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::printCounters().

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

fillDescriptions() [1/4]

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::fillDescriptions ( edm::ParameterSetDescription &  desc )

static

Definition at line 170 of file CAHitNtupletGeneratorOnGPU.cc.

References submitPVResolutionJobs::desc.

                                                                                                 {
  fillDescriptionsCommon(desc);
  edm::LogWarning("CAHitNtupletGeneratorOnGPU::fillDescriptions")
      << "Note: this fillDescriptions is a dummy one. Most probably you are missing some parameters. \n"
         "please implement your TrackerTraits descriptions in CAHitNtupletGeneratorOnGPU. \n";
}

fillDescriptions() [2/4]

template<>

void CAHitNtupletGeneratorOnGPU< pixelTopology::Phase1 >::fillDescriptions

(

edm::ParameterSetDescription &

desc

)

Definition at line 178 of file CAHitNtupletGeneratorOnGPU.cc.

References submitPVResolutionJobs::desc, phase1PixelTopology::phicuts, and HLT_2024v12_cff::trackQualityCuts.

                                                                                                       {
  fillDescriptionsCommon(desc);

  desc.add<bool>("idealConditions", true);
  desc.add<bool>("includeJumpingForwardDoublets", false);
  desc.add<double>("z0Cut", 12.0);
  desc.add<double>("ptCut", 0.5);

  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<std::vector<int>>(
          "phiCuts", std::vector<int>(std::begin(phase1PixelTopology::phicuts), std::end(phase1PixelTopology::phicuts)))
      ->setComment("Cuts in phi for cells");

  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).");
}

fillDescriptions() [3/4]

template<>

void CAHitNtupletGeneratorOnGPU< pixelTopology::HIonPhase1 >::fillDescriptions

(

edm::ParameterSetDescription &

desc

)

Definition at line 211 of file CAHitNtupletGeneratorOnGPU.cc.

References submitPVResolutionJobs::desc, phase1PixelTopology::phicuts, and HLT_2024v12_cff::trackQualityCuts.

                                                                                                           {
  fillDescriptionsCommon(desc);

  desc.add<bool>("idealConditions", false);
  desc.add<bool>("includeJumpingForwardDoublets", false);
  desc.add<double>("z0Cut", 10.0);
  desc.add<double>("ptCut", 0.0);

  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.0)->setComment("Min pT for triplets, in GeV");
  trackQualityCuts.add<double>("tripletMaxTip", 0.1)->setComment("Max |Tip| for triplets, in cm");
  trackQualityCuts.add<double>("tripletMaxZip", 6.)->setComment("Max |Zip| for triplets, in cm");
  trackQualityCuts.add<double>("quadrupletMinPt", 0.0)->setComment("Min pT for quadruplets, in GeV");
  trackQualityCuts.add<double>("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm");
  trackQualityCuts.add<double>("quadrupletMaxZip", 6.)->setComment("Max |Zip| for quadruplets, in cm");

  desc.add<std::vector<int>>(
          "phiCuts", std::vector<int>(std::begin(phase1PixelTopology::phicuts), std::end(phase1PixelTopology::phicuts)))
      ->setComment("Cuts in phi for cells");

  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).");
}

fillDescriptions() [4/4]

template<>

void CAHitNtupletGeneratorOnGPU< pixelTopology::Phase2 >::fillDescriptions

(

edm::ParameterSetDescription &

desc

)

Definition at line 244 of file CAHitNtupletGeneratorOnGPU.cc.

References submitPVResolutionJobs::desc, phase2PixelTopology::phicuts, and HLT_2024v12_cff::trackQualityCuts.

                                                                                                       {
  fillDescriptionsCommon(desc);

  desc.add<bool>("idealConditions", false);
  desc.add<bool>("includeFarForwards", true);
  desc.add<bool>("includeJumpingForwardDoublets", true);
  desc.add<double>("z0Cut", 7.5);
  desc.add<double>("ptCut", 0.85);

  edm::ParameterSetDescription trackQualityCuts;
  trackQualityCuts.add<double>("maxChi2", 5.)->setComment("Max normalized chi2");
  trackQualityCuts.add<double>("minPt", 0.5)->setComment("Min pT in GeV");
  trackQualityCuts.add<double>("maxTip", 0.3)->setComment("Max |Tip| in cm");
  trackQualityCuts.add<double>("maxZip", 12.)->setComment("Max |Zip|, in cm");

  desc.add<std::vector<int>>(
          "phiCuts", std::vector<int>(std::begin(phase2PixelTopology::phicuts), std::end(phase2PixelTopology::phicuts)))
      ->setComment("Cuts in phi for cells");

  desc.add<edm::ParameterSetDescription>("trackQualityCuts", trackQualityCuts)
      ->setComment(
          "Quality cuts based on the results of the track fit:\n  - apply cuts based on the fit results (pT, Tip, "
          "Zip).");
}

fillDescriptionsCommon()

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::fillDescriptionsCommon ( edm::ParameterSetDescription &  desc )

static

Definition at line 270 of file CAHitNtupletGeneratorOnGPU.cc.

References submitPVResolutionJobs::desc, and HLT_2024v12_cff::maxNumberOfDoublets.

                                                                                                       {
  // 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.9)->setComment("Cut on minimum pt");
  desc.add<double>("CAThetaCutBarrel", 0.002)->setComment("Cut on RZ alignement for Barrel");
  desc.add<double>("CAThetaCutForward", 0.003)->setComment("Cut on RZ alignment for Forward");
  desc.add<double>("hardCurvCut", 1. / (0.35 * 87.))->setComment("Cut on minimum curvature");
  desc.add<double>("dcaCutInnerTriplet", 0.15)->setComment("Cut on origin radius when the inner hit is on BPix1");
  desc.add<double>("dcaCutOuterTriplet", 0.25)->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>("fillStatistics", false);
  desc.add<unsigned int>("minHitsPerNtuplet", 4);
  desc.add<unsigned int>("maxNumberOfDoublets", TrackerTraits::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>("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");
}

hitNtuplets()

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::hitNtuplets ( const HitsConstView &  hh, const edm::EventSetup &  es, bool  useRiemannFit, cudaStream_t  cudaStream  )

private

launchKernels()

template<typename TrackerTraits >

void CAHitNtupletGeneratorOnGPU< TrackerTraits >::launchKernels ( const HitsConstView &  hh, bool  useRiemannFit, cudaStream_t  cudaStream  ) const

private

makeTuples()

template<typename TrackerTraits >

TrackSoAHeterogeneousHost< TrackerTraits > CAHitNtupletGeneratorOnGPU< TrackerTraits >::makeTuples	(	HitsOnHost const &	hits_d,
		float	bfield
	)		const

Definition at line 368 of file CAHitNtupletGeneratorOnGPU.cc.

References HelixFitOnGPU< TrackerTraits >::allocateOnGPU(), cms::cuda::assert(), gather_cfg::cout, HelixFitOnGPU< TrackerTraits >::launchBrokenLineKernelsOnCPU(), HelixFitOnGPU< TrackerTraits >::launchRiemannKernelsOnCPU(), DMR_cfg::maxTracks, TrackingRecHitSoAHost< TrackerTraits >::nHits(), beamSpotPI::nTracks, TrackingRecHitSoAHost< TrackerTraits >::offsetBPIX2(), DiMuonV_cfg::tracks, and cms::cuda::PortableHostCollection< T >::view().

                                                                                                                   {
  using HelixFitOnGPU = HelixFitOnGPU<TrackerTraits>;
  using TrackSoA = TrackSoAHeterogeneousHost<TrackerTraits>;
  using CPUKernels = CAHitNtupletGeneratorKernelsCPU<TrackerTraits>;

  TrackSoA tracks;

  CPUKernels kernels(m_params);
  kernels.setCounters(m_counters);
  kernels.allocateOnGPU(hits_h.nHits(), nullptr);

  kernels.buildDoublets(hits_h.view(), hits_h.offsetBPIX2(), nullptr);
  kernels.launchKernels(hits_h.view(), tracks.view(), nullptr);

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

  // now fit
  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
  fitter.allocateOnGPU(kernels.tupleMultiplicity(), tracks.view());

  if (m_params.useRiemannFit_) {
    fitter.launchRiemannKernelsOnCPU(hits_h.view(), hits_h.nHits(), TrackerTraits::maxNumberOfQuadruplets);
  } else {
    fitter.launchBrokenLineKernelsOnCPU(hits_h.view(), hits_h.nHits(), TrackerTraits::maxNumberOfQuadruplets);
  }

  kernels.classifyTuples(hits_h.view(), tracks.view(), 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 maxTracks = tracks.view().metadata().size();
  auto nTracks = tracks.view().nTracks();
  assert(nTracks < maxTracks);
  if (nTracks == maxTracks - 1) {
    edm::LogWarning("PixelTracks") << "Unsorted reconstructed pixel tracks truncated to " << maxTracks - 1
                                   << " candidates";
  }

  return tracks;
}

makeTuplesAsync()

template<typename TrackerTraits >

TrackSoAHeterogeneousDevice< TrackerTraits > CAHitNtupletGeneratorOnGPU< TrackerTraits >::makeTuplesAsync	(	HitsOnDevice const &	hits_d,
		float	bfield,
		cudaStream_t	stream
	)		const

Definition at line 334 of file CAHitNtupletGeneratorOnGPU.cc.

References HelixFitOnGPU< TrackerTraits >::allocateOnGPU(), gather_cfg::cout, cudaCheck, HelixFitOnGPU< TrackerTraits >::launchBrokenLineKernels(), HelixFitOnGPU< TrackerTraits >::launchRiemannKernels(), TrackingRecHitSoADevice< TrackerTraits >::nHits(), TrackingRecHitSoADevice< TrackerTraits >::offsetBPIX2(), cms::cuda::stream, DiMuonV_cfg::tracks, and cms::cuda::PortableDeviceCollection< T >::view().

                                                                         {
  using HelixFitOnGPU = HelixFitOnGPU<TrackerTraits>;
  using TrackSoA = TrackSoAHeterogeneousDevice<TrackerTraits>;
  using GPUKernels = CAHitNtupletGeneratorKernelsGPU<TrackerTraits>;

  TrackSoA tracks(stream);

  GPUKernels kernels(m_params);
  kernels.setCounters(m_counters);
  kernels.allocateOnGPU(hits_d.nHits(), stream);

  kernels.buildDoublets(hits_d.view(), hits_d.offsetBPIX2(), stream);

  kernels.launchKernels(hits_d.view(), tracks.view(), stream);

  HelixFitOnGPU fitter(bfield, m_params.fitNas4_);
  fitter.allocateOnGPU(kernels.tupleMultiplicity(), tracks.view());
  if (m_params.useRiemannFit_) {
    fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), TrackerTraits::maxNumberOfQuadruplets, stream);
  } else {
    fitter.launchBrokenLineKernels(hits_d.view(), hits_d.nHits(), TrackerTraits::maxNumberOfQuadruplets, stream);
  }
  kernels.classifyTuples(hits_d.view(), tracks.view(), 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

template<typename TrackerTraits >

Counters* CAHitNtupletGeneratorOnGPU< TrackerTraits >::m_counters = nullptr

private

Definition at line 79 of file CAHitNtupletGeneratorOnGPU.h.

m_params

template<typename TrackerTraits >

Params CAHitNtupletGeneratorOnGPU< TrackerTraits >::m_params

private

Definition at line 77 of file CAHitNtupletGeneratorOnGPU.h.