3 #include <alpaka/alpaka.hpp> 16 template <
typename TrackerTraits>
18 template <
typename TrackerTraits>
20 template <
typename TrackerTraits>
27 template <
int N,
typename TrackerTraits>
30 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
37 double *__restrict__
phits,
52 const uint32_t
threadIdx(alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u]);
59 auto tuple_idx = local_idx +
offset;
75 for (
unsigned int i = 0;
i < hitsInFit; ++
i) {
81 hits_ge.col(
i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
94 template <
int N,
typename TrackerTraits>
97 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
102 double *__restrict__
phits,
115 auto tuple_idx = local_idx +
offset;
140 template <
int N,
typename TrackerTraits>
143 template <
typename TAcc,
typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
149 double *__restrict__
phits,
162 auto tuple_idx = local_idx +
offset;
189 printf(
"kernelLineFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
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",
210 template <
typename TrackerTraits>
214 uint32_t maxNumberOfTuples,
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);
224 auto hitsDevice = cms::alpakatools::make_device_buffer<double[]>(
226 auto hits_geDevice = cms::alpakatools::make_device_buffer<float[]>(
228 auto fast_fit_resultsDevice = cms::alpakatools::make_device_buffer<double[]>(
230 auto circle_fit_resultsDevice_holder =
235 for (uint32_t
offset = 0;
offset < maxNumberOfTuples;
offset += maxNumberOfConcurrentFits_) {
237 alpaka::exec<Acc1D>(
queue,
246 hits_geDevice.data(),
247 fast_fit_resultsDevice.data(),
250 alpaka::exec<Acc1D>(
queue,
257 hits_geDevice.data(),
258 fast_fit_resultsDevice.data(),
259 circle_fit_resultsDevice_,
262 alpaka::exec<Acc1D>(
queue,
270 hits_geDevice.data(),
271 fast_fit_resultsDevice.data(),
272 circle_fit_resultsDevice_,
276 alpaka::exec<Acc1D>(
queue,
285 hits_geDevice.data(),
286 fast_fit_resultsDevice.data(),
289 alpaka::exec<Acc1D>(
queue,
296 hits_geDevice.data(),
297 fast_fit_resultsDevice.data(),
298 circle_fit_resultsDevice_,
301 alpaka::exec<Acc1D>(
queue,
309 hits_geDevice.data(),
310 fast_fit_resultsDevice.data(),
311 circle_fit_resultsDevice_,
316 alpaka::exec<Acc1D>(
queue,
325 hits_geDevice.data(),
326 fast_fit_resultsDevice.data(),
329 alpaka::exec<Acc1D>(
queue,
336 hits_geDevice.data(),
337 fast_fit_resultsDevice.data(),
338 circle_fit_resultsDevice_,
341 alpaka::exec<Acc1D>(
queue,
349 hits_geDevice.data(),
350 fast_fit_resultsDevice.data(),
351 circle_fit_resultsDevice_,
355 alpaka::exec<Acc1D>(
queue,
364 hits_geDevice.data(),
365 fast_fit_resultsDevice.data(),
368 alpaka::exec<Acc1D>(
queue,
375 hits_geDevice.data(),
376 fast_fit_resultsDevice.data(),
377 circle_fit_resultsDevice_,
380 alpaka::exec<Acc1D>(
queue,
388 hits_geDevice.data(),
389 fast_fit_resultsDevice.data(),
390 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