CAHitNtupletGeneratorKernelsCPU< TrackerTraits > Class Template Reference

#include <CAHitNtupletGeneratorKernels.h>

Inheritance diagram for CAHitNtupletGeneratorKernelsCPU< TrackerTraits >:

Public Member Functions
void	allocateOnGPU (int32_t nHits, cudaStream_t stream)
void	buildDoublets (const HitsConstView &hh, int32_t offsetBPIX2, cudaStream_t stream)
void	classifyTuples (const HitsConstView &hh, TkSoAView &track_view, cudaStream_t cudaStream)
void	launchKernels (const HitsConstView &hh, TkSoAView &track_view, cudaStream_t cudaStream)
Public Member Functions inherited from CAHitNtupletGeneratorKernels< cms::cudacompat::CPUTraits, TrackerTraits >
void	allocateOnGPU (int32_t nHits, cudaStream_t stream)
void	buildDoublets (const HitsConstView &hh, cudaStream_t stream)
	CAHitNtupletGeneratorKernels (Params const &params)
void	classifyTuples (const HitsConstView &hh, TkSoAView &track_view, cudaStream_t cudaStream)
void	cleanup (cudaStream_t cudaStream)
void	launchKernels (const HitsConstView &hh, TkSoAView &track_view, cudaStream_t cudaStream)
void	setCounters (Counters *counters)
TupleMultiplicity const *	tupleMultiplicity () const
	~CAHitNtupletGeneratorKernels ()=default

Static Public Member Functions
static void	printCounters (Counters const *counters)
Static Public Member Functions inherited from CAHitNtupletGeneratorKernels< cms::cudacompat::CPUTraits, TrackerTraits >
static void	printCounters (Counters const *counters)

Private Types
using	CAParams = caHitNtupletGenerator::CAParamsT< TrackerTraits >
using	CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
using	CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
using	Counters = caHitNtupletGenerator::Counters
using	HitContainer = typename TrackSoA< TrackerTraits >::HitContainer
using	HitsConstView = TrackingRecHitSoAConstView< TrackerTraits >
using	HitToTuple = caStructures::HitToTupleT< TrackerTraits >
using	Params = caHitNtupletGenerator::ParamsT< TrackerTraits >
using	TkSoAView = TrackSoAView< TrackerTraits >
using	TupleMultiplicity = caStructures::TupleMultiplicityT< TrackerTraits >

Additional Inherited Members
Public Types inherited from CAHitNtupletGeneratorKernels< cms::cudacompat::CPUTraits, TrackerTraits >
using	CACell = GPUCACellT< TrackerTraits >
using	CAParams = caHitNtupletGenerator::CAParamsT< TrackerTraits >
using	CellCuts = gpuPixelDoublets::CellCutsT< TrackerTraits >
using	CellNeighbors = caStructures::CellNeighborsT< TrackerTraits >
using	CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
using	CellTracks = caStructures::CellTracksT< TrackerTraits >
using	CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
using	Counters = caHitNtupletGenerator::Counters
using	HitContainer = typename TrackSoA< TrackerTraits >::HitContainer
using	HitsConstView = TrackingRecHitSoAConstView< TrackerTraits >
using	HitsView = TrackingRecHitSoAView< TrackerTraits >
using	HitToTuple = caStructures::HitToTupleT< TrackerTraits >
using	OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
using	OuterHitOfCellContainer = caStructures::OuterHitOfCellContainerT< TrackerTraits >
using	Params = caHitNtupletGenerator::ParamsT< TrackerTraits >
using	Quality = pixelTrack::Quality
using	QualityCuts = pixelTrack::QualityCutsT< TrackerTraits >
using	TkSoAView = TrackSoAView< TrackerTraits >
using	TrackerTraits = TrackerTraits
using	Traits = cms::cudacompat::CPUTraits
using	TupleMultiplicity = caStructures::TupleMultiplicityT< TrackerTraits >
using	unique_ptr = typename Traits::template unique_ptr< T >
Protected Member Functions inherited from CAHitNtupletGeneratorKernels< cms::cudacompat::CPUTraits, TrackerTraits >
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...
Protected Attributes inherited from CAHitNtupletGeneratorKernels< cms::cudacompat::CPUTraits, TrackerTraits >
unique_ptr< unsigned char[]>	cellStorage_
Counters *	counters_
unique_ptr< CellCuts >	device_cellCuts_
unique_ptr< HitToTuple >	device_hitToTuple_
cms::cuda::AtomicPairCounter *	device_hitToTuple_apc_
unique_ptr< uint32_t[]>	device_hitToTupleStorage_
cms::cuda::AtomicPairCounter *	device_hitTuple_apc_
unique_ptr< OuterHitOfCellContainer[]>	device_isOuterHitOfCell_
uint32_t *	device_nCells_
unique_ptr< cms::cuda::AtomicPairCounter::c_type[]>	device_storage_
unique_ptr< CellNeighborsVector >	device_theCellNeighbors_
CellNeighbors *	device_theCellNeighborsContainer_
unique_ptr< CACell[]>	device_theCells_
unique_ptr< CellTracksVector >	device_theCellTracks_
CellTracks *	device_theCellTracksContainer_
unique_ptr< TupleMultiplicity >	device_tupleMultiplicity_
HitToTuple::View	hitToTupleView_
OuterHitOfCell	isOuterHitOfCell_
Params	params_
const uint32_t	paramsMaxDoubletes3Quarters_
	Intermediate result avoiding repeated computations. More...

Detailed Description

template<typename TrackerTraits>
class CAHitNtupletGeneratorKernelsCPU< TrackerTraits >

Definition at line 321 of file CAHitNtupletGeneratorKernels.h.

Member Typedef Documentation

CAParams

template<typename TrackerTraits >

using CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::CAParams = caHitNtupletGenerator::CAParamsT<TrackerTraits>

private

Definition at line 325 of file CAHitNtupletGeneratorKernels.h.

CellNeighborsVector

template<typename TrackerTraits >

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

private

Definition at line 329 of file CAHitNtupletGeneratorKernels.h.

CellTracksVector

template<typename TrackerTraits >

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

private

Definition at line 331 of file CAHitNtupletGeneratorKernels.h.

Counters

template<typename TrackerTraits >

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

private

Definition at line 324 of file CAHitNtupletGeneratorKernels.h.

HitContainer

template<typename TrackerTraits >

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

private

Definition at line 327 of file CAHitNtupletGeneratorKernels.h.

HitsConstView

template<typename TrackerTraits >

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

private

Definition at line 334 of file CAHitNtupletGeneratorKernels.h.

HitToTuple

template<typename TrackerTraits >

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

private

Definition at line 330 of file CAHitNtupletGeneratorKernels.h.

Params

template<typename TrackerTraits >

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

private

Definition at line 337 of file CAHitNtupletGeneratorKernels.h.

TkSoAView

template<typename TrackerTraits >

using CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::TkSoAView = TrackSoAView<TrackerTraits>

private

Definition at line 335 of file CAHitNtupletGeneratorKernels.h.

TupleMultiplicity

template<typename TrackerTraits >

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

private

Definition at line 332 of file CAHitNtupletGeneratorKernels.h.

Member Function Documentation

allocateOnGPU()

template<typename TrackerTraits >

void CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::allocateOnGPU	(	int32_t	nHits,
		cudaStream_t	stream
	)

Definition at line 12 of file CAHitNtupletGeneratorKernelsAlloc.cc.

References cms::cuda::assert(), ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), gather_cfg::cout, cudaCheck, nHits, cms::cuda::stream, svgfig::template(), and relativeConstraints::value.

                                                                                                     {
  using Traits = cms::cudacompat::CPUTraits;
#endif

  using CellCuts = gpuPixelDoublets::CellCutsT<TrackerTraits>;

  // ALLOCATIONS FOR THE INTERMEDIATE RESULTS (STAYS ON WORKER)

  this->device_theCellNeighbors_ = Traits::template make_unique<CellNeighborsVector>(stream);
  this->device_theCellTracks_ = Traits::template make_unique<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);
  this->device_hitToTuple_ = Traits::template make_unique<HitToTuple>(stream);
  this->device_hitToTupleStorage_ = Traits::template make_unique<typename HitToTuple::Counter[]>(nHits, stream);
  this->hitToTupleView_.assoc = this->device_hitToTuple_.get();
  this->hitToTupleView_.offStorage = this->device_hitToTupleStorage_.get();
  this->hitToTupleView_.offSize = nHits;

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

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

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

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

buildDoublets()

template<typename TrackerTraits >

void CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::buildDoublets	(	const HitsConstView &	hh,
		int32_t	offsetBPIX2,
		cudaStream_t	stream
	)

Definition at line 17 of file CAHitNtupletGeneratorKernels.cc.

References gpuPixelDoublets::assert(), gather_cfg::cout, gpuPixelDoublets::hh, SiStripPI::max, gpuPixelDoublets::nActualPairs, TrackingDataMCValidation_Standalone_cff::nhits, phase1PixelTopology::nPairs, and mitigatedMETSequence_cff::U.

                                                                                        {
  using namespace gpuPixelDoublets;

  using GPUCACell = GPUCACellT<TrackerTraits>;
  using OuterHitOfCell = typename GPUCACell::OuterHitOfCell;
  using CellNeighbors = typename GPUCACell::CellNeighbors;
  using CellTracks = typename GPUCACell::CellTracks;
  using OuterHitOfCellContainer = typename GPUCACell::OuterHitOfCellContainer;

  auto nhits = hh.nHits();

#ifdef NTUPLE_DEBUG
  std::cout << "building Doublets out of " << nhits << " Hits. BPIX2 offset is " << 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
  //this->device_isOuterHitOfCell_ = Traits::template make_unique<GPUCACell::OuterHitOfCell[]>(std::max(1U, nhits), stream);
  this->device_isOuterHitOfCell_ = std::make_unique<OuterHitOfCellContainer[]>(std::max(1U, nhits));
  assert(this->device_isOuterHitOfCell_.get());
  this->isOuterHitOfCell_ = OuterHitOfCell{this->device_isOuterHitOfCell_.get(), offsetBPIX2};

  auto cellStorageSize = TrackerTraits::maxNumOfActiveDoublets * sizeof(CellNeighbors) +
                         TrackerTraits::maxNumOfActiveDoublets * sizeof(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);
  this->cellStorage_ = std::make_unique<unsigned char[]>(cellStorageSize);
  this->device_theCellNeighborsContainer_ = (CellNeighbors *)this->cellStorage_.get();
  this->device_theCellTracksContainer_ =
      (CellTracks *)(this->cellStorage_.get() + TrackerTraits::maxNumOfActiveDoublets * sizeof(CellNeighbors));

  initDoublets<TrackerTraits>(this->isOuterHitOfCell_,
                              nhits,
                              this->device_theCellNeighbors_.get(),
                              this->device_theCellNeighborsContainer_,
                              this->device_theCellTracks_.get(),
                              this->device_theCellTracksContainer_);

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

  // take all layer pairs into account
  auto nActualPairs = this->params_.nPairs();

  assert(nActualPairs <= TrackerTraits::nPairs);

  getDoubletsFromHisto<TrackerTraits>(this->device_theCells_.get(),
                                      this->device_nCells_,
                                      this->device_theCellNeighbors_.get(),
                                      this->device_theCellTracks_.get(),
                                      hh,
                                      this->isOuterHitOfCell_,
                                      nActualPairs,
                                      this->params_.caParams_.maxNumberOfDoublets_,
                                      this->device_cellCuts_.get());
}

classifyTuples()

template<typename TrackerTraits >

void CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::classifyTuples	(	const HitsConstView &	hh,
		TkSoAView &	track_view,
		cudaStream_t	cudaStream
	)

Definition at line 144 of file CAHitNtupletGeneratorKernels.cc.

References caHitNtupletGeneratorKernels::hh, caHitNtupletGeneratorKernels::iev, CommonMethods::lock(), mutex, caHitNtupletGeneratorKernels::nhits, and caHitNtupletGeneratorKernels::tracks_view.

                                                                                             {
  using namespace caHitNtupletGeneratorKernels;

  int32_t nhits = hh.metadata().size();

  // classify tracks based on kinematics
  kernel_classifyTracks<TrackerTraits>(tracks_view, this->params_.qualityCuts_);
  if (this->params_.lateFishbone_) {
    // apply fishbone cleaning to good tracks
    kernel_fishboneCleaner<TrackerTraits>(this->device_theCells_.get(), this->device_nCells_, tracks_view);
  }

  // remove duplicates (tracks that share a doublet)
  kernel_fastDuplicateRemover<TrackerTraits>(
      this->device_theCells_.get(), this->device_nCells_, tracks_view, this->params_.dupPassThrough_);

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

  // remove duplicates (tracks that share at least one hit)
  if (this->params_.doSharedHitCut_) {
    kernel_rejectDuplicate<TrackerTraits>(tracks_view,
                                          this->params_.minHitsForSharingCut_,
                                          this->params_.dupPassThrough_,
                                          this->device_hitToTuple_.get());

    kernel_sharedHitCleaner<TrackerTraits>(hh,
                                           tracks_view,
                                           this->params_.minHitsForSharingCut_,
                                           this->params_.dupPassThrough_,
                                           this->device_hitToTuple_.get());
    if (this->params_.useSimpleTripletCleaner_) {
      kernel_simpleTripletCleaner<TrackerTraits>(tracks_view,
                                                 this->params_.minHitsForSharingCut_,
                                                 this->params_.dupPassThrough_,
                                                 this->device_hitToTuple_.get());
    } else {
      kernel_tripletCleaner<TrackerTraits>(tracks_view,
                                           this->params_.minHitsForSharingCut_,
                                           this->params_.dupPassThrough_,
                                           this->device_hitToTuple_.get());
    }
  }

  if (this->params_.doStats_) {
    std::lock_guard guard(lock_stat);
    kernel_checkOverflows<TrackerTraits>(tracks_view,
                                         this->device_tupleMultiplicity_.get(),
                                         this->device_hitToTuple_.get(),
                                         this->device_hitTuple_apc_,
                                         this->device_theCells_.get(),
                                         this->device_nCells_,
                                         this->device_theCellNeighbors_.get(),
                                         this->device_theCellTracks_.get(),
                                         this->isOuterHitOfCell_,
                                         nhits,
                                         this->params_.caParams_.maxNumberOfDoublets_,
                                         this->counters_);
  }

  if (this->params_.doStats_) {
    // counters (add flag???)
    std::lock_guard guard(lock_stat);
    kernel_doStatsForHitInTracks<TrackerTraits>(this->device_hitToTuple_.get(), this->counters_);
    kernel_doStatsForTracks<TrackerTraits>(tracks_view, this->counters_);
  }

#ifdef DUMP_GPU_TK_TUPLES
  static std::atomic<int> iev(0);
  static std::mutex lock;
  {
    std::lock_guard<std::mutex> guard(lock);
    ++iev;
    kernel_print_found_ntuplets<TrackerTraits>(hh, tracks_view, this->device_hitToTuple_.get(), 0, 1000000, iev);
  }
#endif
}

launchKernels()

template<typename TrackerTraits >

void CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::launchKernels	(	const HitsConstView &	hh,
		TkSoAView &	track_view,
		cudaStream_t	cudaStream
	)

Definition at line 80 of file CAHitNtupletGeneratorKernels.cc.

References gather_cfg::cout, caHitNtupletGeneratorKernels::hh, caHitNtupletGeneratorKernels::nhits, and caHitNtupletGeneratorKernels::tracks_view.

                                                                                            {
  using namespace caHitNtupletGeneratorKernels;

  // zero tuples
  cms::cuda::launchZero(&tracks_view.hitIndices(), cudaStream);

  uint32_t nhits = hh.metadata().size();

#ifdef NTUPLE_DEBUG
  std::cout << "start tuple building. N hits " << nhits << std::endl;
  if (nhits < 2)
    std::cout << "too few hits " << nhits << std::endl;
#endif

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

  kernel_connect<TrackerTraits>(this->device_hitTuple_apc_,
                                this->device_hitToTuple_apc_,  // needed only to be reset, ready for next kernel
                                hh,
                                this->device_theCells_.get(),
                                this->device_nCells_,
                                this->device_theCellNeighbors_.get(),
                                this->isOuterHitOfCell_,
                                this->params_.caParams_);

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

  kernel_find_ntuplets<TrackerTraits>(hh,
                                      tracks_view,
                                      this->device_theCells_.get(),
                                      this->device_nCells_,
                                      this->device_theCellTracks_.get(),
                                      this->device_hitTuple_apc_,
                                      this->params_.caParams_);
  if (this->params_.doStats_)
    kernel_mark_used(this->device_theCells_.get(), this->device_nCells_);

  cms::cuda::finalizeBulk(this->device_hitTuple_apc_, &tracks_view.hitIndices());

  kernel_fillHitDetIndices<TrackerTraits>(tracks_view, hh);
  kernel_fillNLayers<TrackerTraits>(tracks_view, this->device_hitTuple_apc_);

  // remove duplicates (tracks that share a doublet)
  kernel_earlyDuplicateRemover<TrackerTraits>(
      this->device_theCells_.get(), this->device_nCells_, tracks_view, this->params_.dupPassThrough_);

  kernel_countMultiplicity<TrackerTraits>(tracks_view, this->device_tupleMultiplicity_.get());
  cms::cuda::launchFinalize(this->device_tupleMultiplicity_.get(), cudaStream);
  kernel_fillMultiplicity<TrackerTraits>(tracks_view, this->device_tupleMultiplicity_.get());

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

printCounters()

template<typename TrackerTraits >

void CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::printCounters ( Counters const *  counters )

static

Definition at line 12 of file CAHitNtupletGeneratorKernels.cc.

References caHitNtupletGeneratorKernels::counters.

Referenced by CAHitNtupletGeneratorOnGPU< TrackerTraits >::endJob().

                                                                                           {
  caHitNtupletGeneratorKernels::kernel_printCounters(counters);
}