BrokenLineFitOnGPU.h File Reference

#include <cstdint>
#include <cuda_runtime.h>
#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitsUtilities.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
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_BLFastFit(Tuples< TrackerTraits > const *__restrict__ foundNtuplets
	assert (hitsInFit<=nHitsL)
	assert (nHitsL<=nHitsH)
	assert (phits)
	assert (pfast_fit)
	assert (foundNtuplets)
	assert (tupleMultiplicity)
	assert (totTK<=int(tupleMultiplicity->size()))
	assert (totTK >=0)
	assert (results_view.eta())
	for (int local_idx=local_start, nt=riemannFit::maxNumberOfConcurrentFits;local_idx< nt;local_idx+=gridDim.x *blockDim.x)

Variables
double	bField
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits >	hh
constexpr auto	invalidTkId = std::numeric_limits<typename TrackerTraits::tindex_type>::max()
auto	local_start = blockIdx.x * blockDim.x + threadIdx.x
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t uint32_t	nHitsH
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t	nHitsL
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t uint32_t int32_t	offset
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__	pfast_fit
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__	phits
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__	phits_ge
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__	ptkids
double OutputSoAView< TrackerTraits >	results_view
int	totTK = tupleMultiplicity->end(nHitsH) - tupleMultiplicity->begin(nHitsL)
TupleMultiplicity< TrackerTraits > const *__restrict__	tupleMultiplicity

Typedef Documentation

OutputSoAView

template<typename TrackerTraits >

using OutputSoAView = TrackSoAView<TrackerTraits>

Definition at line 22 of file BrokenLineFitOnGPU.h.

TupleMultiplicity

template<typename TrackerTraits >

using TupleMultiplicity = caStructures::TupleMultiplicityT<TrackerTraits>

Definition at line 24 of file BrokenLineFitOnGPU.h.

Tuples

template<typename TrackerTraits >

using Tuples = typename TrackSoA<TrackerTraits>::HitContainer

Definition at line 20 of file BrokenLineFitOnGPU.h.

Function Documentation

__attribute__()

template<int N, typename TrackerTraits >

__attribute__ ( (always_inline)  ) const

inline

assert() [1/9]

assert

(

hitsInFit<=

nHitsL

)

Referenced by for().

assert() [2/9]

assert

(

nHitsL<=

nHitsH

)

assert() [3/9]

assert

(

phits

)

assert() [4/9]

assert

(

pfast_fit

)

assert() [5/9]

assert

(

foundNtuplets

)

assert() [6/9]

assert

(

tupleMultiplicity

)

assert() [7/9]

assert

(

totTK<=

inttupleMultiplicity->size()

)

assert() [8/9]

assert

(

totTK >=

0

)

assert() [9/9]

assert

(

results_view.

eta()

)

for()

for

(

int

local_idx = local_start

)

Definition at line 62 of file BrokenLineFitOnGPU.h.

References cms::cudacompat::__syncthreads(), funct::abs(), assert(), cms::cudacompat::atomicAdd(), newFWLiteAna::bin, fileCollector::done, Calorimetry_cff::dp, GCP_Ntuples_cfg::dump, PVValHelper::dx, PVValHelper::dy, PVValHelper::dz, f, brokenline::fastFit(), caHitNtupletGeneratorKernels::foundNtuplets, hh, hfClusterShapes_cfi::hits, mps_fire::i, createfilelist::int, invalidTkId, dqmiolumiharvest::j, CPEFastParametrisation::kGenErrorQBins, CPEFastParametrisation::kNumErrorBins, SiStripPI::max, dqmiodumpmetadata::n, nHits, nHitsH, nHitsL, offset, pfast_fit, phits, phits_ge, ptkids, mps_update::status, totTK, and tupleMultiplicity.

                                            {
    int tuple_idx = local_idx + offset;
    if (tuple_idx >= totTK) {
      ptkids[local_idx] = invalidTkId;
      break;
    }
    // get it from the ntuple container (one to one to helix)
    auto tkid = *(tupleMultiplicity->begin(nHitsL) + tuple_idx);
    assert(int(tkid) < foundNtuplets->nOnes());

    ptkids[local_idx] = tkid;

    auto nHits = foundNtuplets->size(tkid);

    assert(nHits >= nHitsL);
    assert(nHits <= nHitsH);

    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 = hh[hitId[hitsInFit - 1]].xGlobal() - hh[hitId[0]].xGlobal();
    auto dy = hh[hitId[hitsInFit - 1]].yGlobal() - hh[hitId[0]].yGlobal();
    auto dz = hh[hitId[hitsInFit - 1]].zGlobal() - hh[hitId[0]].zGlobal();
    float ux, uy, uz;
#endif

    float incr = std::max(1.f, float(nHits) / float(hitsInFit));
    float n = 0;
    for (uint32_t i = 0; i < hitsInFit; ++i) {
      int j = int(n + 0.5f);  // round
      if (hitsInFit - 1 == i)
        j = nHits - 1;  // force last hit to ensure max lever arm.
      assert(j < int(nHits));
      n += incr;
      auto hit = hitId[j];
      float ge[6];

#ifdef YERR_FROM_DC
      auto const &dp = hh.cpeParams().detParams(hh.detectorIndex(hit));
      auto status = hh[hit].chargeAndStatus().status;
      int qbin = CPEFastParametrisation::kGenErrorQBins - 1 - 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 * (float(phase1PixelTopology::pixelThickess) / float(phase1PixelTopology::pixelPitchY)) * 8.f) - 4;
      int low_value = 0;
      int high_value = CPEFastParametrisation::kNumErrorBins - 1;
      // return estimated bin value truncated to [0, 15]
      bin = std::clamp(bin, low_value, high_value);
      float yerr = dp.sigmay[bin] * 1.e-4f;  // toCM
      yerr *= dp.yfact[qbin];                // inflate
      yerr *= yerr;
      yerr += dp.apeYY;
      yerr = nok ? hh[hit].yerrLocal() : yerr;
      dp.frame.toGlobal(hh[hit].xerrLocal(), 0, yerr, ge);
#else
      hh.cpeParams().detParams(hh[hit].detectorIndex()).frame.toGlobal(hh[hit].xerrLocal(), 0, hh[hit].yerrLocal(), ge);
#endif

#ifdef BL_DUMP_HITS
      bool dump = foundNtuplets->size(tkid) == 5;
      if (dump) {
        printf("Track id %d %d Hit %d on %d\nGlobal: hits.col(%d) << %f,%f,%f\n",
               local_idx,
               tkid,
               hit,
               hh[hit].detectorIndex(),
               i,
               hh[hit].xGlobal(),
               hh[hit].yGlobal(),
               hh[hit].zGlobal());
        printf("Error: hits_ge.col(%d) << %e,%e,%e,%e,%e,%e\n", i, ge[0], ge[1], ge[2], ge[3], ge[4], ge[5]);
      }
#endif

      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];
    }
    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));
  }

Variable Documentation

bField

uint32_t double bField

Definition at line 170 of file BrokenLineFitOnGPU.h.

hh

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

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by Histos::addObject(), HtrXmlPattern::analyze(), RecAnalyzerMinbias::beginJob(), FWTEveViewer::CaptureAndSaveImage(), PFEnergyCalibration::energyEmHad(), Histos::fill(), HcalObjRepresent::FillUnphysicalHEHFBins(), TwoTrackMinimumDistance::firstAngle(), for(), MultiHitGeneratorFromChi2::hitSets(), CaloTowerHardcodeGeometryLoader::makeCell(), TwoTrackMinimumDistance::pathLength(), TwoTrackMinimumDistance::pointsHelixHelix(), TrackListMerger::produce(), Histos::put(), FWTEveViewer::SaveJpg(), FWTEveViewer::SavePng(), TwoTrackMinimumDistance::secondAngle(), CmsShowMainFrame::showFWorksInfo(), and DCULVRVoltagesDat::writeArrayDB().

invalidTkId

constexpr auto invalidTkId = std::numeric_limits<typename TrackerTraits::tindex_type>::max()

Definition at line 40 of file BrokenLineFitOnGPU.h.

Referenced by for().

local_start

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

Definition at line 50 of file BrokenLineFitOnGPU.h.

nHitsH

TupleMultiplicity<TrackerTraits> const* __restrict__ TrackingRecHitSoAConstView<TrackerTraits> TrackerTraits::tindex_type* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ uint32_t uint32_t nHitsH

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by for().

nHitsL

TupleMultiplicity<TrackerTraits> const* __restrict__ TrackingRecHitSoAConstView<TrackerTraits> TrackerTraits::tindex_type* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ uint32_t nHitsL

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by for().

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 38 of file BrokenLineFitOnGPU.h.

Referenced by for().

pfast_fit

double OutputSoAView< TrackerTraits > TrackerTraits::tindex_type const *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ pfast_fit

Initial value:

{
  assert(results_view.pt())

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by for().

phits

uint32_t double OutputSoAView< TrackerTraits > double *__restrict__ phits

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by 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

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

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by for().

ptkids

double OutputSoAView< TrackerTraits > TrackerTraits::tindex_type const *__restrict__ ptkids

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by for().

results_view

double OutputSoAView<TrackerTraits> results_view

Definition at line 170 of file BrokenLineFitOnGPU.h.

totTK

int totTK = tupleMultiplicity->end(nHitsH) - tupleMultiplicity->begin(nHitsL)

Definition at line 51 of file BrokenLineFitOnGPU.h.

Referenced by for().

tupleMultiplicity

TupleMultiplicity<TrackerTraits> const* __restrict__ tupleMultiplicity

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by HelixFitOnGPU< TrackerTraits >::allocateOnGPU(), and for().