CMS 3D CMS Logo

RiemannFit.dev.cc
Go to the documentation of this file.
1 #include <cstdint>
2 
3 #include <alpaka/alpaka.hpp>
4 
12 
13 #include "HelixFit.h"
14 #include "CAStructures.h"
15 
16 template <typename TrackerTraits>
18 template <typename TrackerTraits>
20 template <typename TrackerTraits>
22 
24  using namespace alpaka;
25  using namespace cms::alpakatools;
26 
27  template <int N, typename TrackerTraits>
29  public:
30  template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
31  ALPAKA_FN_ACC void operator()(TAcc const &acc,
32  Tuples<TrackerTraits> const *__restrict__ foundNtuplets,
34  uint32_t nHits,
36  pixelCPEforDevice::ParamsOnDeviceT<TrackerTraits> const *__restrict__ cpeParams,
37  double *__restrict__ phits,
38  float *__restrict__ phits_ge,
39  double *__restrict__ pfast_fit,
40  uint32_t offset) const {
41  constexpr uint32_t hitsInFit = N;
42 
43  ALPAKA_ASSERT_ACC(hitsInFit <= nHits);
44 
48 
49  // look in bin for this hit multiplicity
50 
51 #ifdef RIEMANN_DEBUG
52  const uint32_t threadIdx(alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u]);
54  printf("%d Ntuple of size %d for %d hits to fit\n", tupleMultiplicity->size(nHits), nHits, hitsInFit);
55 #endif
56 
58  for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
59  auto tuple_idx = local_idx + offset;
60  if (tuple_idx >= tupleMultiplicity->size(nHits))
61  break;
62 
63  // get it from the ntuple container (one to one to helix)
64  auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
65  ALPAKA_ASSERT_ACC(static_cast<int>(tkid) < foundNtuplets->nOnes());
66 
67  ALPAKA_ASSERT_ACC(foundNtuplets->size(tkid) == nHits);
68 
69  riemannFit::Map3xNd<N> hits(phits + local_idx);
71  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
72 
73  // Prepare data structure
74  auto const *hitId = foundNtuplets->begin(tkid);
75  for (unsigned int i = 0; i < hitsInFit; ++i) {
76  auto hit = hitId[i];
77  float ge[6];
78  cpeParams->detParams(hh[hit].detectorIndex()).frame.toGlobal(hh[hit].xerrLocal(), 0, hh[hit].yerrLocal(), ge);
79 
80  hits.col(i) << hh[hit].xGlobal(), hh[hit].yGlobal(), hh[hit].zGlobal();
81  hits_ge.col(i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
82  }
84 
85  // no NaN here....
90  }
91  }
92  };
93 
94  template <int N, typename TrackerTraits>
96  public:
97  template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
98  ALPAKA_FN_ACC void operator()(TAcc const &acc,
100  uint32_t nHits,
101  double bField,
102  double *__restrict__ phits,
103  float *__restrict__ phits_ge,
104  double *__restrict__ pfast_fit_input,
106  uint32_t offset) const {
109 
110  // same as above...
111 
112  // look in bin for this hit multiplicity
114  for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
115  auto tuple_idx = local_idx + offset;
116  if (tuple_idx >= tupleMultiplicity->size(nHits))
117  break;
118 
119  riemannFit::Map3xNd<N> hits(phits + local_idx);
121  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
122 
123  riemannFit::VectorNd<N> rad = (hits.block(0, 0, 2, N).colwise().norm());
124 
126  riemannFit::loadCovariance2D(acc, hits_ge, hits_cov);
127 
128  circle_fit[local_idx] =
129  riemannFit::circleFit(acc, hits.block(0, 0, 2, N), hits_cov, fast_fit, rad, bField, true);
130 
131 #ifdef RIEMANN_DEBUG
132 // auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
133 // printf("kernelCircleFit circle.par(0,1,2): %d %f,%f,%f\n", tkid,
134 // circle_fit[local_idx].par(0), circle_fit[local_idx].par(1), circle_fit[local_idx].par(2));
135 #endif
136  }
137  }
138  };
139 
140  template <int N, typename TrackerTraits>
142  public:
143  template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
144  ALPAKA_FN_ACC void operator()(TAcc const &acc,
146  uint32_t nHits,
147  double bField,
149  double *__restrict__ phits,
150  float *__restrict__ phits_ge,
151  double *__restrict__ pfast_fit_input,
152  riemannFit::CircleFit *__restrict__ circle_fit,
153  uint32_t offset) const {
156 
157  // same as above...
158 
159  // look in bin for this hit multiplicity
161  for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
162  auto tuple_idx = local_idx + offset;
163  if (tuple_idx >= tupleMultiplicity->size(nHits))
164  break;
165 
166  // get it for the ntuple container (one to one to helix)
167  int32_t tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
168 
169  riemannFit::Map3xNd<N> hits(phits + local_idx);
171  riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
172 
173  auto const &line_fit = riemannFit::lineFit(acc, hits, hits_ge, circle_fit[local_idx], fast_fit, bField, true);
174 
176 
178  circle_fit[local_idx].par,
179  circle_fit[local_idx].cov,
180  line_fit.par,
181  line_fit.cov,
182  1.f / float(bField),
183  tkid);
184  results_view[tkid].pt() = bField / std::abs(circle_fit[local_idx].par(2));
185  results_view[tkid].eta() = asinhf(line_fit.par(0));
186  results_view[tkid].chi2() = (circle_fit[local_idx].chi2 + line_fit.chi2) / (2 * N - 5);
187 
188 #ifdef RIEMANN_DEBUG
189  printf("kernelLineFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
190  N,
191  nHits,
192  tkid,
193  circle_fit[local_idx].par(0),
194  circle_fit[local_idx].par(1),
195  circle_fit[local_idx].par(2));
196  printf("kernelLineFit line.par(0,1): %d %f,%f\n", tkid, line_fit.par(0), line_fit.par(1));
197  printf("kernelLineFit chi2 cov %f/%f %e,%e,%e,%e,%e\n",
198  circle_fit[local_idx].chi2,
199  line_fit.chi2,
200  circle_fit[local_idx].cov(0, 0),
201  circle_fit[local_idx].cov(1, 1),
202  circle_fit[local_idx].cov(2, 2),
203  line_fit.cov(0, 0),
204  line_fit.cov(1, 1));
205 #endif
206  }
207  }
208  };
209 
210  template <typename TrackerTraits>
213  uint32_t nhits,
214  uint32_t maxNumberOfTuples,
215  Queue &queue) {
216  assert(tuples_);
217 
218  auto blockSize = 64;
219  auto numberOfBlocks = (maxNumberOfConcurrentFits_ + blockSize - 1) / blockSize;
220  const auto workDivTriplets = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks, blockSize);
221  const auto workDivQuadsPenta = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks / 4, blockSize);
222 
223  // Fit internals
224  auto hitsDevice = cms::alpakatools::make_device_buffer<double[]>(
225  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Matrix3xNd<4>) / sizeof(double));
226  auto hits_geDevice = cms::alpakatools::make_device_buffer<float[]>(
227  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Matrix6x4f) / sizeof(float));
228  auto fast_fit_resultsDevice = cms::alpakatools::make_device_buffer<double[]>(
229  queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::Vector4d) / sizeof(double));
230  auto circle_fit_resultsDevice_holder =
231  cms::alpakatools::make_device_buffer<char[]>(queue, maxNumberOfConcurrentFits_ * sizeof(riemannFit::CircleFit));
232  riemannFit::CircleFit *circle_fit_resultsDevice_ =
233  (riemannFit::CircleFit *)(circle_fit_resultsDevice_holder.data());
234 
235  for (uint32_t offset = 0; offset < maxNumberOfTuples; offset += maxNumberOfConcurrentFits_) {
236  // triplets
237  alpaka::exec<Acc1D>(queue,
238  workDivTriplets,
240  tuples_,
241  tupleMultiplicity_,
242  3,
243  hv,
244  cpeParams,
245  hitsDevice.data(),
246  hits_geDevice.data(),
247  fast_fit_resultsDevice.data(),
248  offset);
249 
250  alpaka::exec<Acc1D>(queue,
251  workDivTriplets,
253  tupleMultiplicity_,
254  3,
255  bField_,
256  hitsDevice.data(),
257  hits_geDevice.data(),
258  fast_fit_resultsDevice.data(),
259  circle_fit_resultsDevice_,
260  offset);
261 
262  alpaka::exec<Acc1D>(queue,
263  workDivTriplets,
265  tupleMultiplicity_,
266  3,
267  bField_,
268  outputSoa_,
269  hitsDevice.data(),
270  hits_geDevice.data(),
271  fast_fit_resultsDevice.data(),
272  circle_fit_resultsDevice_,
273  offset);
274 
275  // quads
276  alpaka::exec<Acc1D>(queue,
277  workDivQuadsPenta,
279  tuples_,
280  tupleMultiplicity_,
281  4,
282  hv,
283  cpeParams,
284  hitsDevice.data(),
285  hits_geDevice.data(),
286  fast_fit_resultsDevice.data(),
287  offset);
288 
289  alpaka::exec<Acc1D>(queue,
290  workDivQuadsPenta,
292  tupleMultiplicity_,
293  4,
294  bField_,
295  hitsDevice.data(),
296  hits_geDevice.data(),
297  fast_fit_resultsDevice.data(),
298  circle_fit_resultsDevice_,
299  offset);
300 
301  alpaka::exec<Acc1D>(queue,
302  workDivQuadsPenta,
304  tupleMultiplicity_,
305  4,
306  bField_,
307  outputSoa_,
308  hitsDevice.data(),
309  hits_geDevice.data(),
310  fast_fit_resultsDevice.data(),
311  circle_fit_resultsDevice_,
312  offset);
313 
314  if (fitNas4_) {
315  // penta
316  alpaka::exec<Acc1D>(queue,
317  workDivQuadsPenta,
319  tuples_,
320  tupleMultiplicity_,
321  5,
322  hv,
323  cpeParams,
324  hitsDevice.data(),
325  hits_geDevice.data(),
326  fast_fit_resultsDevice.data(),
327  offset);
328 
329  alpaka::exec<Acc1D>(queue,
330  workDivQuadsPenta,
332  tupleMultiplicity_,
333  5,
334  bField_,
335  hitsDevice.data(),
336  hits_geDevice.data(),
337  fast_fit_resultsDevice.data(),
338  circle_fit_resultsDevice_,
339  offset);
340 
341  alpaka::exec<Acc1D>(queue,
342  workDivQuadsPenta,
344  tupleMultiplicity_,
345  5,
346  bField_,
347  outputSoa_,
348  hitsDevice.data(),
349  hits_geDevice.data(),
350  fast_fit_resultsDevice.data(),
351  circle_fit_resultsDevice_,
352  offset);
353  } else {
354  // penta all 5
355  alpaka::exec<Acc1D>(queue,
356  workDivQuadsPenta,
358  tuples_,
359  tupleMultiplicity_,
360  5,
361  hv,
362  cpeParams,
363  hitsDevice.data(),
364  hits_geDevice.data(),
365  fast_fit_resultsDevice.data(),
366  offset);
367 
368  alpaka::exec<Acc1D>(queue,
369  workDivQuadsPenta,
371  tupleMultiplicity_,
372  5,
373  bField_,
374  hitsDevice.data(),
375  hits_geDevice.data(),
376  fast_fit_resultsDevice.data(),
377  circle_fit_resultsDevice_,
378  offset);
379 
380  alpaka::exec<Acc1D>(queue,
381  workDivQuadsPenta,
383  tupleMultiplicity_,
384  5,
385  bField_,
386  outputSoa_,
387  hitsDevice.data(),
388  hits_geDevice.data(),
389  fast_fit_resultsDevice.data(),
390  circle_fit_resultsDevice_,
391  offset);
392  }
393  }
394  }
395 
396  template class HelixFit<pixelTopology::Phase1>;
397  template class HelixFit<pixelTopology::Phase2>;
399 
400 } // namespace ALPAKA_ACCELERATOR_NAMESPACE
const dim3 threadIdx
Definition: cudaCompat.h:29
Eigen::Vector4d Vector4d
Definition: FitResult.h:12
constexpr DetParams const &__restrict__ detParams(int i) const
ALPAKA_FN_ACC auto uniform_elements(TAcc const &acc, TArgs... args)
Definition: workdivision.h:303
ALPAKA_FN_ACC ALPAKA_FN_INLINE LineFit lineFit(const TAcc &acc, const M3xN &hits, const M6xN &hits_ge, const CircleFit &circle, const V4 &fast_fit, const double bField, const bool error)
Perform an ordinary least square fit in the s-z plane to compute the parameters cotTheta and Zip...
Definition: RiemannFit.h:796
ALPAKA_FN_ACC void loadCovariance2D(const TAcc &acc, M6xNf const &ge, M2Nd &hits_cov)
Definition: FitUtils.h:126
ALPAKA_FN_ACC void operator()(TAcc const &acc, Tuples< TrackerTraits > const *__restrict__ foundNtuplets, TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity, uint32_t nHits, TrackingRecHitSoAConstView< TrackerTraits > hh, pixelCPEforDevice::ParamsOnDeviceT< TrackerTraits > const *__restrict__ cpeParams, double *__restrict__ phits, float *__restrict__ phits_ge, double *__restrict__ pfast_fit, uint32_t offset) const
Eigen::Matrix< float, 6, 4 > Matrix6x4f
Definition: HelixFit.h:25
Vector3d par
parameter: (X0,Y0,R)
Definition: FitResult.h:24
assert(be >=bs)
typename reco::TrackSoA< TrackerTraits >::HitContainer Tuples
TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fastFit(const TAcc &acc, const M3xN &hits, V4 &result)
A very fast helix fit: it fits a circle by three points (first, middle and last point) and a line by ...
Definition: RiemannFit.h:385
Eigen::Matrix< double, N, 1 > VectorNd
Definition: FitUtils.h:37
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
uint32_t double double *__restrict__ float *__restrict__ double *__restrict__ pfast_fit_input
Abs< T >::type abs(const T &t)
Definition: Abs.h:22
Eigen::Matrix< double, 3, N > Matrix3xNd
Definition: HelixFit.h:30
ALPAKA_FN_ACC void operator()(TAcc const &acc, TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity, uint32_t nHits, double bField, OutputSoAView< TrackerTraits > results_view, double *__restrict__ phits, float *__restrict__ phits_ge, double *__restrict__ pfast_fit_input, riemannFit::CircleFit *__restrict__ circle_fit, uint32_t offset) const
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fromCircleToPerigee(const TAcc &acc, CircleFit &circle)
Transform circle parameter from (X0,Y0,R) to (phi,Tip,q/R) and consequently covariance matrix...
Definition: FitUtils.h:239
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 void operator()(TAcc const &acc, TupleMultiplicity< TrackerTraits > const *__restrict__ tupleMultiplicity, uint32_t nHits, double bField, double *__restrict__ phits, float *__restrict__ phits_ge, double *__restrict__ pfast_fit_input, riemannFit::CircleFit *circle_fit, uint32_t offset) const
ALPAKA_FN_ACC ALPAKA_FN_INLINE void const M3xN const V4 & fast_fit
Definition: BrokenLine.h:158
Eigen::Matrix< double, 2 *N, 2 *N > Matrix2Nd
Definition: FitUtils.h:25
#define N
Definition: blowfish.cc:9
typename TrackingRecHitSoA< TrackerTraits >::template TrackingRecHitSoALayout<>::ConstView TrackingRecHitSoAConstView
ALPAKA_FN_ACC constexpr bool once_per_grid(TAcc const &acc)
uint32_t double double *__restrict__ float *__restrict__ double *__restrict__ riemannFit::CircleFit * circle_fit
void launchRiemannKernels(const HitConstView &hv, ParamsOnDevice const *cpeParams, uint32_t nhits, uint32_t maxNumberOfTuples, Queue &queue)
ALPAKA_FN_ACC ALPAKA_FN_INLINE CircleFit circleFit(const TAcc &acc, const M2xN &hits2D, const Matrix2Nd< N > &hits_cov2D, const V4 &fast_fit, const VectorNd< N > &rad, const double bField, const bool error)
Fit a generic number of 2D points with a circle using Riemann-Chernov algorithm. Covariance matrix of...
Definition: RiemannFit.h:465
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
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
reco::TrackSoAView< TrackerTraits > OutputSoAView
constexpr uint32_t maxNumberOfConcurrentFits
Definition: HelixFit.h:21
typename reco::TrackSoA< TrackerTraits >::template Layout<>::View TrackSoAView
Definition: TracksSoA.h:45