CAHitNtupletGeneratorKernels< TTraits > Class Template Reference

#include <CAHitNtupletGeneratorKernels.h>

Public Types
using	Counters = cAHitNtupletGenerator::Counters
using	HitContainer = pixelTrack::HitContainer
using	HitsOnCPU = TrackingRecHit2DHeterogeneous< Traits >
using	HitsOnGPU = TrackingRecHit2DSOAView
using	HitsView = TrackingRecHit2DSOAView
using	HitToTuple = caConstants::HitToTuple
using	Params = cAHitNtupletGenerator::Params
using	Quality = pixelTrack::Quality
using	QualityCuts = cAHitNtupletGenerator::QualityCuts
using	TkSoA = pixelTrack::TrackSoA
using	Traits = TTraits
using	TupleMultiplicity = caConstants::TupleMultiplicity
template<typename T >
using	unique_ptr = typename Traits::template unique_ptr< T >

Public Member Functions
void	allocateOnGPU (int32_t nHits, cudaStream_t stream)
void	buildDoublets (HitsOnCPU const &hh, cudaStream_t stream)
	CAHitNtupletGeneratorKernels (Params const &params)
void	classifyTuples (HitsOnCPU const &hh, TkSoA *tuples_d, cudaStream_t cudaStream)
void	cleanup (cudaStream_t cudaStream)
void	launchKernels (HitsOnCPU const &hh, TkSoA *tuples_d, cudaStream_t cudaStream)
void	setCounters (Counters *counters)
TupleMultiplicity const *	tupleMultiplicity () const
	~CAHitNtupletGeneratorKernels ()=default

Static Public Member Functions
static void	printCounters (Counters const *counters)

Private Member Functions
uint32_t	nDoubletBlocks (uint32_t blockSize)
	Compute the number of doublet blocks for block size. More...
uint32_t	nQuadrupletBlocks (uint32_t blockSize)
	Compute the number of quadruplet blocks for block size. More...

Private Attributes
unique_ptr< unsigned char[]>	cellStorage_
Counters *	counters_ = nullptr
unique_ptr< HitToTuple >	device_hitToTuple_
cms::cuda::AtomicPairCounter *	device_hitToTuple_apc_ = nullptr
unique_ptr< HitToTuple::Counter[]>	device_hitToTupleStorage_
cms::cuda::AtomicPairCounter *	device_hitTuple_apc_ = nullptr
unique_ptr < GPUCACell::OuterHitOfCellContainer[]>	device_isOuterHitOfCell_
uint32_t *	device_nCells_ = nullptr
unique_ptr < cms::cuda::AtomicPairCounter::c_type[]>	device_storage_
unique_ptr < caConstants::CellNeighborsVector >	device_theCellNeighbors_
caConstants::CellNeighbors *	device_theCellNeighborsContainer_
unique_ptr< GPUCACell[]>	device_theCells_
unique_ptr < caConstants::CellTracksVector >	device_theCellTracks_
caConstants::CellTracks *	device_theCellTracksContainer_
unique_ptr< TupleMultiplicity >	device_tupleMultiplicity_
HitToTuple::View	hitToTupleView_
GPUCACell::OuterHitOfCell	isOuterHitOfCell_
Params const &	params_
const uint32_t	paramsMaxDoubletes3Quarters_
	Intermediate result avoiding repeated computations. More...

Detailed Description

template<typename TTraits>
class CAHitNtupletGeneratorKernels< TTraits >

Definition at line 156 of file CAHitNtupletGeneratorKernels.h.

Member Typedef Documentation

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::Counters = cAHitNtupletGenerator::Counters

Definition at line 162 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::HitContainer = pixelTrack::HitContainer

Definition at line 176 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::HitsOnCPU = TrackingRecHit2DHeterogeneous<Traits>

Definition at line 169 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 168 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::HitsView = TrackingRecHit2DSOAView

Definition at line 167 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::HitToTuple = caConstants::HitToTuple

Definition at line 171 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::Params = cAHitNtupletGenerator::Params

Definition at line 161 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::Quality = pixelTrack::Quality

Definition at line 174 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::QualityCuts = cAHitNtupletGenerator::QualityCuts

Definition at line 160 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::TkSoA = pixelTrack::TrackSoA

Definition at line 175 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::Traits = TTraits

Definition at line 158 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

using CAHitNtupletGeneratorKernels< TTraits >::TupleMultiplicity = caConstants::TupleMultiplicity

Definition at line 172 of file CAHitNtupletGeneratorKernels.h.

template<typename TTraits >

template<typename T >

using CAHitNtupletGeneratorKernels< TTraits >::unique_ptr = typename Traits::template unique_ptr<T>

Definition at line 165 of file CAHitNtupletGeneratorKernels.h.

Constructor & Destructor Documentation

template<typename TTraits >

CAHitNtupletGeneratorKernels< TTraits >::CAHitNtupletGeneratorKernels ( Params const &  params )

inline

Definition at line 178 of file CAHitNtupletGeneratorKernels.h.

      : params_(params), paramsMaxDoubletes3Quarters_(3 * params.maxNumberOfDoublets_ / 4) {}

template<typename TTraits >

CAHitNtupletGeneratorKernels< TTraits >::~CAHitNtupletGeneratorKernels ( )

default

Member Function Documentation

void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU	(	int32_t	nHits,
		cudaStream_t	stream
	)

Definition at line 9 of file CAHitNtupletGeneratorKernelsAlloc.cc.

References cms::cuda::assert(), cms::cuda::OneToManyAssocView< Assoc >::assoc, gather_cfg::cout, cudaCheck, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTupleStorage_, CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_storage_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_, cms::cuda::AtomicPairCounter::get(), CAHitNtupletGeneratorKernels< TTraits >::hitToTupleView_, nHits, cms::cuda::OneToManyAssocView< Assoc >::offSize, cms::cuda::OneToManyAssocView< Assoc >::offStorage, cms::cuda::stream, svgfig::template(), and relativeConstraints::value.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

                                                                                      {
#endif
  // ALLOCATIONS FOR THE INTERMEDIATE RESULTS (STAYS ON WORKER)

  device_theCellNeighbors_ = Traits::template make_unique<caConstants::CellNeighborsVector>(stream);
  device_theCellTracks_ = Traits::template make_unique<caConstants::CellTracksVector>(stream);

#ifdef GPU_DEBUG
  std::cout << "Allocation for tuple building. N hits " << nHits << std::endl;
#endif

  nHits++;  // storage requires one more counter;
  assert(nHits > 0);
  device_hitToTuple_ = Traits::template make_unique<HitToTuple>(stream);
  device_hitToTupleStorage_ = Traits::template make_unique<HitToTuple::Counter[]>(nHits, stream);
  hitToTupleView_.assoc = device_hitToTuple_.get();
  hitToTupleView_.offStorage = device_hitToTupleStorage_.get();
  hitToTupleView_.offSize = nHits;

  device_tupleMultiplicity_ = Traits::template make_unique<TupleMultiplicity>(stream);

  device_storage_ = Traits::template make_unique<cms::cuda::AtomicPairCounter::c_type[]>(3, stream);

  device_hitTuple_apc_ = (cms::cuda::AtomicPairCounter*)device_storage_.get();
  device_hitToTuple_apc_ = (cms::cuda::AtomicPairCounter*)device_storage_.get() + 1;
  device_nCells_ = (uint32_t*)(device_storage_.get() + 2);

  // FIXME: consider collapsing these 3 in one adhoc kernel
  if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
    cudaCheck(cudaMemsetAsync(device_nCells_, 0, sizeof(uint32_t), stream));
  } else {
    *device_nCells_ = 0;
  }
  cms::cuda::launchZero(device_tupleMultiplicity_.get(), stream);
  cms::cuda::launchZero(hitToTupleView_, stream);  // we may wish to keep it in the edm
#ifdef GPU_DEBUG
  cudaDeviceSynchronize();
  cudaCheck(cudaGetLastError());
#endif
}

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

Definition at line 17 of file CAHitNtupletGeneratorKernels.cc.

References cms::cuda::assert(), CAHitNtupletGeneratorKernels< TTraits >::cellStorage_, gather_cfg::cout, CAHitNtupletGeneratorKernels< TTraits >::device_isOuterHitOfCell_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighborsContainer_, CAHitNtupletGeneratorKernels< TTraits >::device_theCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracksContainer_, cAHitNtupletGenerator::Params::doClusterCut_, cAHitNtupletGenerator::Params::doPtCut_, cAHitNtupletGenerator::Params::doZ0Cut_, cAHitNtupletGenerator::Params::idealConditions_, cAHitNtupletGenerator::Params::includeJumpingForwardDoublets_, CAHitNtupletGeneratorKernels< TTraits >::isOuterHitOfCell_, SiStripPI::max, cAHitNtupletGenerator::Params::maxNumberOfDoublets_, caConstants::maxNumOfActiveDoublets, cAHitNtupletGenerator::Params::minHitsPerNtuplet_, gpuPixelDoublets::nActualPairs, TrackingRecHit2DHeterogeneous< Traits >::nHits(), nhits, gpuPixelDoublets::nPairs, gpuPixelDoublets::nPairsForQuadruplets, gpuPixelDoublets::nPairsForTriplets, TrackingRecHit2DHeterogeneous< Traits >::offsetBPIX2(), CAHitNtupletGeneratorKernels< TTraits >::params_, and TrackingRecHit2DHeterogeneous< Traits >::view().

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

                                                                                            {
  auto nhits = hh.nHits();

#ifdef NTUPLE_DEBUG
  std::cout << "building Doublets out of " << nhits << " Hits. BPIX2 offset is " << hh.offsetBPIX2() << std::endl;
#endif

  // use "nhits" to heuristically dimension the workspace

  // no need to use the Traits allocations, since we know this is being compiled for the CPU
  //device_isOuterHitOfCell_ = Traits::template make_unique<GPUCACell::OuterHitOfCell[]>(std::max(1U, nhits), stream);
  device_isOuterHitOfCell_ = std::make_unique<GPUCACell::OuterHitOfCellContainer[]>(std::max(1U, nhits));
  assert(device_isOuterHitOfCell_.get());
  isOuterHitOfCell_ = GPUCACell::OuterHitOfCell{device_isOuterHitOfCell_.get(), hh.offsetBPIX2()};

  auto cellStorageSize = caConstants::maxNumOfActiveDoublets * sizeof(GPUCACell::CellNeighbors) +
                         caConstants::maxNumOfActiveDoublets * sizeof(GPUCACell::CellTracks);
  // no need to use the Traits allocations, since we know this is being compiled for the CPU
  //cellStorage_ = Traits::template make_unique<unsigned char[]>(cellStorageSize, stream);
  cellStorage_ = std::make_unique<unsigned char[]>(cellStorageSize);
  device_theCellNeighborsContainer_ = (GPUCACell::CellNeighbors *)cellStorage_.get();
  device_theCellTracksContainer_ = (GPUCACell::CellTracks *)(cellStorage_.get() + caConstants::maxNumOfActiveDoublets *
                                                                                      sizeof(GPUCACell::CellNeighbors));

  gpuPixelDoublets::initDoublets(isOuterHitOfCell_,
                                 nhits,
                                 device_theCellNeighbors_.get(),
                                 device_theCellNeighborsContainer_,
                                 device_theCellTracks_.get(),
                                 device_theCellTracksContainer_);

  // no need to use the Traits allocations, since we know this is being compiled for the CPU
  //device_theCells_ = Traits::template make_unique<GPUCACell[]>(params_.maxNumberOfDoublets_, stream);
  device_theCells_ = std::make_unique<GPUCACell[]>(params_.maxNumberOfDoublets_);
  if (0 == nhits)
    return;  // protect against empty events

  // take all layer pairs into account
  auto nActualPairs = gpuPixelDoublets::nPairs;
  if (not params_.includeJumpingForwardDoublets_) {
    // exclude forward "jumping" layer pairs
    nActualPairs = gpuPixelDoublets::nPairsForTriplets;
  }
  if (params_.minHitsPerNtuplet_ > 3) {
    // for quadruplets, exclude all "jumping" layer pairs
    nActualPairs = gpuPixelDoublets::nPairsForQuadruplets;
  }

  assert(nActualPairs <= gpuPixelDoublets::nPairs);
  gpuPixelDoublets::getDoubletsFromHisto(device_theCells_.get(),
                                         device_nCells_,
                                         device_theCellNeighbors_.get(),
                                         device_theCellTracks_.get(),
                                         hh.view(),
                                         isOuterHitOfCell_,
                                         nActualPairs,
                                         params_.idealConditions_,
                                         params_.doClusterCut_,
                                         params_.doZ0Cut_,
                                         params_.doPtCut_,
                                         params_.maxNumberOfDoublets_);
}

void CAHitNtupletGeneratorKernelsCPU::classifyTuples	(	HitsOnCPU const &	hh,
		TkSoA *	tuples_d,
		cudaStream_t	cudaStream
	)

Definition at line 147 of file CAHitNtupletGeneratorKernels.cc.

References CAHitNtupletGeneratorKernels< TTraits >::counters_, cAHitNtupletGenerator::Params::cuts_, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_, CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_, cAHitNtupletGenerator::Params::doSharedHitCut_, cAHitNtupletGenerator::Params::doStats_, cAHitNtupletGenerator::Params::dupPassThrough_, TrackSoAHeterogeneousT< S >::hitIndices, CAHitNtupletGeneratorKernels< TTraits >::hitToTupleView_, iev, CAHitNtupletGeneratorKernels< TTraits >::isOuterHitOfCell_, cAHitNtupletGenerator::Params::lateFishbone_, cAHitNtupletGenerator::Params::maxNumberOfDoublets_, cAHitNtupletGenerator::Params::minHitsForSharingCut_, TrackingRecHit2DHeterogeneous< Traits >::nHits(), nhits, CAHitNtupletGeneratorKernels< TTraits >::params_, TrackSoAHeterogeneousT< S >::qualityData(), cAHitNtupletGenerator::Params::useSimpleTripletCleaner_, and TrackingRecHit2DHeterogeneous< Traits >::view().

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

                                                                                                                  {
  int32_t nhits = hh.nHits();

  auto const *tuples_d = &tracks_d->hitIndices;
  auto *quality_d = tracks_d->qualityData();

  // classify tracks based on kinematics
  kernel_classifyTracks(tuples_d, tracks_d, params_.cuts_, quality_d);

  if (params_.lateFishbone_) {
    // apply fishbone cleaning to good tracks
    kernel_fishboneCleaner(device_theCells_.get(), device_nCells_, quality_d);
  }

  // remove duplicates (tracks that share a doublet)
  kernel_fastDuplicateRemover(device_theCells_.get(), device_nCells_, tracks_d, params_.dupPassThrough_);

  // fill hit->track "map"
  if (params_.doSharedHitCut_ || params_.doStats_) {
    kernel_countHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
    cms::cuda::launchFinalize(hitToTupleView_, cudaStream);
    kernel_fillHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
  }

  // remove duplicates (tracks that share at least one hit)
  if (params_.doSharedHitCut_) {
    kernel_rejectDuplicate(
        tracks_d, quality_d, params_.minHitsForSharingCut_, params_.dupPassThrough_, device_hitToTuple_.get());

    kernel_sharedHitCleaner(hh.view(),
                            tracks_d,
                            quality_d,
                            params_.minHitsForSharingCut_,
                            params_.dupPassThrough_,
                            device_hitToTuple_.get());
    if (params_.useSimpleTripletCleaner_) {
      kernel_simpleTripletCleaner(
          tracks_d, quality_d, params_.minHitsForSharingCut_, params_.dupPassThrough_, device_hitToTuple_.get());
    } else {
      kernel_tripletCleaner(
          tracks_d, quality_d, params_.minHitsForSharingCut_, params_.dupPassThrough_, device_hitToTuple_.get());
    }
  }

  if (params_.doStats_) {
    std::lock_guard guard(lock_stat);
    kernel_checkOverflows(tuples_d,
                          device_tupleMultiplicity_.get(),
                          device_hitToTuple_.get(),
                          device_hitTuple_apc_,
                          device_theCells_.get(),
                          device_nCells_,
                          device_theCellNeighbors_.get(),
                          device_theCellTracks_.get(),
                          isOuterHitOfCell_,
                          nhits,
                          params_.maxNumberOfDoublets_,
                          counters_);
  }

  if (params_.doStats_) {
    // counters (add flag???)
    std::lock_guard guard(lock_stat);
    kernel_doStatsForHitInTracks(device_hitToTuple_.get(), counters_);
    kernel_doStatsForTracks(tuples_d, quality_d, counters_);
  }

#ifdef DUMP_GPU_TK_TUPLES
  static std::atomic<int> iev(0);
  ++iev;
  kernel_print_found_ntuplets(hh.view(), tuples_d, tracks_d, quality_d, device_hitToTuple_.get(), 100, iev);
#endif
}

template<typename TTraits >

void CAHitNtupletGeneratorKernels< TTraits >::cleanup

(

cudaStream_t

cudaStream

)

void CAHitNtupletGeneratorKernelsCPU::launchKernels	(	HitsOnCPU const &	hh,
		TkSoA *	tuples_d,
		cudaStream_t	cudaStream
	)

Definition at line 81 of file CAHitNtupletGeneratorKernels.cc.

References cms::cuda::assert(), cAHitNtupletGenerator::Params::CAThetaCutBarrel_, cAHitNtupletGenerator::Params::CAThetaCutForward_, cAHitNtupletGenerator::Params::dcaCutInnerTriplet_, cAHitNtupletGenerator::Params::dcaCutOuterTriplet_, TrackSoAHeterogeneousT< S >::detIndices, CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_, CAHitNtupletGeneratorKernels< TTraits >::device_nCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_, CAHitNtupletGeneratorKernels< TTraits >::device_theCells_, CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_, CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_, cAHitNtupletGenerator::Params::doStats_, cAHitNtupletGenerator::Params::dupPassThrough_, cAHitNtupletGenerator::Params::earlyFishbone_, cAHitNtupletGenerator::Params::hardCurvCut_, TrackSoAHeterogeneousT< S >::hitIndices, CAHitNtupletGeneratorKernels< TTraits >::isOuterHitOfCell_, cAHitNtupletGenerator::Params::lateFishbone_, cAHitNtupletGenerator::Params::minHitsPerNtuplet_, TrackingRecHit2DHeterogeneous< Traits >::nHits(), nhits, CAHitNtupletGeneratorKernels< TTraits >::params_, cAHitNtupletGenerator::Params::ptmin_, TrackSoAHeterogeneousT< S >::qualityData(), and TrackingRecHit2DHeterogeneous< Traits >::view().

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

                                                                                                                 {
  auto *tuples_d = &tracks_d->hitIndices;
  auto *detId_d = &tracks_d->detIndices;
  auto *quality_d = tracks_d->qualityData();

  assert(tuples_d && quality_d);

  // zero tuples
  cms::cuda::launchZero(tuples_d, cudaStream);

  auto nhits = hh.nHits();

  // std::cout << "N hits " << nhits << std::endl;
  // if (nhits<2) std::cout << "too few hits " << nhits << std::endl;

  //
  // applying conbinatoric cleaning such as fishbone at this stage is too expensive
  //

  kernel_connect(device_hitTuple_apc_,
                 device_hitToTuple_apc_,  // needed only to be reset, ready for next kernel
                 hh.view(),
                 device_theCells_.get(),
                 device_nCells_,
                 device_theCellNeighbors_.get(),
                 isOuterHitOfCell_,
                 params_.hardCurvCut_,
                 params_.ptmin_,
                 params_.CAThetaCutBarrel_,
                 params_.CAThetaCutForward_,
                 params_.dcaCutInnerTriplet_,
                 params_.dcaCutOuterTriplet_);

  if (nhits > 1 && params_.earlyFishbone_) {
    gpuPixelDoublets::fishbone(hh.view(), device_theCells_.get(), device_nCells_, isOuterHitOfCell_, nhits, false);
  }

  kernel_find_ntuplets(hh.view(),
                       device_theCells_.get(),
                       device_nCells_,
                       device_theCellTracks_.get(),
                       tuples_d,
                       device_hitTuple_apc_,
                       quality_d,
                       params_.minHitsPerNtuplet_);
  if (params_.doStats_)
    kernel_mark_used(device_theCells_.get(), device_nCells_);

  cms::cuda::finalizeBulk(device_hitTuple_apc_, tuples_d);

  kernel_fillHitDetIndices(tuples_d, hh.view(), detId_d);
  kernel_fillNLayers(tracks_d);

  // remove duplicates (tracks that share a doublet)
  kernel_earlyDuplicateRemover(device_theCells_.get(), device_nCells_, tracks_d, quality_d, params_.dupPassThrough_);

  kernel_countMultiplicity(tuples_d, quality_d, device_tupleMultiplicity_.get());
  cms::cuda::launchFinalize(device_tupleMultiplicity_.get(), cudaStream);
  kernel_fillMultiplicity(tuples_d, quality_d, device_tupleMultiplicity_.get());

  if (nhits > 1 && params_.lateFishbone_) {
    gpuPixelDoublets::fishbone(hh.view(), device_theCells_.get(), device_nCells_, isOuterHitOfCell_, nhits, true);
  }
}

template<typename TTraits >

uint32_t CAHitNtupletGeneratorKernels< TTraits >::nDoubletBlocks ( uint32_t  blockSize )

inlineprivate

Compute the number of doublet blocks for block size.

Definition at line 226 of file CAHitNtupletGeneratorKernels.h.

References CAHitNtupletGeneratorKernels< TTraits >::paramsMaxDoubletes3Quarters_.

                                                     {
    // We want (3 * params_.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize, but first part is pre-computed.
    return (paramsMaxDoubletes3Quarters_ + blockSize - 1) / blockSize;
  }

template<typename TTraits >

uint32_t CAHitNtupletGeneratorKernels< TTraits >::nQuadrupletBlocks ( uint32_t  blockSize )

inlineprivate

Compute the number of quadruplet blocks for block size.

Definition at line 232 of file CAHitNtupletGeneratorKernels.h.

References caConstants::maxNumberOfQuadruplets.

                                                        {
    // caConstants::maxNumberOfQuadruplets is a constexpr, so the compiler will pre compute the 3*max/4
    return (3 * caConstants::maxNumberOfQuadruplets / 4 + blockSize - 1) / blockSize;
  }

void CAHitNtupletGeneratorKernelsCPU::printCounters ( Counters const *  counters )

static

Definition at line 12 of file CAHitNtupletGeneratorKernels.cc.

Referenced by CAHitNtupletGeneratorOnGPU::~CAHitNtupletGeneratorOnGPU().

                                                                            {
  kernel_printCounters(counters);
}

template<typename TTraits >

void CAHitNtupletGeneratorKernels< TTraits >::setCounters ( Counters *  counters )

inline

Definition at line 193 of file CAHitNtupletGeneratorKernels.h.

References counters, and CAHitNtupletGeneratorKernels< TTraits >::counters_.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

{ counters_ = counters; }

template<typename TTraits >

TupleMultiplicity const* CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity ( ) const

inline

Definition at line 182 of file CAHitNtupletGeneratorKernels.h.

References CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_.

Referenced by CAHitNtupletGeneratorOnGPU::makeTuples(), and CAHitNtupletGeneratorOnGPU::makeTuplesAsync().

{ return device_tupleMultiplicity_.get(); }

Member Data Documentation

template<typename TTraits >

unique_ptr<unsigned char[]> CAHitNtupletGeneratorKernels< TTraits >::cellStorage_

private

Definition at line 199 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

template<typename TTraits >

Counters* CAHitNtupletGeneratorKernels< TTraits >::counters_ = nullptr

private

Definition at line 196 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::setCounters().

template<typename TTraits >

unique_ptr<HitToTuple> CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_

private

Definition at line 210 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), and CAHitNtupletGeneratorKernels< TTraits >::classifyTuples().

template<typename TTraits >

cms::cuda::AtomicPairCounter* CAHitNtupletGeneratorKernels< TTraits >::device_hitToTuple_apc_ = nullptr

private

Definition at line 214 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

unique_ptr<HitToTuple::Counter[]> CAHitNtupletGeneratorKernels< TTraits >::device_hitToTupleStorage_

private

Definition at line 211 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU().

template<typename TTraits >

cms::cuda::AtomicPairCounter* CAHitNtupletGeneratorKernels< TTraits >::device_hitTuple_apc_ = nullptr

private

Definition at line 216 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

unique_ptr<GPUCACell::OuterHitOfCellContainer[]> CAHitNtupletGeneratorKernels< TTraits >::device_isOuterHitOfCell_

private

Definition at line 206 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

template<typename TTraits >

uint32_t* CAHitNtupletGeneratorKernels< TTraits >::device_nCells_ = nullptr

private

Definition at line 208 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

unique_ptr<cms::cuda::AtomicPairCounter::c_type[]> CAHitNtupletGeneratorKernels< TTraits >::device_storage_

private

Definition at line 220 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU().

template<typename TTraits >

unique_ptr<caConstants::CellNeighborsVector> CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighbors_

private

Definition at line 200 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

caConstants::CellNeighbors* CAHitNtupletGeneratorKernels< TTraits >::device_theCellNeighborsContainer_

private

Definition at line 201 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

template<typename TTraits >

unique_ptr<GPUCACell[]> CAHitNtupletGeneratorKernels< TTraits >::device_theCells_

private

Definition at line 205 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

unique_ptr<caConstants::CellTracksVector> CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracks_

private

Definition at line 202 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

caConstants::CellTracks* CAHitNtupletGeneratorKernels< TTraits >::device_theCellTracksContainer_

private

Definition at line 203 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets().

template<typename TTraits >

unique_ptr<TupleMultiplicity> CAHitNtupletGeneratorKernels< TTraits >::device_tupleMultiplicity_

private

Definition at line 218 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), CAHitNtupletGeneratorKernels< TTraits >::launchKernels(), and CAHitNtupletGeneratorKernels< TTraits >::tupleMultiplicity().

template<typename TTraits >

HitToTuple::View CAHitNtupletGeneratorKernels< TTraits >::hitToTupleView_

private

Definition at line 212 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), and CAHitNtupletGeneratorKernels< TTraits >::classifyTuples().

template<typename TTraits >

GPUCACell::OuterHitOfCell CAHitNtupletGeneratorKernels< TTraits >::isOuterHitOfCell_

private

Definition at line 207 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

Params const& CAHitNtupletGeneratorKernels< TTraits >::params_

private

Definition at line 222 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::buildDoublets(), CAHitNtupletGeneratorKernels< TTraits >::classifyTuples(), and CAHitNtupletGeneratorKernels< TTraits >::launchKernels().

template<typename TTraits >

const uint32_t CAHitNtupletGeneratorKernels< TTraits >::paramsMaxDoubletes3Quarters_

private

Intermediate result avoiding repeated computations.

Definition at line 224 of file CAHitNtupletGeneratorKernels.h.

Referenced by CAHitNtupletGeneratorKernels< TTraits >::nDoubletBlocks().