CMS 3D CMS Logo

BrokenLineFit.dev.cc
Go to the documentation of this file.
1 //#define BROKENLINE_DEBUG
2 //#define BL_DUMP_HITS
3 
4 #include <cstdint>
5 
6 #include <alpaka/alpaka.hpp>
7 
13 
14 #include "HelixFit.h"
15 
16 template <typename TrackerTraits>
18 template <typename TrackerTraits>
20 template <typename TrackerTraits>
22 
23 // #define BL_DUMP_HITS
24 
26  template <int N, typename TrackerTraits>
28  public:
29  template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
30  ALPAKA_FN_ACC void operator()(TAcc const &acc,
31  Tuples<TrackerTraits> const *__restrict__ foundNtuplets,
34  pixelCPEforDevice::ParamsOnDeviceT<TrackerTraits> const *__restrict__ cpeParams,
35  typename TrackerTraits::tindex_type *__restrict__ ptkids,
36  double *__restrict__ phits,
37  float *__restrict__ phits_ge,
38  double *__restrict__ pfast_fit,
39  uint32_t nHitsL,
40  uint32_t nHitsH,
41  int32_t offset) const {
42  constexpr uint32_t hitsInFit = N;
44 
45  ALPAKA_ASSERT_ACC(hitsInFit <= nHitsL);
51 
52  // look in bin for this hit multiplicity
54  ALPAKA_ASSERT_ACC(totTK <= int(tupleMultiplicity->size()));
56 
57 #ifdef BROKENLINE_DEBUG
58  const uint32_t threadIdx(alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u]);
60  printf("%d total Ntuple\n", tupleMultiplicity->size());
61  printf("%d Ntuple of size %d/%d for %d hits to fit\n", totTK, nHitsL, nHitsH, hitsInFit);
62  }
63 #endif
65  for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
66  auto tuple_idx = local_idx + offset;
67  if ((int)tuple_idx >= totTK) {
68  ptkids[local_idx] = invalidTkId;
69  break;
70  }
71  // get it from the ntuple container (one to one to helix)
72  auto tkid = *(tupleMultiplicity->begin(nHitsL) + tuple_idx);
73  ALPAKA_ASSERT_ACC(static_cast<int>(tkid) < foundNtuplets->nOnes());
74 
75  ptkids[local_idx] = tkid;
76 
77  auto nHits = foundNtuplets->size(tkid);
78 
81 
82  riemannFit::Map3xNd<N> hits(phits + local_idx);
84  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
85 
86 #ifdef BL_DUMP_HITS
87  auto &&done = alpaka::declareSharedVar<int, __COUNTER__>(acc);
88  done = 0;
89  alpaka::syncBlockThreads(acc);
90  bool dump =
91  (foundNtuplets->size(tkid) == 5 && 0 == alpaka::atomicAdd(acc, &done, 1, alpaka::hierarchy::Blocks{}));
92 #endif
93 
94  // Prepare data structure
95  auto const *hitId = foundNtuplets->begin(tkid);
96 
97  // #define YERR_FROM_DC
98 #ifdef YERR_FROM_DC
99  // try to compute more precise error in y
100  auto dx = hh[hitId[hitsInFit - 1]].xGlobal() - hh[hitId[0]].xGlobal();
101  auto dy = hh[hitId[hitsInFit - 1]].yGlobal() - hh[hitId[0]].yGlobal();
102  auto dz = hh[hitId[hitsInFit - 1]].zGlobal() - hh[hitId[0]].zGlobal();
103  float ux, uy, uz;
104 #endif
105 
106  float incr = std::max(1.f, float(nHits) / float(hitsInFit));
107  float n = 0;
108  for (uint32_t i = 0; i < hitsInFit; ++i) {
109  int j = int(n + 0.5f); // round
110  if (hitsInFit - 1 == i)
111  j = nHits - 1; // force last hit to ensure max lever arm.
112  ALPAKA_ASSERT_ACC(j < int(nHits));
113  n += incr;
114  auto hit = hitId[j];
115  float ge[6];
116 
117 #ifdef YERR_FROM_DC
118  auto const &dp = cpeParams->detParams(hh.detectorIndex(hit));
119  auto status = hh[hit].chargeAndStatus().status;
120  int qbin = CPEFastParametrisation::kGenErrorQBins - 1 - status.qBin;
121  ALPAKA_ASSERT_ACC(qbin >= 0 && qbin < 5);
122  bool nok = (status.isBigY | status.isOneY);
123  // compute cotanbeta and use it to recompute error
124  dp.frame.rotation().multiply(dx, dy, dz, ux, uy, uz);
125  auto cb = std::abs(uy / uz);
126  int bin =
127  int(cb * (float(phase1PixelTopology::pixelThickess) / float(phase1PixelTopology::pixelPitchY)) * 8.f) - 4;
128  int low_value = 0;
129  int high_value = CPEFastParametrisation::kNumErrorBins - 1;
130  // return estimated bin value truncated to [0, 15]
131  bin = std::clamp(bin, low_value, high_value);
132  float yerr = dp.sigmay[bin] * 1.e-4f; // toCM
133  yerr *= dp.yfact[qbin]; // inflate
134  yerr *= yerr;
135  yerr += dp.apeYY;
136  yerr = nok ? hh[hit].yerrLocal() : yerr;
137  dp.frame.toGlobal(hh[hit].xerrLocal(), 0, yerr, ge);
138 #else
139  cpeParams->detParams(hh[hit].detectorIndex()).frame.toGlobal(hh[hit].xerrLocal(), 0, hh[hit].yerrLocal(), ge);
140 #endif
141 
142 #ifdef BL_DUMP_HITS
143  bool dump = foundNtuplets->size(tkid) == 5;
144  if (dump) {
145  printf("Track id %d %d Hit %d on %d\nGlobal: hits.col(%d) << %f,%f,%f\n",
146  local_idx,
147  tkid,
148  hit,
149  hh[hit].detectorIndex(),
150  i,
151  hh[hit].xGlobal(),
152  hh[hit].yGlobal(),
153  hh[hit].zGlobal());
154  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]);
155  }
156 #endif
157 
158  hits.col(i) << hh[hit].xGlobal(), hh[hit].yGlobal(), hh[hit].zGlobal();
159  hits_ge.col(i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
160  }
162 
163 #ifdef BROKENLINE_DEBUG
164  // any NaN value should cause the track to be rejected at a later stage
169 #endif
170  }
171  }
172  };
173 
174  template <int N, typename TrackerTraits>
175  struct Kernel_BLFit {
176  public:
177  template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
178  ALPAKA_FN_ACC void operator()(TAcc const &acc,
180  double bField,
182  typename TrackerTraits::tindex_type const *__restrict__ ptkids,
183  double *__restrict__ phits,
184  float *__restrict__ phits_ge,
185  double *__restrict__ pfast_fit) const {
191 
192  // same as above...
193  // look in bin for this hit multiplicity
195  for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
196  if (invalidTkId == ptkids[local_idx])
197  break;
198  auto tkid = ptkids[local_idx];
199 
200  ALPAKA_ASSERT_ACC(tkid < TrackerTraits::maxNumberOfTuples);
201 
202  riemannFit::Map3xNd<N> hits(phits + local_idx);
203  riemannFit::Map4d fast_fit(pfast_fit + local_idx);
204  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
205 
207 
210 
212  brokenline::lineFit(acc, hits_ge, fast_fit, bField, data, line);
213  brokenline::circleFit(acc, hits, hits_ge, fast_fit, bField, data, circle);
214 
216  results_view, circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid);
217  results_view[tkid].pt() = float(bField) / float(std::abs(circle.par(2)));
218  results_view[tkid].eta() = alpaka::math::asinh(acc, line.par(0));
219  results_view[tkid].chi2() = (circle.chi2 + line.chi2) / (2 * N - 5);
220 
221 #ifdef BROKENLINE_DEBUG
222  if (!(circle.chi2 >= 0) || !(line.chi2 >= 0))
223  printf("kernelBLFit failed! %f/%f\n", circle.chi2, line.chi2);
224  printf("kernelBLFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
225  N,
226  N,
227  tkid,
228  circle.par(0),
229  circle.par(1),
230  circle.par(2));
231  printf("kernelBLHits line.par(0,1): %d %f,%f\n", tkid, line.par(0), line.par(1));
232  printf("kernelBLHits chi2 cov %f/%f %e,%e,%e,%e,%e\n",
233  circle.chi2,
234  line.chi2,
235  circle.cov(0, 0),
236  circle.cov(1, 1),
237  circle.cov(2, 2),
238  line.cov(0, 0),
239  line.cov(1, 1));
240 #endif
241  }
242  }
243  };
244 
245  template <typename TrackerTraits>
249  uint32_t hitsInFit,
250  uint32_t maxNumberOfTuples,
251  Queue &queue) {
252  ALPAKA_ASSERT_ACC(tuples_);
253 
254  uint32_t blockSize = 64;
255  uint32_t numberOfBlocks = cms::alpakatools::divide_up_by(maxNumberOfConcurrentFits_, blockSize);
256  const WorkDiv1D workDivTriplets = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks, blockSize);
257  const WorkDiv1D workDivQuadsPenta = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks / 4, blockSize);
258 
259  // Fit internals
260  auto tkidDevice =
261  cms::alpakatools::make_device_buffer<typename TrackerTraits::tindex_type[]>(queue, maxNumberOfConcurrentFits_);
262  auto hitsDevice = cms::alpakatools::make_device_buffer<double[]>(
263  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Matrix3xNd<6>) / sizeof(double));
264  auto hits_geDevice = cms::alpakatools::make_device_buffer<float[]>(
265  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Matrix6xNf<6>) / sizeof(float));
266  auto fast_fit_resultsDevice = cms::alpakatools::make_device_buffer<double[]>(
267  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Vector4d) / sizeof(double));
268 
269  for (uint32_t offset = 0; offset < maxNumberOfTuples; offset += maxNumberOfConcurrentFits_) {
270  // fit triplets
271 
272  alpaka::exec<Acc1D>(queue,
273  workDivTriplets,
275  tuples_,
276  tupleMultiplicity_,
277  hv,
278  cpeParams,
279  tkidDevice.data(),
280  hitsDevice.data(),
281  hits_geDevice.data(),
282  fast_fit_resultsDevice.data(),
283  3,
284  3,
285  offset);
286 
287  alpaka::exec<Acc1D>(queue,
288  workDivTriplets,
290  tupleMultiplicity_,
291  bField_,
292  outputSoa_,
293  tkidDevice.data(),
294  hitsDevice.data(),
295  hits_geDevice.data(),
296  fast_fit_resultsDevice.data());
297 
298  if (fitNas4_) {
299  // fit all as 4
300  riemannFit::rolling_fits<4, TrackerTraits::maxHitsOnTrack, 1>([this,
301  &hv,
302  &cpeParams,
303  &tkidDevice,
304  &hitsDevice,
305  &hits_geDevice,
306  &fast_fit_resultsDevice,
307  &offset,
308  &queue,
309  &workDivQuadsPenta](auto i) {
310  alpaka::exec<Acc1D>(queue,
311  workDivQuadsPenta,
313  tuples_,
314  tupleMultiplicity_,
315  hv,
316  cpeParams,
317  tkidDevice.data(),
318  hitsDevice.data(),
319  hits_geDevice.data(),
320  fast_fit_resultsDevice.data(),
321  4,
322  4,
323  offset);
324 
325  alpaka::exec<Acc1D>(queue,
326  workDivQuadsPenta,
328  tupleMultiplicity_,
329  bField_,
330  outputSoa_,
331  tkidDevice.data(),
332  hitsDevice.data(),
333  hits_geDevice.data(),
334  fast_fit_resultsDevice.data());
335  });
336 
337  } else {
338  riemannFit::rolling_fits<4, TrackerTraits::maxHitsOnTrackForFullFit, 1>([this,
339  &hv,
340  &cpeParams,
341  &tkidDevice,
342  &hitsDevice,
343  &hits_geDevice,
344  &fast_fit_resultsDevice,
345  &offset,
346  &queue,
347  &workDivQuadsPenta](auto i) {
348  alpaka::exec<Acc1D>(queue,
349  workDivQuadsPenta,
351  tuples_,
352  tupleMultiplicity_,
353  hv,
354  cpeParams,
355  tkidDevice.data(),
356  hitsDevice.data(),
357  hits_geDevice.data(),
358  fast_fit_resultsDevice.data(),
359  i,
360  i,
361  offset);
362 
363  alpaka::exec<Acc1D>(queue,
364  workDivQuadsPenta,
366  tupleMultiplicity_,
367  bField_,
368  outputSoa_,
369  tkidDevice.data(),
370  hitsDevice.data(),
371  hits_geDevice.data(),
372  fast_fit_resultsDevice.data());
373  });
374 
375  static_assert(TrackerTraits::maxHitsOnTrackForFullFit < TrackerTraits::maxHitsOnTrack);
376 
377  //Fit all the rest using the maximum from previous call
378  alpaka::exec<Acc1D>(queue,
379  workDivQuadsPenta,
381  tuples_,
382  tupleMultiplicity_,
383  hv,
384  cpeParams,
385  tkidDevice.data(),
386  hitsDevice.data(),
387  hits_geDevice.data(),
388  fast_fit_resultsDevice.data(),
389  TrackerTraits::maxHitsOnTrackForFullFit,
390  TrackerTraits::maxHitsOnTrack - 1,
391  offset);
392 
393  alpaka::exec<Acc1D>(queue,
394  workDivQuadsPenta,
396  tupleMultiplicity_,
397  bField_,
398  outputSoa_,
399  tkidDevice.data(),
400  hitsDevice.data(),
401  hits_geDevice.data(),
402  fast_fit_resultsDevice.data());
403  }
404 
405  } // loop on concurrent fits
406  }
407 
408  template class HelixFit<pixelTopology::Phase1>;
409  template class HelixFit<pixelTopology::Phase2>;
411 
412 } // namespace ALPAKA_ACCELERATOR_NAMESPACE
const dim3 threadIdx
Definition: cudaCompat.h:29
Eigen::Vector4d Vector4d
Definition: FitResult.h:12
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fastFit(const TAcc &acc, const M3xN &hits, V4 &result)
A very fast helix fit.
Definition: BrokenLine.h:266
def isnan(num)
constexpr DetParams const &__restrict__ detParams(int i) const
ALPAKA_FN_ACC auto uniform_elements(TAcc const &acc, TArgs... args)
Definition: workdivision.h:311
constexpr Idx divide_up_by(Idx value, Idx divisor)
Definition: workdivision.h:20
void launchBrokenLineKernels(const HitConstView &hv, ParamsOnDevice const *cpeParams, uint32_t nhits, uint32_t maxNumberOfTuples, Queue &queue)
constexpr int kGenErrorQBins
Vector3d par
parameter: (X0,Y0,R)
Definition: FitResult.h:24
typename reco::TrackSoA< TrackerTraits >::HitContainer Tuples
TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity
ALPAKA_FN_ACC void operator()(TAcc const &acc, TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity, double bField, OutputSoAView< TrackerTraits > results_view, typename TrackerTraits::tindex_type const *__restrict__ ptkids, double *__restrict__ phits, float *__restrict__ phits_ge, double *__restrict__ pfast_fit) const
ALPAKA_FN_ACC ALPAKA_FN_INLINE void lineFit(const TAcc &acc, const M6xN &hits_ge, const V4 &fast_fit, const double bField, const PreparedBrokenLineData< n > &data, riemannFit::LineFit &line_results)
Performs the Broken Line fit in the straight track case (that is, the fit parameters are only the int...
Definition: BrokenLine.h:472
__host__ __device__ void prepareBrokenLineData(const M3xN &hits, const V4 &fast_fit, const double bField, PreparedBrokenLineData< n > &results)
Computes the data needed for the Broken Line fit procedure that are mainly common for the circle and ...
Definition: BrokenLine.h:151
Eigen::Matrix< float, 6, N > Matrix6xNf
Definition: HelixFit.h:35
WorkDiv< Dim1D > WorkDiv1D
Definition: config.h:32
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< Vector4d, 0, Eigen::InnerStride< stride > > Map4d
Definition: HelixFit.h:39
Eigen::Map< Matrix3xNd< N >, 0, Eigen::Stride< 3 *stride, stride > > Map3xNd
Definition: HelixFit.h:32
ALPAKA_FN_ACC void operator()(TAcc const &acc, Tuples< TrackerTraits > const *__restrict__ foundNtuplets, TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity, TrackingRecHitSoAConstView< TrackerTraits > hh, pixelCPEforDevice::ParamsOnDeviceT< TrackerTraits > const *__restrict__ cpeParams, typename TrackerTraits::tindex_type *__restrict__ ptkids, double *__restrict__ phits, float *__restrict__ phits_ge, double *__restrict__ pfast_fit, uint32_t nHitsL, uint32_t nHitsH, int32_t offset) const
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ ptkids
std::vector< Block > Blocks
Definition: Block.h:99
Abs< T >::type abs(const T &t)
Definition: Abs.h:22
double f[11][100]
Eigen::Matrix< double, 3, N > Matrix3xNd
Definition: HelixFit.h:30
double OutputSoAView< TrackerTraits > results_view
Eigen::Map< Matrix6xNf< N >, 0, Eigen::Stride< 6 *stride, stride > > Map6xNf
Definition: HelixFit.h:37
ALPAKA_FN_ACC ALPAKA_FN_INLINE void uint32_t const uint32_t CACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > hh
int nt
Definition: AMPTWrapper.h:42
ALPAKA_FN_ACC ALPAKA_FN_INLINE void const M3xN const V4 & fast_fit
Definition: BrokenLine.h:158
#define N
Definition: blowfish.cc:9
int totTK
constexpr auto invalidTkId
typename TrackingRecHitSoA< TrackerTraits >::template TrackingRecHitSoALayout<>::ConstView TrackingRecHitSoAConstView
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t uint32_t nHitsH
ALPAKA_FN_ACC constexpr bool once_per_grid(TAcc const &acc)
char data[epos_bytes_allocation]
Definition: EPOS_Wrapper.h:80
data needed for the Broken Line fit procedure.
Definition: BrokenLine.h:30
static constexpr __host__ __device__ void copyFromCircle(TrackSoAView &tracks, V3 const &cp, M3 const &ccov, V2 const &lp, M2 const &lcov, float b, int32_t i)
TupleMultiplicity< TrackerTraits > const *__restrict__ uint32_t nHits
ALPAKA_FN_ACC ALPAKA_FN_INLINE void circleFit(const TAcc &acc, const M3xN &hits, const M6xN &hits_ge, const V4 &fast_fit, const double bField, PreparedBrokenLineData< n > &data, karimaki_circle_fit &circle_results)
Performs the Broken Line fit in the curved track case (that is, the fit parameters are the intercepti...
Definition: BrokenLine.h:323
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
reco::TrackSoAView< TrackerTraits > OutputSoAView
constexpr uint32_t maxNumberOfConcurrentFits
Definition: HelixFit.h:21
typename reco::TrackSoA< TrackerTraits >::template Layout<>::View TrackSoAView
Definition: TracksSoA.h:45