RiemannFitOnGPU.h File Reference

#include <cstdint>
#include <cuda_runtime.h>
#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
#include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
#include "RecoPixelVertexing/PixelTrackFitting/interface/RiemannFit.h"
#include "HelixFitOnGPU.h"

Go to the source code of this file.

Typedefs
using	HitsOnGPU = TrackingRecHit2DSOAView
using	OutputSoA = pixelTrack::TrackSoA
using	Tuples = pixelTrack::HitContainer

Functions
template<int N>
	__attribute__ ((always_inline)) void kernel_FastFit(Tuples const *__restrict__ foundNtuplets
	assert (foundNtuplets)
	assert (hitsInFit<=nHits)
	assert (pfast_fit)
	assert (tupleMultiplicity)
	for (int local_idx=local_start, nt=riemannFit::maxNumberOfConcurrentFits;local_idx< nt;local_idx+=gridDim.x *blockDim.x)

Variables
const caConstants::TupleMultiplicity *__restrict__ uint32_t const HitsOnGPU *__restrict__	hhp
auto	local_start = blockIdx.x * blockDim.x + threadIdx.x
const caConstants::TupleMultiplicity *__restrict__ uint32_t	nHits
const caConstants::TupleMultiplicity *__restrict__ uint32_t const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__ uint32_t	offset
const caConstants::TupleMultiplicity *__restrict__ uint32_t const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__ double *__restrict__	pfast_fit
const caConstants::TupleMultiplicity *__restrict__ uint32_t const HitsOnGPU *__restrict__ double *__restrict__	phits
const caConstants::TupleMultiplicity *__restrict__ uint32_t const HitsOnGPU *__restrict__ double *__restrict__ float *__restrict__	phits_ge
const caConstants::TupleMultiplicity *__restrict__	tupleMultiplicity

Typedef Documentation

HitsOnGPU

using HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 17 of file RiemannFitOnGPU.h.

OutputSoA

using OutputSoA = pixelTrack::TrackSoA

Definition at line 19 of file RiemannFitOnGPU.h.

Tuples

using Tuples = pixelTrack::HitContainer

Definition at line 18 of file RiemannFitOnGPU.h.

Function Documentation

__attribute__()

template<int N>

__attribute__ ( (always_inline)  ) const

inline

Definition at line 87 of file RiemannFitOnGPU.h.

                                                  {
  assert(circle_fit);
  assert(N <= nHits);
 
  // same as above...
 
  // look in bin for this hit multiplicity
  auto local_start = blockIdx.x * blockDim.x + threadIdx.x;
  for (int local_idx = local_start, nt = riemannFit::maxNumberOfConcurrentFits; local_idx < nt;
       local_idx += gridDim.x * blockDim.x) {
    auto tuple_idx = local_idx + offset;
    if (tuple_idx >= tupleMultiplicity->size(nHits))
      break;
 
    riemannFit::Map3xNd<N> hits(phits + local_idx);
    riemannFit::Map4d fast_fit(pfast_fit_input + local_idx);
    riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
 
    riemannFit::VectorNd<N> rad = (hits.block(0, 0, 2, N).colwise().norm());
 
    riemannFit::Matrix2Nd<N> hits_cov = riemannFit::Matrix2Nd<N>::Zero();
    riemannFit::loadCovariance2D(hits_ge, hits_cov);
 
    circle_fit[local_idx] = riemannFit::circleFit(hits.block(0, 0, 2, N), hits_cov, fast_fit, rad, bField, true);
 
#ifdef RIEMANN_DEBUG
//    auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
//  printf("kernelCircleFit circle.par(0,1,2): %d %f,%f,%f\n", tkid,
//         circle_fit[local_idx].par(0), circle_fit[local_idx].par(1), circle_fit[local_idx].par(2));
#endif
  }
}

References assert(), Calorimetry_cff::bField, cms::cudacompat::blockDim, cms::cudacompat::blockIdx, riemannFit::circleFit(), cms::cudacompat::gridDim, hfClusterShapes_cfi::hits, riemannFit::loadCovariance2D(), local_start, riemannFit::maxNumberOfConcurrentFits, N, nHits, nt, offset, phits, phits_ge, cms::cudacompat::threadIdx, and tupleMultiplicity.

assert() [1/4]

assert

(

foundNtuplets

)

assert() [2/4]

assert

(

hitsInFit<=

nHits

)

Referenced by __attribute__(), and for().

assert() [3/4]

assert

(

pfast_fit

)

assert() [4/4]

assert

(

tupleMultiplicity

)

for()

for

(

int

local_idx = local_start

)

Definition at line 46 of file RiemannFitOnGPU.h.

                                            {
    auto tuple_idx = local_idx + offset;
    if (tuple_idx >= tupleMultiplicity->size(nHits))
      break;
 
    // get it from the ntuple container (one to one to helix)
    auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
    assert(tkid < foundNtuplets->nbins());
 
    assert(foundNtuplets->size(tkid) == nHits);
 
    riemannFit::Map3xNd<N> hits(phits + local_idx);
    riemannFit::Map4d fast_fit(pfast_fit + local_idx);
    riemannFit::Map6xNf<N> hits_ge(phits_ge + local_idx);
 
    // Prepare data structure
    auto const *hitId = foundNtuplets->begin(tkid);
    for (unsigned int i = 0; i < hitsInFit; ++i) {
      auto hit = hitId[i];
      // printf("Hit global: %f,%f,%f\n", hhp->xg_d[hit],hhp->yg_d[hit],hhp->zg_d[hit]);
      float ge[6];
      hhp->cpeParams()
          .detParams(hhp->detectorIndex(hit))
          .frame.toGlobal(hhp->xerrLocal(hit), 0, hhp->yerrLocal(hit), ge);
      // printf("Error: %d: %f,%f,%f,%f,%f,%f\n",hhp->detInd_d[hit],ge[0],ge[1],ge[2],ge[3],ge[4],ge[5]);
 
      hits.col(i) << hhp->xGlobal(hit), hhp->yGlobal(hit), hhp->zGlobal(hit);
      hits_ge.col(i) << ge[0], ge[1], ge[2], ge[3], ge[4], ge[5];
    }
    riemannFit::fastFit(hits, fast_fit);
 
    // no NaN here....
    assert(fast_fit(0) == fast_fit(0));
    assert(fast_fit(1) == fast_fit(1));
    assert(fast_fit(2) == fast_fit(2));
    assert(fast_fit(3) == fast_fit(3));
  }

References assert(), riemannFit::fastFit(), foundNtuplets, amptDefault_cfi::frame, hhp, hfClusterShapes_cfi::hits, mps_fire::i, LaserClient_cfi::nbins, nHits, offset, pfast_fit, phits, phits_ge, and tupleMultiplicity.

Variable Documentation

hhp

const caConstants::TupleMultiplicity* __restrict__ uint32_t const HitsOnGPU* __restrict__ hhp

Definition at line 23 of file RiemannFitOnGPU.h.

Referenced by for().

local_start

auto local_start = blockIdx.x * blockDim.x + threadIdx.x

Definition at line 39 of file RiemannFitOnGPU.h.

Referenced by __attribute__().

nHits

const caConstants::TupleMultiplicity* __restrict__ uint32_t nHits

Definition at line 23 of file RiemannFitOnGPU.h.

Referenced by __attribute__(), and for().

offset

const caConstants::TupleMultiplicity* __restrict__ uint32_t const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ uint32_t offset

Initial value:

{
  constexpr uint32_t hitsInFit = N

Definition at line 29 of file RiemannFitOnGPU.h.

Referenced by __attribute__(), and for().

pfast_fit

const caConstants::TupleMultiplicity* __restrict__ uint32_t const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ double* __restrict__ pfast_fit

Definition at line 23 of file RiemannFitOnGPU.h.

Referenced by for().

phits

const caConstants::TupleMultiplicity* __restrict__ uint32_t const HitsOnGPU* __restrict__ double* __restrict__ phits

Definition at line 23 of file RiemannFitOnGPU.h.

Referenced by __attribute__(), and for().

phits_ge

const caConstants::TupleMultiplicity* __restrict__ uint32_t const HitsOnGPU* __restrict__ double* __restrict__ float* __restrict__ phits_ge

Definition at line 23 of file RiemannFitOnGPU.h.

Referenced by __attribute__(), and for().

tupleMultiplicity

const caConstants::TupleMultiplicity* __restrict__ tupleMultiplicity

Definition at line 23 of file RiemannFitOnGPU.h.

Referenced by __attribute__(), and for().