3 #include <alpaka/alpaka.hpp> 15 template <
typename TrackerTraits>
17 template <
typename TrackerTraits>
19 template <
typename TrackerTraits>
26 template <
int N,
typename TrackerTraits>
29 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
36 double *__restrict__
phits,
51 const uint32_t
threadIdx(alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u]);
58 auto tuple_idx = local_idx +
offset;
74 for (
unsigned int i = 0;
i < hitsInFit; ++
i) {
80 hits_ge.col(
i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
95 template <
int N,
typename TrackerTraits>
98 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
103 double *__restrict__
phits,
116 auto tuple_idx = local_idx +
offset;
141 template <
int N,
typename TrackerTraits>
144 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
150 double *__restrict__
phits,
163 auto tuple_idx = local_idx +
offset;
190 printf(
"kernelLineFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
197 printf(
"kernelLineFit line.par(0,1): %d %f,%f\n", tkid, line_fit.par(0), line_fit.par(1));
198 printf(
"kernelLineFit chi2 cov %f/%f %e,%e,%e,%e,%e\n",
211 template <
typename TrackerTraits>
215 uint32_t maxNumberOfTuples,
220 auto numberOfBlocks = (maxNumberOfConcurrentFits_ + blockSize - 1) / blockSize;
221 const auto workDivTriplets = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks, blockSize);
222 const auto workDivQuadsPenta = cms::alpakatools::make_workdiv<Acc1D>(numberOfBlocks / 4, blockSize);
225 auto hitsDevice = cms::alpakatools::make_device_buffer<double[]>(
227 auto hits_geDevice = cms::alpakatools::make_device_buffer<float[]>(
229 auto fast_fit_resultsDevice = cms::alpakatools::make_device_buffer<double[]>(
231 auto circle_fit_resultsDevice_holder =
236 for (uint32_t
offset = 0;
offset < maxNumberOfTuples;
offset += maxNumberOfConcurrentFits_) {
238 alpaka::exec<Acc1D>(
queue,
247 hits_geDevice.data(),
248 fast_fit_resultsDevice.data(),
251 alpaka::exec<Acc1D>(
queue,
258 hits_geDevice.data(),
259 fast_fit_resultsDevice.data(),
260 circle_fit_resultsDevice_,
263 alpaka::exec<Acc1D>(
queue,
271 hits_geDevice.data(),
272 fast_fit_resultsDevice.data(),
273 circle_fit_resultsDevice_,
277 alpaka::exec<Acc1D>(
queue,
286 hits_geDevice.data(),
287 fast_fit_resultsDevice.data(),
290 alpaka::exec<Acc1D>(
queue,
297 hits_geDevice.data(),
298 fast_fit_resultsDevice.data(),
299 circle_fit_resultsDevice_,
302 alpaka::exec<Acc1D>(
queue,
310 hits_geDevice.data(),
311 fast_fit_resultsDevice.data(),
312 circle_fit_resultsDevice_,
317 alpaka::exec<Acc1D>(
queue,
326 hits_geDevice.data(),
327 fast_fit_resultsDevice.data(),
330 alpaka::exec<Acc1D>(
queue,
337 hits_geDevice.data(),
338 fast_fit_resultsDevice.data(),
339 circle_fit_resultsDevice_,
342 alpaka::exec<Acc1D>(
queue,
350 hits_geDevice.data(),
351 fast_fit_resultsDevice.data(),
352 circle_fit_resultsDevice_,
356 alpaka::exec<Acc1D>(
queue,
365 hits_geDevice.data(),
366 fast_fit_resultsDevice.data(),
369 alpaka::exec<Acc1D>(
queue,
376 hits_geDevice.data(),
377 fast_fit_resultsDevice.data(),
378 circle_fit_resultsDevice_,
381 alpaka::exec<Acc1D>(
queue,
389 hits_geDevice.data(),
390 fast_fit_resultsDevice.data(),
391 circle_fit_resultsDevice_,
constexpr DetParams const &__restrict__ detParams(int i) const
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...
ALPAKA_FN_ACC void loadCovariance2D(const TAcc &acc, M6xNf const &ge, M2Nd &hits_cov)
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
Vector3d par
parameter: (X0,Y0,R)
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 ...
Eigen::Matrix< double, N, 1 > VectorNd
TupleMultiplicity< TrackerTraits > const *__restrict__ TrackingRecHitSoAConstView< TrackerTraits > TrackerTraits::tindex_type *__restrict__ double *__restrict__ phits
Eigen::Map< Vector4d, 0, Eigen::InnerStride< stride > > Map4d
Eigen::Map< Matrix3xNd< N >, 0, Eigen::Stride< 3 *stride, stride > > Map3xNd
uint32_t double double *__restrict__ float *__restrict__ double *__restrict__ pfast_fit_input
Abs< T >::type abs(const T &t)
Eigen::Matrix< double, 3, N > Matrix3xNd
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...
double OutputSoAView< TrackerTraits > results_view
Eigen::Map< Matrix6xNf< N >, 0, Eigen::Stride< 6 *stride, stride > > Map6xNf
ALPAKA_FN_ACC ALPAKA_FN_INLINE void uint32_t const uint32_t CACellT< TrackerTraits > uint32_t CellNeighborsVector< TrackerTraits > CellTracksVector< TrackerTraits > HitsConstView< TrackerTraits > hh
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
Eigen::Matrix< double, 2 *N, 2 *N > Matrix2Nd
typename TrackingRecHitSoA< TrackerTraits >::template TrackingRecHitSoALayout<>::ConstView TrackingRecHitSoAConstView
auto const & foundNtuplets
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...
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
ALPAKA_ASSERT_ACC(offsets)
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
typename reco::TrackSoA< TrackerTraits >::template Layout<>::View TrackSoAView