CMS 3D CMS Logo
Namespaces | Classes | Typedefs | Enumerations | Functions | Variables
ecal::multifit Namespace Reference

Namespaces

 v1
 

Classes

struct  ConditionsProducts
 
struct  conf_data
 
struct  ConfigurationParameters
 
struct  EventDataForScratchGPU
 
struct  EventInputDataGPU
 
struct  EventOutputDataGPU
 
struct  xyz
 

Typedefs

typedef Eigen::Matrix< char, Eigen::Dynamic, 1, 0, PulseVectorSize, 1 > BXVector
 
typedef Eigen::Matrix< char, SampleVectorSize, 1 > BXVectorType
 
using data_type = ::ecal::reco::ComputationScalarType
 
typedef Eigen::Matrix< data_type, FullSampleVectorSize, FullSampleVectorSize > FullSampleMatrix
 
typedef Eigen::Matrix< data_type, FullSampleVectorSize, 1 > FullSampleVector
 
using PermutationMatrix = Eigen::PermutationMatrix< SampleMatrix::RowsAtCompileTime >
 
typedef Eigen::LDLT< PulseMatrix > PulseDecompLDLT
 
typedef Eigen::LLT< PulseMatrix > PulseDecompLLT
 
typedef Eigen::Matrix< data_type, Eigen::Dynamic, Eigen::Dynamic, 0, PulseVectorSize, PulseVectorSize > PulseMatrix
 
typedef Eigen::Matrix< data_type, SampleVectorSize, SampleVectorSize > PulseMatrixType
 
typedef Eigen::Matrix< data_type, Eigen::Dynamic, 1, 0, PulseVectorSize, 1 > PulseVector
 
typedef Eigen::LLT< SampleMatrix > SampleDecompLLT
 
typedef Eigen::LLT< SampleMatrixD > SampleDecompLLTD
 
typedef Eigen::Matrix< char, SampleVectorSize, 1 > SampleGainVector
 
typedef Eigen::Matrix< data_type, SampleVectorSize, SampleVectorSize > SampleMatrix
 
using SampleMatrixD = Eigen::Matrix< double, SampleVectorSize, SampleVectorSize >
 
typedef std::array< SampleMatrixD, NGains > SampleMatrixGainArray
 
typedef Eigen::Matrix< data_type, SampleVectorSize, Eigen::Dynamic, 0, SampleVectorSize, PulseVectorSize > SamplePulseMatrix
 
typedef Eigen::Matrix< data_type, SampleVectorSize, 1 > SampleVector
 
typedef Eigen::Matrix< data_type, 1, 1 > SingleMatrix
 
typedef Eigen::Matrix< data_type, 1, 1 > SingleVector
 

Enumerations

enum  MinimizationState : char { MinimizationState::NotFinished = 0, MinimizationState::Finished = 1, MinimizationState::Precomputed = 2 }
 
enum  TimeComputationState : char { TimeComputationState::NotFinished = 0, TimeComputationState::Finished = 1 }
 

Functions

void entryPoint (EventInputDataGPU const &, EventOutputDataGPU &, EventDataForScratchGPU &, ConditionsProducts const &, ConfigurationParameters const &, cudaStream_t)
 
template<typename EigenM >
constexpr auto getLength () -> uint32_t
 
__global__ void kernel_permute_results (SampleVector *amplitudes, BXVectorType const *activeBXs, ::ecal::reco::StorageScalarType *energies, char const *acState, int const nchannels)
 
__global__ void kernel_prep_1d_and_initialize (EcalPulseShape const *shapes_in, uint16_t const *digis_in_eb, uint32_t const *dids_eb, uint16_t const *digis_in_ee, uint32_t const *dids_ee, SampleVector *amplitudes, SampleVector *amplitudesForMinimizationEB, SampleVector *amplitudesForMinimizationEE, SampleGainVector *gainsNoise, float const *mean_x1, float const *mean_x12, float const *rms_x12, float const *mean_x6, float const *gain6Over1, float const *gain12Over6, bool *hasSwitchToGain6, bool *hasSwitchToGain1, bool *isSaturated, ::ecal::reco::StorageScalarType *energiesEB, ::ecal::reco::StorageScalarType *energiesEE, ::ecal::reco::StorageScalarType *chi2EB, ::ecal::reco::StorageScalarType *chi2EE, ::ecal::reco::StorageScalarType *pedestalEB, ::ecal::reco::StorageScalarType *pedestalEE, uint32_t *dids_outEB, uint32_t *dids_outEE, uint32_t *flagsEB, uint32_t *flagsEE, char *acState, BXVectorType *bxs, uint32_t const offsetForHashes, uint32_t const offsetForInputs, bool const gainSwitchUseMaxSampleEB, bool const gainSwitchUseMaxSampleEE, int const nchannels)
 
__global__ void kernel_prep_2d (SampleGainVector const *gainNoise, uint32_t const *dids_eb, uint32_t const *dids_ee, float const *rms_x12, float const *rms_x6, float const *rms_x1, float const *gain12Over6, float const *gain6Over1, double const *G12SamplesCorrelationEB, double const *G6SamplesCorrelationEB, double const *G1SamplesCorrelationEB, double const *G12SamplesCorrelationEE, double const *G6SamplesCorrelationEE, double const *G1SamplesCorrelationEE, SampleMatrix *noisecov, PulseMatrixType *pulse_matrix, EcalPulseShape const *pulse_shape, bool const *hasSwitchToGain6, bool const *hasSwitchToGain1, bool const *isSaturated, uint32_t const offsetForHashes, uint32_t const offsetForInputs)
 
__global__ void kernel_time_computation_init (uint16_t const *digis_eb, uint32_t const *dids_eb, uint16_t const *digis_ee, uint32_t const *dids_ee, float const *rms_x12, float const *rms_x6, float const *rms_x1, float const *mean_x12, float const *mean_x6, float const *mean_x1, float const *gain12Over6, float const *gain6Over1, SampleVector::Scalar *sample_values, SampleVector::Scalar *sample_value_errors, SampleVector::Scalar *ampMaxError, bool *useless_sample_values, char *pedestal_nums, uint32_t const offsetForHashes, uint32_t const offsetForInputs, unsigned int const sample_maskEB, unsigned int const sample_maskEE, int nchannels)
 
__global__ void kernel_time_compute_ampl (SampleVector::Scalar const *sample_values, SampleVector::Scalar const *sample_value_errors, uint32_t const *dids_eb, uint32_t const *dids_ed, bool const *useless_samples, SampleVector::Scalar const *g_timeMax, SampleVector::Scalar const *amplitudeFitParametersEB, SampleVector::Scalar const *amplitudeFitParametersEE, SampleVector::Scalar *g_amplitudeMax, int const nchannels, uint32_t const offsetForInputs)
 
__global__ void kernel_time_compute_findamplchi2_and_finish (SampleVector::Scalar const *sample_values, SampleVector::Scalar const *sample_value_errors, uint32_t const *dids_eb, uint32_t const *dids_ee, bool const *useless_samples, SampleVector::Scalar const *g_tMaxAlphaBeta, SampleVector::Scalar const *g_tMaxErrorAlphaBeta, SampleVector::Scalar const *g_accTimeMax, SampleVector::Scalar const *g_accTimeWgt, ConfigurationParameters::type const *amplitudeFitParametersEB, ConfigurationParameters::type const *amplitudeFitParametersEE, SampleVector::Scalar const *sumAAsNullHypot, SampleVector::Scalar const *sum0sNullHypot, SampleVector::Scalar const *chi2sNullHypot, TimeComputationState *g_state, SampleVector::Scalar *g_ampMaxAlphaBeta, SampleVector::Scalar *g_ampMaxError, SampleVector::Scalar *g_timeMax, SampleVector::Scalar *g_timeError, int const nchannels, uint32_t const offsetForInputs)
 
__global__ void kernel_time_compute_fixMGPAslew (uint16_t const *digis_eb, uint16_t const *digis_ee, SampleVector::Scalar *sample_values, SampleVector::Scalar *sample_value_errors, bool *useless_sample_values, unsigned int const sample_mask, int const nchannels, uint32_t const offsetForInputs)
 
__global__ void kernel_time_compute_makeratio (SampleVector::Scalar const *sample_values, SampleVector::Scalar const *sample_value_errors, uint32_t const *dids_eb, uint32_t const *dids_ee, bool const *useless_sample_values, char const *pedestal_nums, ConfigurationParameters::type const *amplitudeFitParametersEB, ConfigurationParameters::type const *amplitudeFitParametersEE, ConfigurationParameters::type const *timeFitParametersEB, ConfigurationParameters::type const *timeFitParametersEE, SampleVector::Scalar const *sumAAsNullHypot, SampleVector::Scalar const *sum0sNullHypot, SampleVector::Scalar *tMaxAlphaBetas, SampleVector::Scalar *tMaxErrorAlphaBetas, SampleVector::Scalar *g_accTimeMax, SampleVector::Scalar *g_accTimeWgt, TimeComputationState *g_state, unsigned int const timeFitParameters_sizeEB, unsigned int const timeFitParameters_sizeEE, ConfigurationParameters::type const timeFitLimits_firstEB, ConfigurationParameters::type const timeFitLimits_firstEE, ConfigurationParameters::type const timeFitLimits_secondEB, ConfigurationParameters::type const timeFitLimits_secondEE, int const nchannels, uint32_t const offsetForInputs)
 
__global__ void kernel_time_compute_nullhypot (SampleVector::Scalar const *sample_values, SampleVector::Scalar const *sample_value_errors, bool const *useless_sample_values, SampleVector::Scalar *chi2s, SampleVector::Scalar *sum0s, SampleVector::Scalar *sumAAs, int const nchannels)
 
__global__ void kernel_time_correction_and_finalize (::ecal::reco::StorageScalarType const *g_amplitudeEB, ::ecal::reco::StorageScalarType const *g_amplitudeEE, uint16_t const *digis_eb, uint32_t const *dids_eb, uint16_t const *digis_ee, uint32_t const *dids_ee, float const *amplitudeBinsEB, float const *amplitudeBinsEE, float const *shiftBinsEB, float const *shiftBinsEE, SampleVector::Scalar const *g_timeMax, SampleVector::Scalar const *g_timeError, float const *g_rms_x12, float const *timeCalibConstant, ::ecal::reco::StorageScalarType *g_jitterEB, ::ecal::reco::StorageScalarType *g_jitterEE, ::ecal::reco::StorageScalarType *g_jitterErrorEB, ::ecal::reco::StorageScalarType *g_jitterErrorEE, uint32_t *flagsEB, uint32_t *flagsEE, int const amplitudeBinsSizeEB, int const amplitudeBinsSizeEE, ConfigurationParameters::type const timeConstantTermEB, ConfigurationParameters::type const timeConstantTermEE, float const offsetTimeValueEB, float const offsetTimeValueEE, ConfigurationParameters::type const timeNconstEB, ConfigurationParameters::type const timeNconstEE, ConfigurationParameters::type const amplitudeThresholdEB, ConfigurationParameters::type const amplitudeThresholdEE, ConfigurationParameters::type const outOfTimeThreshG12pEB, ConfigurationParameters::type const outOfTimeThreshG12pEE, ConfigurationParameters::type const outOfTimeThreshG12mEB, ConfigurationParameters::type const outOfTimeThreshG12mEE, ConfigurationParameters::type const outOfTimeThreshG61pEB, ConfigurationParameters::type const outOfTimeThreshG61pEE, ConfigurationParameters::type const outOfTimeThreshG61mEB, ConfigurationParameters::type const outOfTimeThreshG61mEE, uint32_t const offsetForHashes, uint32_t const offsetForInputs, int const nchannels)
 

Variables

constexpr int FullSampleVectorSize = 19
 
constexpr int NGains = 3
 
constexpr int PulseVectorSize = 12
 
constexpr int SampleVectorSize = 10
 

Typedef Documentation

◆ BXVector

typedef Eigen::Matrix<char, Eigen::Dynamic, 1, 0, PulseVectorSize, 1> ecal::multifit::BXVector

Definition at line 26 of file EigenMatrixTypes_gpu.h.

◆ BXVectorType

typedef Eigen::Matrix<char, SampleVectorSize, 1> ecal::multifit::BXVectorType

Definition at line 20 of file EigenMatrixTypes_gpu.h.

◆ data_type

using ecal::multifit::data_type = typedef ::ecal::reco::ComputationScalarType

Definition at line 17 of file EigenMatrixTypes_gpu.h.

◆ FullSampleMatrix

typedef Eigen::Matrix<data_type, FullSampleVectorSize, FullSampleVectorSize> ecal::multifit::FullSampleMatrix

Definition at line 29 of file EigenMatrixTypes_gpu.h.

◆ FullSampleVector

typedef Eigen::Matrix<data_type, FullSampleVectorSize, 1> ecal::multifit::FullSampleVector

Definition at line 24 of file EigenMatrixTypes_gpu.h.

◆ PermutationMatrix

using ecal::multifit::PermutationMatrix = typedef Eigen::PermutationMatrix<SampleMatrix::RowsAtCompileTime>

Definition at line 43 of file EigenMatrixTypes_gpu.h.

◆ PulseDecompLDLT

typedef Eigen::LDLT<PulseMatrix> ecal::multifit::PulseDecompLDLT

Definition at line 36 of file EigenMatrixTypes_gpu.h.

◆ PulseDecompLLT

typedef Eigen::LLT<PulseMatrix> ecal::multifit::PulseDecompLLT

Definition at line 35 of file EigenMatrixTypes_gpu.h.

◆ PulseMatrix

typedef Eigen::Matrix<data_type, Eigen::Dynamic, Eigen::Dynamic, 0, PulseVectorSize, PulseVectorSize> ecal::multifit::PulseMatrix

Definition at line 30 of file EigenMatrixTypes_gpu.h.

◆ PulseMatrixType

typedef Eigen::Matrix<data_type, SampleVectorSize, SampleVectorSize> ecal::multifit::PulseMatrixType

Definition at line 19 of file EigenMatrixTypes_gpu.h.

◆ PulseVector

typedef Eigen::Matrix<data_type, Eigen::Dynamic, 1, 0, PulseVectorSize, 1> ecal::multifit::PulseVector

Definition at line 25 of file EigenMatrixTypes_gpu.h.

◆ SampleDecompLLT

typedef Eigen::LLT<SampleMatrix> ecal::multifit::SampleDecompLLT

Definition at line 33 of file EigenMatrixTypes_gpu.h.

◆ SampleDecompLLTD

typedef Eigen::LLT<SampleMatrixD> ecal::multifit::SampleDecompLLTD

Definition at line 34 of file EigenMatrixTypes_gpu.h.

◆ SampleGainVector

typedef Eigen::Matrix<char, SampleVectorSize, 1> ecal::multifit::SampleGainVector

Definition at line 27 of file EigenMatrixTypes_gpu.h.

◆ SampleMatrix

typedef Eigen::Matrix<data_type, SampleVectorSize, SampleVectorSize> ecal::multifit::SampleMatrix

Definition at line 28 of file EigenMatrixTypes_gpu.h.

◆ SampleMatrixD

using ecal::multifit::SampleMatrixD = typedef Eigen::Matrix<double, SampleVectorSize, SampleVectorSize>

Definition at line 21 of file EigenMatrixTypes_gpu.h.

◆ SampleMatrixGainArray

typedef std::array<SampleMatrixD, NGains> ecal::multifit::SampleMatrixGainArray

Definition at line 41 of file EigenMatrixTypes_gpu.h.

◆ SamplePulseMatrix

typedef Eigen::Matrix<data_type, SampleVectorSize, Eigen::Dynamic, 0, SampleVectorSize, PulseVectorSize> ecal::multifit::SamplePulseMatrix

Definition at line 32 of file EigenMatrixTypes_gpu.h.

◆ SampleVector

typedef Eigen::Matrix<data_type, SampleVectorSize, 1> ecal::multifit::SampleVector

Definition at line 23 of file EigenMatrixTypes_gpu.h.

◆ SingleMatrix

typedef Eigen::Matrix<data_type, 1, 1> ecal::multifit::SingleMatrix

Definition at line 38 of file EigenMatrixTypes_gpu.h.

◆ SingleVector

typedef Eigen::Matrix<data_type, 1, 1> ecal::multifit::SingleVector

Definition at line 39 of file EigenMatrixTypes_gpu.h.

Enumeration Type Documentation

◆ MinimizationState

enum ecal::multifit::MinimizationState : char
strong
Enumerator
NotFinished 
Finished 
Precomputed 

Definition at line 52 of file DeclsForKernels.h.

52  : char {
53  NotFinished = 0,
54  Finished = 1,
55  Precomputed = 2,
56  };
ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::TimeComputationState::Finished
ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::TimeComputationState::NotFinished
ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::MinimizationState::Precomputed

◆ TimeComputationState

enum ecal::multifit::TimeComputationState : char
strong
Enumerator
NotFinished 
Finished 

Definition at line 51 of file DeclsForKernels.h.

51 : char { NotFinished = 0, Finished = 1 };
ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::TimeComputationState::Finished
ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::TimeComputationState::NotFinished

Function Documentation

◆ entryPoint()

void ecal::multifit::entryPoint ( EventInputDataGPU const &  ,
EventOutputDataGPU &  ,
EventDataForScratchGPU &  ,
ConditionsProducts const &  ,
ConfigurationParameters const &  ,
cudaStream_t   
)

Referenced by EcalUncalibRecHitProducerGPU::produce().

◆ getLength()

template<typename EigenM >
constexpr auto ecal::multifit::getLength ( ) -> uint32_t

Definition at line 132 of file DeclsForKernels.h.

132  {
133  return EigenM::RowsAtCompileTime * EigenM::ColsAtCompileTime;
134  }

◆ kernel_permute_results()

__global__ void ecal::multifit::kernel_permute_results ( SampleVector *  amplitudes,
BXVectorType const *  activeBXs,
::ecal::reco::StorageScalarType *  energies,
char const *  acState,
int const  nchannels 
)

◆ kernel_prep_1d_and_initialize()

__global__ void ecal::multifit::kernel_prep_1d_and_initialize ( EcalPulseShape const *  shapes_in,
uint16_t const *  digis_in_eb,
uint32_t const *  dids_eb,
uint16_t const *  digis_in_ee,
uint32_t const *  dids_ee,
SampleVector *  amplitudes,
SampleVector *  amplitudesForMinimizationEB,
SampleVector *  amplitudesForMinimizationEE,
SampleGainVector *  gainsNoise,
float const *  mean_x1,
float const *  mean_x12,
float const *  rms_x12,
float const *  mean_x6,
float const *  gain6Over1,
float const *  gain12Over6,
bool *  hasSwitchToGain6,
bool *  hasSwitchToGain1,
bool *  isSaturated,
::ecal::reco::StorageScalarType *  energiesEB,
::ecal::reco::StorageScalarType *  energiesEE,
::ecal::reco::StorageScalarType *  chi2EB,
::ecal::reco::StorageScalarType *  chi2EE,
::ecal::reco::StorageScalarType *  pedestalEB,
::ecal::reco::StorageScalarType *  pedestalEE,
uint32_t *  dids_outEB,
uint32_t *  dids_outEE,
uint32_t *  flagsEB,
uint32_t *  flagsEE,
char *  acState,
BXVectorType *  bxs,
uint32_t const  offsetForHashes,
uint32_t const  offsetForInputs,
bool const  gainSwitchUseMaxSampleEB,
bool const  gainSwitchUseMaxSampleEE,
int const  nchannels 
)

assume kernel launch configuration is (MAXSAMPLES * nchannels, blocks) TODO: is there a point to split this kernel further to separate reductions

◆ kernel_prep_2d()

__global__ void ecal::multifit::kernel_prep_2d ( SampleGainVector const *  gainNoise,
uint32_t const *  dids_eb,
uint32_t const *  dids_ee,
float const *  rms_x12,
float const *  rms_x6,
float const *  rms_x1,
float const *  gain12Over6,
float const *  gain6Over1,
double const *  G12SamplesCorrelationEB,
double const *  G6SamplesCorrelationEB,
double const *  G1SamplesCorrelationEB,
double const *  G12SamplesCorrelationEE,
double const *  G6SamplesCorrelationEE,
double const *  G1SamplesCorrelationEE,
SampleMatrix *  noisecov,
PulseMatrixType *  pulse_matrix,
EcalPulseShape const *  pulse_shape,
bool const *  hasSwitchToGain6,
bool const *  hasSwitchToGain1,
bool const *  isSaturated,
uint32_t const  offsetForHashes,
uint32_t const  offsetForInputs 
)

assume kernel launch configuration is ([MAXSAMPLES, MAXSAMPLES], nchannels)

◆ kernel_time_computation_init()

__global__ void ecal::multifit::kernel_time_computation_init ( uint16_t const *  digis_eb,
uint32_t const *  dids_eb,
uint16_t const *  digis_ee,
uint32_t const *  dids_ee,
float const *  rms_x12,
float const *  rms_x6,
float const *  rms_x1,
float const *  mean_x12,
float const *  mean_x6,
float const *  mean_x1,
float const *  gain12Over6,
float const *  gain6Over1,
SampleVector::Scalar *  sample_values,
SampleVector::Scalar *  sample_value_errors,
SampleVector::Scalar *  ampMaxError,
bool *  useless_sample_values,
char *  pedestal_nums,
uint32_t const  offsetForHashes,
uint32_t const  offsetForInputs,
unsigned int const  sample_maskEB,
unsigned int const  sample_maskEE,
int  nchannels 
)

◆ kernel_time_compute_ampl()

__global__ void ecal::multifit::kernel_time_compute_ampl ( SampleVector::Scalar const *  sample_values,
SampleVector::Scalar const *  sample_value_errors,
uint32_t const *  dids_eb,
uint32_t const *  dids_ed,
bool const *  useless_samples,
SampleVector::Scalar const *  g_timeMax,
SampleVector::Scalar const *  amplitudeFitParametersEB,
SampleVector::Scalar const *  amplitudeFitParametersEE,
SampleVector::Scalar *  g_amplitudeMax,
int const  nchannels,
uint32_t const  offsetForInputs 
)

◆ kernel_time_compute_findamplchi2_and_finish()

__global__ void ecal::multifit::kernel_time_compute_findamplchi2_and_finish ( SampleVector::Scalar const *  sample_values,
SampleVector::Scalar const *  sample_value_errors,
uint32_t const *  dids_eb,
uint32_t const *  dids_ee,
bool const *  useless_samples,
SampleVector::Scalar const *  g_tMaxAlphaBeta,
SampleVector::Scalar const *  g_tMaxErrorAlphaBeta,
SampleVector::Scalar const *  g_accTimeMax,
SampleVector::Scalar const *  g_accTimeWgt,
ConfigurationParameters::type const *  amplitudeFitParametersEB,
ConfigurationParameters::type const *  amplitudeFitParametersEE,
SampleVector::Scalar const *  sumAAsNullHypot,
SampleVector::Scalar const *  sum0sNullHypot,
SampleVector::Scalar const *  chi2sNullHypot,
TimeComputationState *  g_state,
SampleVector::Scalar *  g_ampMaxAlphaBeta,
SampleVector::Scalar *  g_ampMaxError,
SampleVector::Scalar *  g_timeMax,
SampleVector::Scalar *  g_timeError,
int const  nchannels,
uint32_t const  offsetForInputs 
)

launch ctx parameters are 10 threads per channel, N channels per block, Y blocks TODO: do we need to keep the state around or can be removed?!

◆ kernel_time_compute_fixMGPAslew()

__global__ void ecal::multifit::kernel_time_compute_fixMGPAslew ( uint16_t const *  digis_eb,
uint16_t const *  digis_ee,
SampleVector::Scalar *  sample_values,
SampleVector::Scalar *  sample_value_errors,
bool *  useless_sample_values,
unsigned int const  sample_mask,
int const  nchannels,
uint32_t const  offsetForInputs 
)

◆ kernel_time_compute_makeratio()

__global__ void ecal::multifit::kernel_time_compute_makeratio ( SampleVector::Scalar const *  sample_values,
SampleVector::Scalar const *  sample_value_errors,
uint32_t const *  dids_eb,
uint32_t const *  dids_ee,
bool const *  useless_sample_values,
char const *  pedestal_nums,
ConfigurationParameters::type const *  amplitudeFitParametersEB,
ConfigurationParameters::type const *  amplitudeFitParametersEE,
ConfigurationParameters::type const *  timeFitParametersEB,
ConfigurationParameters::type const *  timeFitParametersEE,
SampleVector::Scalar const *  sumAAsNullHypot,
SampleVector::Scalar const *  sum0sNullHypot,
SampleVector::Scalar *  tMaxAlphaBetas,
SampleVector::Scalar *  tMaxErrorAlphaBetas,
SampleVector::Scalar *  g_accTimeMax,
SampleVector::Scalar *  g_accTimeWgt,
TimeComputationState *  g_state,
unsigned int const  timeFitParameters_sizeEB,
unsigned int const  timeFitParameters_sizeEE,
ConfigurationParameters::type const  timeFitLimits_firstEB,
ConfigurationParameters::type const  timeFitLimits_firstEE,
ConfigurationParameters::type const  timeFitLimits_secondEB,
ConfigurationParameters::type const  timeFitLimits_secondEE,
int const  nchannels,
uint32_t const  offsetForInputs 
)

◆ kernel_time_compute_nullhypot()

__global__ void ecal::multifit::kernel_time_compute_nullhypot ( SampleVector::Scalar const *  sample_values,
SampleVector::Scalar const *  sample_value_errors,
bool const *  useless_sample_values,
SampleVector::Scalar *  chi2s,
SampleVector::Scalar *  sum0s,
SampleVector::Scalar *  sumAAs,
int const  nchannels 
)

◆ kernel_time_correction_and_finalize()

__global__ void ecal::multifit::kernel_time_correction_and_finalize ( ::ecal::reco::StorageScalarType const *  g_amplitudeEB,
::ecal::reco::StorageScalarType const *  g_amplitudeEE,
uint16_t const *  digis_eb,
uint32_t const *  dids_eb,
uint16_t const *  digis_ee,
uint32_t const *  dids_ee,
float const *  amplitudeBinsEB,
float const *  amplitudeBinsEE,
float const *  shiftBinsEB,
float const *  shiftBinsEE,
SampleVector::Scalar const *  g_timeMax,
SampleVector::Scalar const *  g_timeError,
float const *  g_rms_x12,
float const *  timeCalibConstant,
::ecal::reco::StorageScalarType *  g_jitterEB,
::ecal::reco::StorageScalarType *  g_jitterEE,
::ecal::reco::StorageScalarType *  g_jitterErrorEB,
::ecal::reco::StorageScalarType *  g_jitterErrorEE,
uint32_t *  flagsEB,
uint32_t *  flagsEE,
int const  amplitudeBinsSizeEB,
int const  amplitudeBinsSizeEE,
ConfigurationParameters::type const  timeConstantTermEB,
ConfigurationParameters::type const  timeConstantTermEE,
float const  offsetTimeValueEB,
float const  offsetTimeValueEE,
ConfigurationParameters::type const  timeNconstEB,
ConfigurationParameters::type const  timeNconstEE,
ConfigurationParameters::type const  amplitudeThresholdEB,
ConfigurationParameters::type const  amplitudeThresholdEE,
ConfigurationParameters::type const  outOfTimeThreshG12pEB,
ConfigurationParameters::type const  outOfTimeThreshG12pEE,
ConfigurationParameters::type const  outOfTimeThreshG12mEB,
ConfigurationParameters::type const  outOfTimeThreshG12mEE,
ConfigurationParameters::type const  outOfTimeThreshG61pEB,
ConfigurationParameters::type const  outOfTimeThreshG61pEE,
ConfigurationParameters::type const  outOfTimeThreshG61mEB,
ConfigurationParameters::type const  outOfTimeThreshG61mEE,
uint32_t const  offsetForHashes,
uint32_t const  offsetForInputs,
int const  nchannels 
)

launch context parameters: 1 thread per channel

Variable Documentation

◆ FullSampleVectorSize

constexpr int ecal::multifit::FullSampleVectorSize = 19

Definition at line 13 of file EigenMatrixTypes_gpu.h.

◆ NGains

constexpr int ecal::multifit::NGains = 3

Definition at line 15 of file EigenMatrixTypes_gpu.h.

◆ PulseVectorSize

constexpr int ecal::multifit::PulseVectorSize = 12

Definition at line 14 of file EigenMatrixTypes_gpu.h.

◆ SampleVectorSize

constexpr int ecal::multifit::SampleVectorSize = 10

Definition at line 12 of file EigenMatrixTypes_gpu.h.

Generated for CMSSW Reference Manual by doxygen 1.8.14