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 (hitsInFit<=nHits)
	assert (pfast_fit)
	assert (foundNtuplets)
	assert (tupleMultiplicity)
	for (int local_idx=local_start, nt=riemannFit::maxNumberOfConcurrentFits;local_idx< nt;local_idx+=gridDim.x *blockDim.x)

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

Typedef Documentation

using HitsOnGPU = TrackingRecHit2DSOAView

Definition at line 17 of file RiemannFitOnGPU.h.

using OutputSoA = pixelTrack::TrackSoA

Definition at line 19 of file RiemannFitOnGPU.h.

using Tuples = pixelTrack::HitContainer

Definition at line 18 of file RiemannFitOnGPU.h.

Function Documentation

template<int N>

__attribute__ ( (always_inline)  ) const

inline

assert

(

hitsInFit<=

nHits

)

assert

(

pfast_fit

)

assert

(

foundNtuplets

)

assert

(

tupleMultiplicity

)

for

(

int

local_idx = local_start

)

Definition at line 46 of file RiemannFitOnGPU.h.

References cms::cuda::assert(), riemannFit::fastFit(), mps_fire::i, nHits, and hltrates_dqm_sourceclient-live_cfg::offset.

                                            {
    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->nOnes());

    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));
  }

Variable Documentation

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

Definition at line 23 of file RiemannFitOnGPU.h.

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

Definition at line 39 of file RiemannFitOnGPU.h.

caConstants::TupleMultiplicity const* __restrict__ uint32_t nHits

Definition at line 23 of file RiemannFitOnGPU.h.

caConstants::TupleMultiplicity const* __restrict__ uint32_t HitsOnGPU const* __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.

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

Definition at line 23 of file RiemannFitOnGPU.h.

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

Definition at line 23 of file RiemannFitOnGPU.h.

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

Definition at line 23 of file RiemannFitOnGPU.h.

caConstants::TupleMultiplicity const* __restrict__ tupleMultiplicity

Definition at line 23 of file RiemannFitOnGPU.h.