CAHitNtupletGeneratorKernelsImpl.h File Reference

#include <cmath>
#include <cstdint>
#include <limits>
#include <cuda_runtime.h>
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
#include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitsUtilities.h"
#include "CAStructures.h"
#include "CAHitNtupletGeneratorKernels.h"
#include "GPUCACell.h"
#include "gpuFishbone.h"
#include "gpuPixelDoublets.h"

Go to the source code of this file.

Namespaces
	caHitNtupletGeneratorKernels

Typedefs
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::CAParams = caHitNtupletGenerator::CAParamsT< TrackerTraits >
using	caHitNtupletGeneratorKernels::Cell = GPUCACellT< TrackerTraits >
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::CellNeighborsVector = caStructures::CellNeighborsVectorT< TrackerTraits >
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::CellTracksVector = caStructures::CellTracksVectorT< TrackerTraits >
using	caHitNtupletGeneratorKernels::Counters = caHitNtupletGenerator::Counters
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::HitContainer = typename TrackSoA< TrackerTraits >::HitContainer
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::HitsConstView = typename GPUCACellT< TrackerTraits >::HitsConstView
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::HitToTuple = caStructures::HitToTupleT< TrackerTraits >
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::OuterHitOfCell = caStructures::OuterHitOfCellT< TrackerTraits >
using	caHitNtupletGeneratorKernels::Quality = pixelTrack::Quality
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::QualityCuts = pixelTrack::QualityCutsT< TrackerTraits >
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::TkSoAView = TrackSoAView< TrackerTraits >
template<typename TrackerTraits >
using	caHitNtupletGeneratorKernels::TupleMultiplicity = caStructures::TupleMultiplicityT< TrackerTraits >

Functions
template<typename TrackerTraits >
	caHitNtupletGeneratorKernels::__attribute__ ((always_inline)) void kernel_classifyTracks(TkSoAView< TrackerTraits > tracks_view
	TODO : why there was quality here? More...
	caHitNtupletGeneratorKernels::assert (nCells)
	cannot be loose More...
	caHitNtupletGeneratorKernels::for (int idx=first, nt=(*nCells);idx< nt;idx+=gridDim.x *blockDim.x)
	caHitNtupletGeneratorKernels::if (0==first) = ntracks

Variables
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter *	caHitNtupletGeneratorKernels::apc
auto &	caHitNtupletGeneratorKernels::c = *counters
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const *	caHitNtupletGeneratorKernels::cellNeighbors
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__	caHitNtupletGeneratorKernels::cells
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const *	caHitNtupletGeneratorKernels::cellTracks
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const int32_t uint32_t Counters *	caHitNtupletGeneratorKernels::counters
QualityCuts< TrackerTraits >	caHitNtupletGeneratorKernels::cuts
uint32_t const *__restrict__ TkSoAView< TrackerTraits > bool	caHitNtupletGeneratorKernels::dupPassThrough
auto	caHitNtupletGeneratorKernels::first = threadIdx.x + blockIdx.x * blockDim.x
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t	caHitNtupletGeneratorKernels::firstPrint
auto const &	caHitNtupletGeneratorKernels::foundNtuplets = *ptuples
auto const	caHitNtupletGeneratorKernels::good = pixelTrack::Quality::strict
	min quality of good More...
HitsConstView< TrackerTraits >	caHitNtupletGeneratorKernels::hh
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const *	caHitNtupletGeneratorKernels::hitToTuple
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t int32_t int	caHitNtupletGeneratorKernels::iev
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const	caHitNtupletGeneratorKernels::isOuterHitOfCell
int	caHitNtupletGeneratorKernels::l1end = hh.hitsLayerStart()[1]
TkSoAView< TrackerTraits > HitToTuple< TrackerTraits > const *__restrict__ int32_t int32_t	caHitNtupletGeneratorKernels::lastPrint
auto const	caHitNtupletGeneratorKernels::longTqual = pixelTrack::Quality::highPurity
constexpr auto	caHitNtupletGeneratorKernels::loose = pixelTrack::Quality::loose
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const int32_t uint32_t	caHitNtupletGeneratorKernels::maxNumberOfDoublets
constexpr float	caHitNtupletGeneratorKernels::maxScore = std::numeric_limits<float>::max()
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__	caHitNtupletGeneratorKernels::nCells
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter GPUCACellT< TrackerTraits > const *__restrict__ uint32_t const *__restrict__ CellNeighborsVector< TrackerTraits > const CellTracksVector< TrackerTraits > const OuterHitOfCell< TrackerTraits > const int32_t	caHitNtupletGeneratorKernels::nHits
auto	caHitNtupletGeneratorKernels::nhits = hh.nHits()
uint16_t	caHitNtupletGeneratorKernels::nmin
constexpr float	caHitNtupletGeneratorKernels::nSigma2 = 25.f
auto	caHitNtupletGeneratorKernels::ntracks = std::min<int>(apc->get().m, tracks_view.metadata().size() - 1)
TkSoAView< TrackerTraits > GPUCACellT< TrackerTraits > *__restrict__ uint32_t const CellTracksVector< TrackerTraits > cms::cuda::AtomicPairCounter CAParams< TrackerTraits >	caHitNtupletGeneratorKernels::params
HitContainer< TrackerTraits > const *__restrict__ Quality const *__restrict__ HitToTuple< TrackerTraits > const *__restrict__	caHitNtupletGeneratorKernels::phitToTuple
HitContainer< TrackerTraits > const *__restrict__	caHitNtupletGeneratorKernels::ptuples
HitContainer< TrackerTraits > const *__restrict__ Quality const *__restrict__	caHitNtupletGeneratorKernels::quality
auto const	caHitNtupletGeneratorKernels::reject = dupPassThrough ? loose : dup
constexpr uint32_t	caHitNtupletGeneratorKernels::tkNotFound = std::numeric_limits<uint16_t>::max()
uint32_t const *__restrict__ TkSoAView< TrackerTraits >	caHitNtupletGeneratorKernels::tracks_view
TupleMultiplicity< TrackerTraits > const *	caHitNtupletGeneratorKernels::tupleMultiplicity
HitContainer< TrackerTraits > const *__restrict__	caHitNtupletGeneratorKernels::tuples