CMS 3D CMS Logo

Typedefs | Functions | Variables
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 "RecoTracker/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.

63  {
64  int tuple_idx = local_idx + offset;
65  if (tuple_idx >= totTK) {
66  ptkids[local_idx] = invalidTkId;
67  break;
68  }
69  // get it from the ntuple container (one to one to helix)
70  auto tkid = *(tupleMultiplicity->begin(nHitsL) + tuple_idx);
71  assert(int(tkid) < foundNtuplets->nOnes());
72 
73  ptkids[local_idx] = tkid;
74 
75  auto nHits = foundNtuplets->size(tkid);
76 
77  assert(nHits >= nHitsL);
78  assert(nHits <= nHitsH);
79 
80  riemannFit::Map3xNd<N> hits(phits + local_idx);
81  riemannFit::Map4d fast_fit(pfast_fit + local_idx);
82  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
83 
84 #ifdef BL_DUMP_HITS
85  __shared__ int done;
86  done = 0;
87  __syncthreads();
88  bool dump = (foundNtuplets->size(tkid) == 5 && 0 == atomicAdd(&done, 1));
89 #endif
90 
91  // Prepare data structure
92  auto const *hitId = foundNtuplets->begin(tkid);
93 
94  // #define YERR_FROM_DC
95 #ifdef YERR_FROM_DC
96  // try to compute more precise error in y
97  auto dx = hh[hitId[hitsInFit - 1]].xGlobal() - hh[hitId[0]].xGlobal();
98  auto dy = hh[hitId[hitsInFit - 1]].yGlobal() - hh[hitId[0]].yGlobal();
99  auto dz = hh[hitId[hitsInFit - 1]].zGlobal() - hh[hitId[0]].zGlobal();
100  float ux, uy, uz;
101 #endif
102 
103  float incr = std::max(1.f, float(nHits) / float(hitsInFit));
104  float n = 0;
105  for (uint32_t i = 0; i < hitsInFit; ++i) {
106  int j = int(n + 0.5f); // round
107  if (hitsInFit - 1 == i)
108  j = nHits - 1; // force last hit to ensure max lever arm.
109  assert(j < int(nHits));
110  n += incr;
111  auto hit = hitId[j];
112  float ge[6];
113 
114 #ifdef YERR_FROM_DC
115  auto const &dp = hh.cpeParams().detParams(hh.detectorIndex(hit));
116  auto status = hh[hit].chargeAndStatus().status;
117  int qbin = CPEFastParametrisation::kGenErrorQBins - 1 - status.qBin;
118  assert(qbin >= 0 && qbin < 5);
119  bool nok = (status.isBigY | status.isOneY);
120  // compute cotanbeta and use it to recompute error
121  dp.frame.rotation().multiply(dx, dy, dz, ux, uy, uz);
122  auto cb = std::abs(uy / uz);
123  int bin =
124  int(cb * (float(phase1PixelTopology::pixelThickess) / float(phase1PixelTopology::pixelPitchY)) * 8.f) - 4;
125  int low_value = 0;
126  int high_value = CPEFastParametrisation::kNumErrorBins - 1;
127  // return estimated bin value truncated to [0, 15]
128  bin = std::clamp(bin, low_value, high_value);
129  float yerr = dp.sigmay[bin] * 1.e-4f; // toCM
130  yerr *= dp.yfact[qbin]; // inflate
131  yerr *= yerr;
132  yerr += dp.apeYY;
133  yerr = nok ? hh[hit].yerrLocal() : yerr;
134  dp.frame.toGlobal(hh[hit].xerrLocal(), 0, yerr, ge);
135 #else
136  hh.cpeParams().detParams(hh[hit].detectorIndex()).frame.toGlobal(hh[hit].xerrLocal(), 0, hh[hit].yerrLocal(), ge);
137 #endif
138 
139 #ifdef BL_DUMP_HITS
140  bool dump = foundNtuplets->size(tkid) == 5;
141  if (dump) {
142  printf("Track id %d %d Hit %d on %d\nGlobal: hits.col(%d) << %f,%f,%f\n",
143  local_idx,
144  tkid,
145  hit,
146  hh[hit].detectorIndex(),
147  i,
148  hh[hit].xGlobal(),
149  hh[hit].yGlobal(),
150  hh[hit].zGlobal());
151  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]);
152  }
153 #endif
154 
155  hits.col(i) << hh[hit].xGlobal(), hh[hit].yGlobal(), hh[hit].zGlobal();
156  hits_ge.col(i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
157  }
158  brokenline::fastFit(hits, fast_fit);
159 
160  // no NaN here....
161  assert(fast_fit(0) == fast_fit(0));
162  assert(fast_fit(1) == fast_fit(1));
163  assert(fast_fit(2) == fast_fit(2));
164  assert(fast_fit(3) == fast_fit(3));
165  }
constexpr int kGenErrorQBins
TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity
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__ phits
Eigen::Map< Matrix3xNd< N >, 0, Eigen::Stride< 3 *stride, stride > > Map3xNd
Definition: HelixFitOnGPU.h:24
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ ptkids
__host__ __device__ void fastFit(const M3xN &hits, V4 &result)
A very fast helix fit.
Definition: BrokenLine.h:258
Abs< T >::type abs(const T &t)
Definition: Abs.h:22
double f[11][100]
Eigen::Map< Matrix6xNf< N >, 0, Eigen::Stride< 6 *stride, stride > > Map6xNf
Definition: HelixFitOnGPU.h:29
assert(hitsInFit<=nHitsL)
int totTK
constexpr auto invalidTkId
Eigen::Map< Vector4d, 0, Eigen::InnerStride< stride > > Map4d
Definition: HelixFitOnGPU.h:31
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t uint32_t nHitsH
void __syncthreads()
Definition: cudaCompat.h:132
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 > hh
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t nHits
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__ double *__restrict__ float *__restrict__ double *__restrict__ pfast_fit
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61
constexpr int kNumErrorBins

Variable Documentation

◆ bField

uint32_t double bField

Definition at line 170 of file BrokenLineFitOnGPU.h.

◆ hh

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

◆ 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
#define N
Definition: blowfish.cc:9

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:
{
double OutputSoAView< TrackerTraits > results_view
assert(hitsInFit<=nHitsL)

Definition at line 30 of file BrokenLineFitOnGPU.h.

Referenced by for().

◆ phits

uint32_t double OutputSoAView< TrackerTraits > double *__restrict__ phits

◆ 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().