RiemannFitOnGPU.h File Reference

#include <cstdint>
#include <cuda_runtime.h>
#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitsUtilities.h"
#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
#include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
#include "RecoTracker/PixelTrackFitting/interface/RiemannFit.h"
#include "HelixFitOnGPU.h"

Go to the source code of this file.

Typedefs
template<typename TrackerTraits >
using	OutputSoAView = TrackSoAView< TrackerTraits >
template<typename TrackerTraits >
using	TupleMultiplicity = caStructures::TupleMultiplicityT< TrackerTraits >
template<typename TrackerTraits >
using	Tuples = typename TrackSoA< TrackerTraits >::HitContainer

Functions
template<int N, typename TrackerTraits >
	__attribute__ ((always_inline)) void kernel_FastFit(Tuples< TrackerTraits > const *__restrict__ foundNtuplets
	assert (hitsInFit<=nHits)
	assert (pfast_fit)
	assert (foundNtuplets)
	assert (tupleMultiplicity)
	assert (N<=nHits)
	for (int local_idx=local_start, nt=riemannFit::maxNumberOfConcurrentFits;local_idx< nt;local_idx+=gridDim.x *blockDim.x)

Variables
uint32_t double	bField
uint32_t double double *__restrict__ float *__restrict__ double *__restrict__ riemannFit::CircleFit *	circle_fit
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t TrackingRecHitSoAConstView< TrackerTraits >	hh
auto	local_start = blockIdx.x * blockDim.x + threadIdx.x
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t	nHits
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t TrackingRecHitSoAConstView< TrackerTraits > double *__restrict__ float *__restrict__ double *__restrict__ uint32_t	offset
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t TrackingRecHitSoAConstView< TrackerTraits > double *__restrict__ float *__restrict__ double *__restrict__	pfast_fit
uint32_t double double *__restrict__ float *__restrict__ double *__restrict__	pfast_fit_input
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t TrackingRecHitSoAConstView< TrackerTraits > double *__restrict__	phits
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t TrackingRecHitSoAConstView< TrackerTraits > double *__restrict__ float *__restrict__	phits_ge
uint32_t double OutputSoAView< TrackerTraits >	results_view
TupleMultiplicity< TrackerTraits > const *__restrict__	tupleMultiplicity

Typedef Documentation

OutputSoAView

template<typename TrackerTraits >

using OutputSoAView = TrackSoAView<TrackerTraits>

Definition at line 21 of file RiemannFitOnGPU.h.

TupleMultiplicity

template<typename TrackerTraits >

using TupleMultiplicity = caStructures::TupleMultiplicityT<TrackerTraits>

Definition at line 23 of file RiemannFitOnGPU.h.

Tuples

template<typename TrackerTraits >

using Tuples = typename TrackSoA<TrackerTraits>::HitContainer

Definition at line 19 of file RiemannFitOnGPU.h.

Function Documentation

__attribute__()

template<int N, typename TrackerTraits >

__attribute__ ( (always_inline)  ) const

inline

assert() [1/5]

assert

(

hitsInFit<=

nHits

)

Referenced by for().

assert() [2/5]

assert

(

pfast_fit

)

assert() [3/5]

assert

(

foundNtuplets

)

assert() [4/5]

assert

(

tupleMultiplicity

)

assert() [5/5]

assert

(

N<=

nHits

)

for()

for

(

int

local_idx = local_start

)

Definition at line 50 of file RiemannFitOnGPU.h.

References assert(), ALPAKA_ACCELERATOR_NAMESPACE::brokenline::fast_fit, riemannFit::fastFit(), caHitNtupletGeneratorKernels::foundNtuplets, hh, hfClusterShapes_cfi::hits, mps_fire::i, nHits, offset, pfast_fit, phits, phits_ge, and tupleMultiplicity.

                                            {
    auto tuple_idx = local_idx + offset;
    if (tuple_idx >= tupleMultiplicity->size(nHits))
      break;

    // get it from the ntuple container (one to one to helix)
    auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
    assert(int(tkid) < foundNtuplets->nOnes());

    assert(foundNtuplets->size(tkid) == nHits);

    riemannFit::Map3xNd<N> hits(phits + local_idx);
    riemannFit::Map4d fast_fit(pfast_fit + local_idx);
    riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);

    // Prepare data structure
    auto const *hitId = foundNtuplets->begin(tkid);
    for (unsigned int i = 0; i < hitsInFit; ++i) {
      auto hit = hitId[i];
      float ge[6];
      hh.cpeParams().detParams(hh[hit].detectorIndex()).frame.toGlobal(hh[hit].xerrLocal(), 0, hh[hit].yerrLocal(), ge);

      hits.col(i) << hh[hit].xGlobal(), hh[hit].yGlobal(), hh[hit].zGlobal();
      hits_ge.col(i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
    }
    riemannFit::fastFit(hits, fast_fit);

    // no NaN here....
    assert(fast_fit(0) == fast_fit(0));
    assert(fast_fit(1) == fast_fit(1));
    assert(fast_fit(2) == fast_fit(2));
    assert(fast_fit(3) == fast_fit(3));
  }

Variable Documentation

bField

uint32_t double bField

Definition at line 88 of file RiemannFitOnGPU.h.

circle_fit

uint32_t double OutputSoAView< TrackerTraits > double *__restrict__ float *__restrict__ double *__restrict__ riemannFit::CircleFit *__restrict__ circle_fit

Definition at line 88 of file RiemannFitOnGPU.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::Kernel_CircleFit< N, TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::Kernel_LineFit< N, TrackerTraits >::operator()().

hh

TupleMultiplicity<TrackerTraits> const* __restrict__ uint32_t TrackingRecHitSoAConstView<TrackerTraits> hh

Definition at line 27 of file RiemannFitOnGPU.h.

Referenced by for().

local_start

auto local_start = blockIdx.x * blockDim.x + threadIdx.x

Definition at line 43 of file RiemannFitOnGPU.h.

nHits

uint32_t nHits

Definition at line 27 of file RiemannFitOnGPU.h.

Referenced by DTMeantimerPatternReco::addHits(), SeedFinder::addHitSelector(), CAHitNtupletGeneratorKernelsCPU< TrackerTraits >::allocateOnGPU(), DTSegmentsTask::analyze(), DTOccupancyEfficiency::analyze(), SiPixelMonitorTrackSoA< T >::analyze(), SiPixelMonitorTrackSoAAlpaka< T >::analyze(), DTSegmentAnalysisTask::analyze(), BtlLocalRecoValidation::analyze(), SiPixelCompareTrackSoAAlpaka< T >::analyze(), SiPixelCompareTrackSoA< T >::analyze(), HcalHitValidation::analyzeLayer(), SiPixelCompareTracks< T >::analyzeSeparate(), RecoDisplacedMuonValidator::bookHistograms(), RecoMuonValidator::bookHistograms(), RecoDisplacedMuonValidator::MuonME::bookHistos(), RecoMuonValidator::MuonME::bookHistos(), FWTracksterLayersProxyBuilder::build(), PtAssignmentEngine2017::calculate_address(), PtAssignmentEngine2017::calculate_pt_xml(), RPCSeedOverlapper::CheckOverlap(), mkfit::Event::clean_cms_seedtracks(), mkfit::StdSeq::clean_cms_seedtracks_iter(), OMTFResult::empty(), SimG4HcalValidation::fetchHits(), DTSegmentAnalysisTask::fillHistos(), L2TauNNProducerAlpaka::fillPatatracks(), L2TauNNProducer::fillPatatracks(), TrackingNtuple::fillSeeds(), CastorShowerLibraryMaker::FillShowerEvent(), MuGEMMuonExtTableProducer::fillTable(), MuDTSegmentExtTableProducer::fillTable(), TrackingNtuple::fillTracks(), gpuVertexFinder::for(), for(), DTChamberEfficiencyTask::getBestSegment(), MuonSeedCleaner::GroupSeeds(), MillePedeMonitor::init(), DTChamberEfficiencyTask::isGoodSegment(), pixelTrack::QualityCutsT< TrackerTraits, pixelTopology::isPhase1Topology< TrackerTraits > >::isHP(), PFDisplacedVertexHelper::isTrackSelected(), main(), ALPAKA_ACCELERATOR_NAMESPACE::pixelgpudetails::PixelRecHitKernel< TrackerTraits >::makeHitsAsync(), SCEnergyCorrectorDRN::makeInput(), BTLUncalibRecHitAlgo::makeRecHit(), LowPtConversion::match(), CovarianceParameterization::meanValue(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::InitDoublets< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_BLFastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_FastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::LoadTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_CircleFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_LineFit< N, TrackerTraits >::operator()(), CovarianceParameterization::pack(), SETSeedFinder::pre_prune(), CtfSpecialSeedGenerator::preliminaryCheck(), PixelTemplateSmearerBase::process(), SeedGeneratorFromProtoTracksEDProducer::produce(), QualityFilter::produce(), TrackListCombiner::produce(), PixelTracksProducer::produce(), HiBadParticleCleaner::produce(), TSGFromL1Muon::produce(), PixelTrackSoAFromCUDAT< TrackerTraits >::produce(), SiPixelRecHitFromSoAAlpaka< TrackerTraits >::produce(), SeedProducerFromSoAT< TrackerTraits >::produce(), PixelTrackProducerFromSoAAlpaka< TrackerTraits >::produce(), PixelTrackProducerFromSoAT< TrackerTraits >::produce(), PixelVertexProducerCUDAT< TrackerTraits >::produceOnCPU(), CSCSegAlgoShowering::pruneFromResidual(), CSCSegAlgoDF::pruneFromResidual(), mkfit::StdSeq::qfilter_pixelLessBkwd(), mkfit::StdSeq::qfilter_pixelLessFwd(), SeedFromGenericPairOrTriplet::qualityFilter(), PFAlgo::recoTracksNotHCAL(), ReferenceTrajectoryBase::ReferenceTrajectoryBase(), mkfit::Track::reserveHits(), reco::TrackResiduals::resize(), mkfit::Track::resizeHits(), PixelTrackReconstruction::run(), ConversionTrackPairFinder::run(), CSCSegAlgoDF::run(), EEBadScFilter::scan5x5(), MuonSeedCleaner::SeedCandidates(), SeedFromProtoTrack::SeedFromProtoTrack(), DTCombinatorialPatternReco4D::segmentSpecialZed(), L2TauNNProducerAlpaka::selectGoodTracksAndVertices(), L2TauNNProducer::selectGoodTracksAndVertices(), pat::PackedCandidate::setTrackPropertiesLite(), CSCBaseElectronicsSim::simulate(), storeTracks(), mkfit::Track::Track(), hi::EPCuts::trackQuality_Pixel(), and CovarianceParameterization::unpack().

offset

uint32_t double OutputSoAView<TrackerTraits> double* __restrict__ float* __restrict__ double* __restrict__ riemannFit::CircleFit* __restrict__ uint32_t offset

Initial value:

{
  constexpr uint32_t hitsInFit = N

Definition at line 33 of file RiemannFitOnGPU.h.

Referenced by for().

pfast_fit

TupleMultiplicity<TrackerTraits> const* __restrict__ uint32_t TrackingRecHitSoAConstView<TrackerTraits> double* __restrict__ float* __restrict__ double* __restrict__ pfast_fit

Definition at line 27 of file RiemannFitOnGPU.h.

Referenced by for().

pfast_fit_input

uint32_t double OutputSoAView< TrackerTraits > double *__restrict__ float *__restrict__ double *__restrict__ pfast_fit_input

Definition at line 88 of file RiemannFitOnGPU.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::Kernel_CircleFit< N, TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::Kernel_LineFit< N, TrackerTraits >::operator()().

phits

uint32_t double OutputSoAView<TrackerTraits> double* __restrict__ phits

Definition at line 27 of file RiemannFitOnGPU.h.

Referenced by for().

phits_ge

uint32_t double OutputSoAView<TrackerTraits> double* __restrict__ float* __restrict__ phits_ge

Definition at line 27 of file RiemannFitOnGPU.h.

Referenced by for().

results_view

uint32_t double OutputSoAView<TrackerTraits> results_view

Definition at line 129 of file RiemannFitOnGPU.h.

tupleMultiplicity

TupleMultiplicity<TrackerTraits> const* __restrict__ tupleMultiplicity

Definition at line 27 of file RiemannFitOnGPU.h.

Referenced by for().