BrokenLineFitOnGPU.h File Reference

#include <cstdint>
#include <cuda_runtime.h>
#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
#include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
#include "RecoPixelVertexing/PixelTrackFitting/interface/BrokenLine.h"
#include "HelixFitOnGPU.h"

Go to the source code of this file.

Typedefs
using	HitsOnGPU = TrackingRecHit2DSOAView
using	OutputSoA = pixelTrack::TrackSoA
using	Tuples = pixelTrack::HitContainer

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

Variables
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__	hhp
auto	local_start = blockIdx.x * blockDim.x + threadIdx.x
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t	nHits
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t uint32_t	offset
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__	pfast_fit
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__ double *__restrict__	phits
const caConstants::TupleMultiplicity *__restrict__ const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__	phits_ge
const caConstants::TupleMultiplicity *__restrict__	tupleMultiplicity

Typedef Documentation

HitsOnGPU

using HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 19 of file BrokenLineFitOnGPU.h.

OutputSoA

using OutputSoA = pixelTrack::TrackSoA

Definition at line 21 of file BrokenLineFitOnGPU.h.

Tuples

using Tuples = pixelTrack::HitContainer

Definition at line 20 of file BrokenLineFitOnGPU.h.

Function Documentation

__attribute__()

template<int N>

__attribute__ ( (always_inline)  ) const

inline

Definition at line 148 of file BrokenLineFitOnGPU.h.

                                              {
  assert(N <= nHits);
 
  assert(results);
  assert(pfast_fit);
 
  // same as above...
 
  // look in bin for this hit multiplicity
  auto local_start = blockIdx.x * blockDim.x + threadIdx.x;
  for (int local_idx = local_start, nt = riemannFit::maxNumberOfConcurrentFits; local_idx < nt;
       local_idx += gridDim.x * blockDim.x) {
    auto tuple_idx = local_idx + offset;
    if (tuple_idx >= tupleMultiplicity->size(nHits))
      break;
 
    // get it for the ntuple container (one to one to helix)
    auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
 
    riemannFit::Map3xNd<N> hits(phits + local_idx);
    riemannFit::Map4d fast_fit(pfast_fit + local_idx);
    riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
 
    brokenline::PreparedBrokenLineData<N> data;
 
    brokenline::karimaki_circle_fit circle;
    riemannFit::LineFit line;
 
    brokenline::prepareBrokenLineData(hits, fast_fit, bField, data);
    brokenline::lineFit(hits_ge, fast_fit, bField, data, line);
    brokenline::circleFit(hits, hits_ge, fast_fit, bField, data, circle);
 
    results->stateAtBS.copyFromCircle(circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid);
    results->pt(tkid) = float(bField) / float(std::abs(circle.par(2)));
    results->eta(tkid) = asinhf(line.par(0));
    results->chi2(tkid) = (circle.chi2 + line.chi2) / (2 * N - 5);
 
#ifdef BROKENLINE_DEBUG
    if (!(circle.chi2 >= 0) || !(line.chi2 >= 0))
      printf("kernelBLFit failed! %f/%f\n", circle.chi2, line.chi2);
    printf("kernelBLFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
           N,
           nHits,
           tkid,
           circle.par(0),
           circle.par(1),
           circle.par(2));
    printf("kernelBLHits line.par(0,1): %d %f,%f\n", tkid, line.par(0), line.par(1));
    printf("kernelBLHits chi2 cov %f/%f  %e,%e,%e,%e,%e\n",
           circle.chi2,
           line.chi2,
           circle.cov(0, 0),
           circle.cov(1, 1),
           circle.cov(2, 2),
           line.cov(0, 0),
           line.cov(1, 1));
#endif
  }
}

References funct::abs(), assert(), Calorimetry_cff::bField, cms::cudacompat::blockDim, cms::cudacompat::blockIdx, riemannFit::CircleFit::chi2, brokenline::circleFit(), riemannFit::CircleFit::cov, data, dqmMemoryStats::float, cms::cudacompat::gridDim, hfClusterShapes_cfi::hits, mps_splice::line, brokenline::lineFit(), local_start, riemannFit::maxNumberOfConcurrentFits, N, nHits, nt, offset, riemannFit::CircleFit::par, pfast_fit, phits, phits_ge, brokenline::prepareBrokenLineData(), bookConverter::results, cms::cudacompat::threadIdx, and tupleMultiplicity.

assert() [1/5]

assert

(

foundNtuplets

)

assert() [2/5]

assert

(

hhp

)

assert() [3/5]

assert

(

hitsInFit<=

nHits

)

Referenced by __attribute__(), and for().

assert() [4/5]

assert

(

pfast_fit

)

assert() [5/5]

assert

(

tupleMultiplicity

)

for()

for

(

int

local_idx = local_start

)

Definition at line 53 of file BrokenLineFitOnGPU.h.

                                            {
    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(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);
 
#ifdef BL_DUMP_HITS
    __shared__ int done;
    done = 0;
    __syncthreads();
    bool dump = (foundNtuplets->size(tkid) == 5 && 0 == atomicAdd(&done, 1));
#endif
 
    // Prepare data structure
    auto const *hitId = foundNtuplets->begin(tkid);
 
    // #define YERR_FROM_DC
#ifdef YERR_FROM_DC
    // try to compute more precise error in y
    auto dx = hhp->xGlobal(hitId[hitsInFit - 1]) - hhp->xGlobal(hitId[0]);
    auto dy = hhp->yGlobal(hitId[hitsInFit - 1]) - hhp->yGlobal(hitId[0]);
    auto dz = hhp->zGlobal(hitId[hitsInFit - 1]) - hhp->zGlobal(hitId[0]);
    float ux, uy, uz;
#endif
    for (unsigned int i = 0; i < hitsInFit; ++i) {
      auto hit = hitId[i];
      float ge[6];
#ifdef YERR_FROM_DC
      auto const &dp = hhp->cpeParams().detParams(hhp->detectorIndex(hit));
      auto status = hhp->status(hit);
      int qbin = 4 - status.qBin;
      assert(qbin >= 0 && qbin < 5);
      bool nok = (status.isBigY | status.isOneY);
      // compute cotanbeta and use it to recompute error
      dp.frame.rotation().multiply(dx, dy, dz, ux, uy, uz);
      auto cb = std::abs(uy / uz);
      int bin = int(cb * (285.f / 150.f) * 8.f) - 4;
      bin = std::max(0, std::min(15, bin));
      float yerr = dp.sigmay[bin] * 1.e-4f;
      yerr *= dp.yfact[qbin];  // inflate
      yerr *= yerr;
      yerr += dp.apeYY;
      yerr = nok ? hhp->yerrLocal(hit) : yerr;
      dp.frame.toGlobal(hhp->xerrLocal(hit), 0, yerr, ge);
#else
      hhp->cpeParams()
          .detParams(hhp->detectorIndex(hit))
          .frame.toGlobal(hhp->xerrLocal(hit), 0, hhp->yerrLocal(hit), ge);
#endif
 
#ifdef BL_DUMP_HITS
      if (dump) {
        printf("Hit global: %d: %d hits.col(%d) << %f,%f,%f\n",
               tkid,
               hhp->detectorIndex(hit),
               i,
               hhp->xGlobal(hit),
               hhp->yGlobal(hit),
               hhp->zGlobal(hit));
        printf("Error: %d: %d  hits_ge.col(%d) << %e,%e,%e,%e,%e,%e\n",
               tkid,
               hhp->detetectorIndex(hit),
               i,
               ge[0],
               ge[1],
               ge[2],
               ge[3],
               ge[4],
               ge[5]);
      }
#endif
      hits.col(i) << hhp->xGlobal(hit), hhp->yGlobal(hit), hhp->zGlobal(hit);
      hits_ge.col(i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
    }
    brokenline::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));
  }

References cms::cudacompat::__syncthreads(), funct::abs(), assert(), cms::cudacompat::atomicAdd(), newFWLiteAna::bin, fileCollector::done, Calorimetry_cff::dp, submitPVValidationJobs::dump, PVValHelper::dx, PVValHelper::dy, PVValHelper::dz, f, brokenline::fastFit(), foundNtuplets, amptDefault_cfi::frame, hhp, hfClusterShapes_cfi::hits, mps_fire::i, createfilelist::int, SiStripPI::max, min(), nHits, offset, pfast_fit, phits, phits_ge, mps_update::status, and tupleMultiplicity.

Variable Documentation

hhp

const caConstants::TupleMultiplicity* __restrict__ const HitsOnGPU* __restrict__ hhp

Definition at line 27 of file BrokenLineFitOnGPU.h.

Referenced by for().

local_start

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

Definition at line 44 of file BrokenLineFitOnGPU.h.

Referenced by __attribute__().

nHits

const caConstants::TupleMultiplicity* __restrict__ const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ uint32_t nHits

Definition at line 27 of file BrokenLineFitOnGPU.h.

Referenced by gpuPixelDoublets::__attribute__(), __attribute__(), DTMeantimerPatternReco::addHits(), SeedFinder::addHitSelector(), CAHitNtupletGeneratorKernels< TTraits >::allocateOnGPU(), DTSegmentsTask::analyze(), DTOccupancyEfficiency::analyze(), DTSegmentAnalysisTask::analyze(), BtlLocalRecoValidation::analyze(), PhotonValidator::analyze(), HcalHitValidation::analyzeLayer(), RecoMuonValidator::bookHistograms(), RecoMuonValidator::MuonME::bookHistos(), FWTracksterLayersProxyBuilder::build(), PtAssignmentEngine2017::calculate_address(), PtAssignmentEngine2017::calculate_pt_xml(), RPCSeedOverlapper::CheckOverlap(), DAFTrackProducerAlgorithm::collectHits(), OMTFResult::empty(), SimG4HcalValidation::fetchHits(), DTSegmentAnalysisTask::fillHistos(), TrackingNtuple::fillSeeds(), CastorShowerLibraryMaker::FillShowerEvent(), TrackingNtuple::fillTracks(), gpuVertexFinder::for(), for(), DTChamberEfficiencyTask::getBestSegment(), MuonSeedCleaner::GroupSeeds(), MillePedeMonitor::init(), DTChamberEfficiencyTask::isGoodSegment(), PFDisplacedVertexHelper::isTrackSelected(), MuonSeedCleaner::LengthFilter(), BTLUncalibRecHitAlgo::makeRecHit(), LowPtConversion::match(), CovarianceParameterization::meanValue(), CovarianceParameterization::pack(), SETSeedFinder::pre_prune(), CtfSpecialSeedGenerator::preliminaryCheck(), PixelTemplateSmearerBase::process(), QualityFilter::produce(), SeedGeneratorFromProtoTracksEDProducer::produce(), TrackListCombiner::produce(), PixelTracksProducer::produce(), TSGFromL1Muon::produce(), HiBadParticleCleaner::produce(), PixelTrackSoAFromCUDA::produce(), SeedProducerFromSoA::produce(), PixelTrackProducerFromSoA::produce(), PixelVertexProducerCUDA::produceOnCPU(), CSCSegAlgoShowering::pruneFromResidual(), CSCSegAlgoDF::pruneFromResidual(), SeedFromGenericPairOrTriplet::qualityFilter(), PFAlgo::recoTracksNotHCAL(), ReferenceTrajectoryBase::ReferenceTrajectoryBase(), reco::TrackResiduals::resize(), PixelTrackReconstruction::run(), ConversionTrackPairFinder::run(), CSCSegAlgoDF::run(), EEBadScFilter::scan5x5(), MuonSeedCleaner::SeedCandidates(), SeedFromProtoTrack::SeedFromProtoTrack(), DTCombinatorialPatternReco4D::segmentSpecialZed(), CSCBaseElectronicsSim::simulate(), storeTracks(), hi::EPCuts::trackQuality_Pixel(), and CovarianceParameterization::unpack().

offset

const caConstants::TupleMultiplicity* __restrict__ const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ uint32_t uint32_t offset

Initial value:

{
  constexpr uint32_t hitsInFit = N

Definition at line 33 of file BrokenLineFitOnGPU.h.

Referenced by __attribute__(), and for().

pfast_fit

const caConstants::TupleMultiplicity* __restrict__ const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ pfast_fit

Definition at line 27 of file BrokenLineFitOnGPU.h.

Referenced by __attribute__(), and for().

phits

const caConstants::TupleMultiplicity* __restrict__ const HitsOnGPU* __restrict__ double* __restrict__ phits

Definition at line 27 of file BrokenLineFitOnGPU.h.

Referenced by gpuPixelRecHits::__attribute__(), __attribute__(), EcalTBValidation::analyze(), EcalSimpleTBAnalyzer::analyze(), EcalSimple2007H4TBAnalyzer::analyze(), ElectronCalibration::analyze(), FWTrackingParticleProxyBuilderFullFramework::build(), ZeeCalibration::duringLoop(), ElectronCalibration::findMaxHit(), for(), JetMatchingTools::getGenParticles(), PhotonConversionTrajectorySeedProducerFromQuadrupletsAlgo::inspect(), JetMatchingTools::lostEnergyFraction(), and SeedForPhotonConversionFromQuadruplets::trajectorySeed().

phits_ge

const caConstants::TupleMultiplicity* __restrict__ const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ phits_ge

Definition at line 27 of file BrokenLineFitOnGPU.h.

Referenced by __attribute__(), and for().

tupleMultiplicity

const Quality *__restrict__ caConstants::TupleMultiplicity * tupleMultiplicity

Initial value:

{
  auto first = blockIdx.x * blockDim.x + threadIdx.x

Definition at line 27 of file BrokenLineFitOnGPU.h.

Referenced by __attribute__(), HelixFitOnGPU::allocateOnGPU(), and for().