9 #include <cuda_runtime.h>
32 double *__restrict__
phits,
38 constexpr uint32_t hitsInFit =
N;
40 assert(hitsInFit <= nHitsL);
50 int totTK = tupleMultiplicity->end(nHitsH) - tupleMultiplicity->begin(nHitsL);
51 assert(totTK <=
int(tupleMultiplicity->size()));
54 #ifdef BROKENLINE_DEBUG
56 printf(
"%d total Ntuple\n", tupleMultiplicity->size());
57 printf(
"%d Ntuple of size %d/%d for %d hits to fit\n", totTK, nHitsL, nHitsH, hitsInFit);
63 int tuple_idx = local_idx +
offset;
64 if (tuple_idx >= totTK) {
69 auto tkid = *(tupleMultiplicity->begin(nHitsL) + tuple_idx);
70 assert(tkid < foundNtuplets->nOnes());
72 ptkids[local_idx] = tkid;
74 auto nHits = foundNtuplets->size(tkid);
87 bool dump = (foundNtuplets->size(tkid) == 5 && 0 ==
atomicAdd(&done, 1));
91 auto const *hitId = foundNtuplets->begin(tkid);
96 auto dx = hhp->xGlobal(hitId[hitsInFit - 1]) - hhp->xGlobal(hitId[0]);
97 auto dy = hhp->yGlobal(hitId[hitsInFit - 1]) - hhp->yGlobal(hitId[0]);
98 auto dz = hhp->zGlobal(hitId[hitsInFit - 1]) - hhp->zGlobal(hitId[0]);
104 for (uint32_t
i = 0;
i < hitsInFit; ++
i) {
105 int j = int(n + 0.5
f);
106 if (hitsInFit - 1 ==
i)
114 auto const &dp = hhp->cpeParams().detParams(hhp->detectorIndex(
hit));
117 assert(qbin >= 0 && qbin < 5);
120 dp.frame.rotation().multiply(
dx,
dy,
dz, ux, uy, uz);
127 bin = std::clamp(bin, low_value, high_value);
128 float yerr = dp.sigmay[
bin] * 1.e-4
f;
129 yerr *= dp.yfact[qbin];
132 yerr = nok ? hhp->yerrLocal(
hit) : yerr;
133 dp.frame.toGlobal(hhp->xerrLocal(
hit), 0, yerr, ge);
136 .detParams(hhp->detectorIndex(
hit))
137 .frame.toGlobal(hhp->xerrLocal(
hit), 0, hhp->yerrLocal(
hit), ge);
141 bool dump = foundNtuplets->size(tkid) == 5;
143 printf(
"Track id %d %d Hit %d on %d\nGlobal: hits.col(%d) << %f,%f,%f\n",
147 hhp->detectorIndex(
hit),
152 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]);
156 hits.col(
i) << hhp->xGlobal(
hit), hhp->yGlobal(
hit), hhp->zGlobal(
hit);
157 hits_ge.col(
i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
162 assert(fast_fit(0) == fast_fit(0));
163 assert(fast_fit(1) == fast_fit(1));
164 assert(fast_fit(2) == fast_fit(2));
165 assert(fast_fit(3) == fast_fit(3));
174 double *__restrict__ phits,
175 float *__restrict__ phits_ge,
176 double *__restrict__ pfast_fit) {
189 auto tkid = ptkids[local_idx];
206 results->
stateAtBS.copyFromCircle(circle.
par, circle.
cov, line.
par, line.
cov, 1.f /
float(bField), tkid);
207 results->
pt(tkid) = float(bField) / float(
std::abs(circle.
par(2)));
208 results->
eta(tkid) = asinhf(line.
par(0));
209 results->
chi2(tkid) = (circle.
chi2 + line.
chi2) / (2 *
N - 5);
211 #ifdef BROKENLINE_DEBUG
212 if (!(circle.
chi2 >= 0) || !(line.
chi2 >= 0))
214 printf(
"kernelBLFit size %d for %d hits circle.par(0,1,2): %d %f,%f,%f\n",
221 printf(
"kernelBLHits line.par(0,1): %d %f,%f\n", tkid, line.
par(0), line.
par(1));
222 printf(
"kernelBLHits chi2 cov %f/%f %e,%e,%e,%e,%e\n",
__host__ __device__ void lineFit(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...
caConstants::tindex_type tindex_type
constexpr uint16_t pixelPitchY
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ tindex_type *__restrict__ ptkids
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ tindex_type *__restrict__ double *__restrict__ float *__restrict__ phits_ge
constexpr uint32_t maxTuples
constexpr int kGenErrorQBins
Vector3d par
parameter: (X0,Y0,R)
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t nHitsL
eigenSoA::ScalarSoA< float, S > pt
__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 ...
eigenSoA::ScalarSoA< float, S > chi2
auto const & foundNtuplets
Eigen::Map< Matrix3xNd< N >, 0, Eigen::Stride< 3 *stride, stride > > Map3xNd
caConstants::TupleMultiplicity const *__restrict__ tupleMultiplicity
printf("params %d %f %f %f\n", minT, eps, errmax, chi2max)
TrajectoryStateSoAT< S > stateAtBS
__host__ __device__ void fastFit(const M3xN &hits, V4 &result)
A very fast helix fit.
Abs< T >::type abs(const T &t)
TrackSoA::HitContainer HitContainer
TrackSoAHeterogeneousT< maxNumber()> TrackSoA
Eigen::Map< Matrix6xNf< N >, 0, Eigen::Stride< 6 *stride, stride > > Map6xNf
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ pfast_fit
constexpr auto invalidTkId
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ hhp
Eigen::Map< Vector4d, 0, Eigen::InnerStride< stride > > Map4d
caConstants::TupleMultiplicity const CAHitNtupletGeneratorKernelsGPU::HitToTuple const cms::cuda::AtomicPairCounter GPUCACell const *__restrict__ uint32_t const *__restrict__ gpuPixelDoublets::CellNeighborsVector const gpuPixelDoublets::CellTracksVector const GPUCACell::OuterHitOfCell const int32_t nHits
char data[epos_bytes_allocation]
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ tindex_type *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t uint32_t nHitsH
Vector2d par
(cotan(theta),Zip)
__host__ __device__ void circleFit(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...
eigenSoA::ScalarSoA< float, S > eta
data needed for the Broken Line fit procedure.
T1 atomicAdd(T1 *a, T2 b)
constexpr int kNumErrorBins
tuple dump
OutputFilePath = cms.string('/tmp/zhokin/'), OutputFileExt = cms.string(''),.
caConstants::TupleMultiplicity const *__restrict__ HitsOnGPU const *__restrict__ tindex_type *__restrict__ double *__restrict__ phits
constexpr uint32_t maxNumberOfConcurrentFits