CMS 3D CMS Logo

 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Properties Friends Macros Groups Pages
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< PulseMatrixPulseDecompLDLT
 
typedef Eigen::LLT< PulseMatrixPulseDecompLLT
 
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< SampleMatrixSampleDecompLLT
 
typedef Eigen::LLT< SampleMatrixDSampleDecompLLTD
 
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

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

Definition at line 27 of file EigenMatrixTypes_gpu.h.

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

Definition at line 21 of file EigenMatrixTypes_gpu.h.

Definition at line 18 of file EigenMatrixTypes_gpu.h.

Definition at line 30 of file EigenMatrixTypes_gpu.h.

Definition at line 25 of file EigenMatrixTypes_gpu.h.

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

Definition at line 44 of file EigenMatrixTypes_gpu.h.

Definition at line 37 of file EigenMatrixTypes_gpu.h.

Definition at line 36 of file EigenMatrixTypes_gpu.h.

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

Definition at line 31 of file EigenMatrixTypes_gpu.h.

Definition at line 20 of file EigenMatrixTypes_gpu.h.

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

Definition at line 26 of file EigenMatrixTypes_gpu.h.

Definition at line 34 of file EigenMatrixTypes_gpu.h.

Definition at line 35 of file EigenMatrixTypes_gpu.h.

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

Definition at line 28 of file EigenMatrixTypes_gpu.h.

Definition at line 29 of file EigenMatrixTypes_gpu.h.

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

Definition at line 22 of file EigenMatrixTypes_gpu.h.

Definition at line 42 of file EigenMatrixTypes_gpu.h.

Definition at line 33 of file EigenMatrixTypes_gpu.h.

Definition at line 24 of file EigenMatrixTypes_gpu.h.

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

Definition at line 39 of file EigenMatrixTypes_gpu.h.

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

Definition at line 40 of file EigenMatrixTypes_gpu.h.

Enumeration Type Documentation

Enumerator
NotFinished 
Finished 
Precomputed 

Definition at line 53 of file DeclsForKernels.h.

Enumerator
NotFinished 
Finished 

Definition at line 52 of file DeclsForKernels.h.

Function Documentation

void ecal::multifit::entryPoint ( EventInputDataGPU const &  ,
EventOutputDataGPU &  ,
EventDataForScratchGPU &  ,
ConditionsProducts const &  ,
ConfigurationParameters const &  ,
cudaStream_t   
)
template<typename EigenM >
constexpr auto ecal::multifit::getLength ( ) -> uint32_t

Definition at line 135 of file DeclsForKernels.h.

135  {
136  return EigenM::RowsAtCompileTime * EigenM::ColsAtCompileTime;
137  }
__global__ void ecal::multifit::kernel_permute_results ( SampleVector amplitudes,
BXVectorType const *  activeBXs,
::ecal::reco::StorageScalarType energies,
char const *  acState,
int const  nchannels 
)
__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

__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)

__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 
)
__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 
)
__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?!

__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 
)
__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 
)
__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 
)
__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

constexpr int ecal::multifit::FullSampleVectorSize = 19

Definition at line 14 of file EigenMatrixTypes_gpu.h.

constexpr int ecal::multifit::NGains = 3

Definition at line 16 of file EigenMatrixTypes_gpu.h.

constexpr int ecal::multifit::PulseVectorSize = 12

Definition at line 15 of file EigenMatrixTypes_gpu.h.

constexpr int ecal::multifit::SampleVectorSize = 10

Definition at line 13 of file EigenMatrixTypes_gpu.h.