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  // no NaN here....
168  }
169  }
170  };
171 
172  template <int N, typename TrackerTraits>
173  struct Kernel_BLFit {
174  public:
175  template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
176  ALPAKA_FN_ACC void operator()(TAcc const &acc,
178  double bField,
180  typename TrackerTraits::tindex_type const *__restrict__ ptkids,
181  double *__restrict__ phits,
182  float *__restrict__ phits_ge,
183  double *__restrict__ pfast_fit) const {
189 
190  // same as above...
191  // look in bin for this hit multiplicity
193  for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
194  if (invalidTkId == ptkids[local_idx])
195  break;
196  auto tkid = ptkids[local_idx];
197 
198  ALPAKA_ASSERT_ACC(tkid < TrackerTraits::maxNumberOfTuples);
199 
200  riemannFit::Map3xNd<N> hits(phits + local_idx);
201  riemannFit::Map4d fast_fit(pfast_fit + local_idx);
202  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
203 
205 
208 
210  brokenline::lineFit(acc, hits_ge, fast_fit, bField, data, line);
211  brokenline::circleFit(acc, hits, hits_ge, fast_fit, bField, data, circle);
212 
214  results_view, circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid);
215  results_view[tkid].pt() = float(bField) / float(std::abs(circle.par(2)));
216  results_view[tkid].eta() = alpaka::math::asinh(acc, line.par(0));
217  results_view[tkid].chi2() = (circle.chi2 + line.chi2) / (2 * N - 5);
218 
219 #ifdef BROKENLINE_DEBUG
220  if (!(circle.chi2 >= 0) || !(line.chi2 >= 0))
221  printf("kernelBLFit failed! %f/%f\n", circle.chi2, line.chi2);
222  printf("kernelBLFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
223  N,
224  N,
225  tkid,
226  circle.par(0),
227  circle.par(1),
228  circle.par(2));
229  printf("kernelBLHits line.par(0,1): %d %f,%f\n", tkid, line.par(0), line.par(1));
230  printf("kernelBLHits chi2 cov %f/%f %e,%e,%e,%e,%e\n",
231  circle.chi2,
232  line.chi2,
233  circle.cov(0, 0),
234  circle.cov(1, 1),
235  circle.cov(2, 2),
236  line.cov(0, 0),
237  line.cov(1, 1));
238 #endif
239  }
240  }
241  };
242 
243  template <typename TrackerTraits>
247  uint32_t hitsInFit,
248  uint32_t maxNumberOfTuples,
249  Queue &queue) {
250  ALPAKA_ASSERT_ACC(tuples_);
251 
252  uint32_t blockSize = 64;
253  uint32_t numberOfBlocks = cms::alpakatools::divide_up_by(maxNumberOfConcurrentFits_, blockSize);
254  const WorkDiv1D workDivTriplets = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks, blockSize);
255  const WorkDiv1D workDivQuadsPenta = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks / 4, blockSize);
256 
257  // Fit internals
258  auto tkidDevice =
259  cms::alpakatools::make_device_buffer<typename TrackerTraits::tindex_type[]>(queue, maxNumberOfConcurrentFits_);
260  auto hitsDevice = cms::alpakatools::make_device_buffer<double[]>(
261  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Matrix3xNd<6>) / sizeof(double));
262  auto hits_geDevice = cms::alpakatools::make_device_buffer<float[]>(
263  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Matrix6xNf<6>) / sizeof(float));
264  auto fast_fit_resultsDevice = cms::alpakatools::make_device_buffer<double[]>(
265  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Vector4d) / sizeof(double));
266 
267  for (uint32_t offset = 0; offset < maxNumberOfTuples; offset += maxNumberOfConcurrentFits_) {
268  // fit triplets
269 
270  alpaka::exec<Acc1D>(queue,
271  workDivTriplets,
273  tuples_,
274  tupleMultiplicity_,
275  hv,
276  cpeParams,
277  tkidDevice.data(),
278  hitsDevice.data(),
279  hits_geDevice.data(),
280  fast_fit_resultsDevice.data(),
281  3,
282  3,
283  offset);
284 
285  alpaka::exec<Acc1D>(queue,
286  workDivTriplets,
288  tupleMultiplicity_,
289  bField_,
290  outputSoa_,
291  tkidDevice.data(),
292  hitsDevice.data(),
293  hits_geDevice.data(),
294  fast_fit_resultsDevice.data());
295 
296  if (fitNas4_) {
297  // fit all as 4
298  riemannFit::rolling_fits<4, TrackerTraits::maxHitsOnTrack, 1>([this,
299  &hv,
300  &cpeParams,
301  &tkidDevice,
302  &hitsDevice,
303  &hits_geDevice,
304  &fast_fit_resultsDevice,
305  &offset,
306  &queue,
307  &workDivQuadsPenta](auto i) {
308  alpaka::exec<Acc1D>(queue,
309  workDivQuadsPenta,
311  tuples_,
312  tupleMultiplicity_,
313  hv,
314  cpeParams,
315  tkidDevice.data(),
316  hitsDevice.data(),
317  hits_geDevice.data(),
318  fast_fit_resultsDevice.data(),
319  4,
320  4,
321  offset);
322 
323  alpaka::exec<Acc1D>(queue,
324  workDivQuadsPenta,
326  tupleMultiplicity_,
327  bField_,
328  outputSoa_,
329  tkidDevice.data(),
330  hitsDevice.data(),
331  hits_geDevice.data(),
332  fast_fit_resultsDevice.data());
333  });
334 
335  } else {
336  riemannFit::rolling_fits<4, TrackerTraits::maxHitsOnTrackForFullFit, 1>([this,
337  &hv,
338  &cpeParams,
339  &tkidDevice,
340  &hitsDevice,
341  &hits_geDevice,
342  &fast_fit_resultsDevice,
343  &offset,
344  &queue,
345  &workDivQuadsPenta](auto i) {
346  alpaka::exec<Acc1D>(queue,
347  workDivQuadsPenta,
349  tuples_,
350  tupleMultiplicity_,
351  hv,
352  cpeParams,
353  tkidDevice.data(),
354  hitsDevice.data(),
355  hits_geDevice.data(),
356  fast_fit_resultsDevice.data(),
357  i,
358  i,
359  offset);
360 
361  alpaka::exec<Acc1D>(queue,
362  workDivQuadsPenta,
364  tupleMultiplicity_,
365  bField_,
366  outputSoa_,
367  tkidDevice.data(),
368  hitsDevice.data(),
369  hits_geDevice.data(),
370  fast_fit_resultsDevice.data());
371  });
372 
373  static_assert(TrackerTraits::maxHitsOnTrackForFullFit < TrackerTraits::maxHitsOnTrack);
374 
375  //Fit all the rest using the maximum from previous call
376  alpaka::exec<Acc1D>(queue,
377  workDivQuadsPenta,
379  tuples_,
380  tupleMultiplicity_,
381  hv,
382  cpeParams,
383  tkidDevice.data(),
384  hitsDevice.data(),
385  hits_geDevice.data(),
386  fast_fit_resultsDevice.data(),
387  TrackerTraits::maxHitsOnTrackForFullFit,
388  TrackerTraits::maxHitsOnTrack - 1,
389  offset);
390 
391  alpaka::exec<Acc1D>(queue,
392  workDivQuadsPenta,
394  tupleMultiplicity_,
395  bField_,
396  outputSoa_,
397  tkidDevice.data(),
398  hitsDevice.data(),
399  hits_geDevice.data(),
400  fast_fit_resultsDevice.data());
401  }
402 
403  } // loop on concurrent fits
404  }
405 
406  template class HelixFit<pixelTopology::Phase1>;
407  template class HelixFit<pixelTopology::Phase2>;
409 
410 } // 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
constexpr DetParams const &__restrict__ detParams(int i) const
ALPAKA_FN_ACC auto uniform_elements(TAcc const &acc, TArgs... args)
Definition: workdivision.h:303
constexpr Idx divide_up_by(Idx value, Idx divisor)
Definition: workdivision.h:19
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