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 27 of file EigenMatrixTypes_gpu.h.

◆ BXVectorType

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

Definition at line 21 of file EigenMatrixTypes_gpu.h.

◆ data_type

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

Definition at line 18 of file EigenMatrixTypes_gpu.h.

◆ FullSampleMatrix

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

Definition at line 30 of file EigenMatrixTypes_gpu.h.

◆ FullSampleVector

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

Definition at line 25 of file EigenMatrixTypes_gpu.h.

◆ PermutationMatrix

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

Definition at line 44 of file EigenMatrixTypes_gpu.h.

◆ PulseDecompLDLT

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

Definition at line 37 of file EigenMatrixTypes_gpu.h.

◆ PulseDecompLLT

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

Definition at line 36 of file EigenMatrixTypes_gpu.h.

◆ PulseMatrix

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

Definition at line 31 of file EigenMatrixTypes_gpu.h.

◆ PulseMatrixType

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

Definition at line 20 of file EigenMatrixTypes_gpu.h.

◆ PulseVector

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

Definition at line 26 of file EigenMatrixTypes_gpu.h.

◆ SampleDecompLLT

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

Definition at line 34 of file EigenMatrixTypes_gpu.h.

◆ SampleDecompLLTD

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

Definition at line 35 of file EigenMatrixTypes_gpu.h.

◆ SampleGainVector

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

Definition at line 28 of file EigenMatrixTypes_gpu.h.

◆ SampleMatrix

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

Definition at line 29 of file EigenMatrixTypes_gpu.h.

◆ SampleMatrixD

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

Definition at line 22 of file EigenMatrixTypes_gpu.h.

◆ SampleMatrixGainArray

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

Definition at line 42 of file EigenMatrixTypes_gpu.h.

◆ SamplePulseMatrix

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

Definition at line 33 of file EigenMatrixTypes_gpu.h.

◆ SampleVector

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

Definition at line 24 of file EigenMatrixTypes_gpu.h.

◆ SingleMatrix

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

Definition at line 39 of file EigenMatrixTypes_gpu.h.

◆ SingleVector

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

Definition at line 40 of file EigenMatrixTypes_gpu.h.

Enumeration Type Documentation

◆ MinimizationState

enum ecal::multifit::MinimizationState : char

strong

Enumerator
NotFinished
Finished
Precomputed

Definition at line 53 of file DeclsForKernels.h.

                                  : char {
       NotFinished = 0,
       Finished = 1,
       Precomputed = 2,
     };

◆ TimeComputationState

enum ecal::multifit::TimeComputationState : char

strong

Enumerator
NotFinished
Finished

Definition at line 52 of file DeclsForKernels.h.

52 : char { NotFinished = 0, Finished = 1 };

Function Documentation

◆ entryPoint()

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

Referenced by EcalUncalibRecHitProducerGPU::acquire().

◆ getLength()

template<typename EigenM >

constexpr auto ecal::multifit::getLength ( ) -> uint32_t

constexpr

Definition at line 135 of file DeclsForKernels.h.

                                            {
       return EigenM::RowsAtCompileTime * EigenM::ColsAtCompileTime;
     }

◆ 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

constexpr

Definition at line 14 of file EigenMatrixTypes_gpu.h.

◆ NGains

constexpr int ecal::multifit::NGains = 3

constexpr

Definition at line 16 of file EigenMatrixTypes_gpu.h.

◆ PulseVectorSize

constexpr int ecal::multifit::PulseVectorSize = 12

constexpr

Definition at line 15 of file EigenMatrixTypes_gpu.h.

◆ SampleVectorSize

constexpr int ecal::multifit::SampleVectorSize = 10

constexpr

Definition at line 13 of file EigenMatrixTypes_gpu.h.

Namespaces

Classes

Typedefs

Enumerations

Functions

Variables

Typedef Documentation

◆ BXVector

◆ BXVectorType

◆ data_type

◆ FullSampleMatrix

◆ FullSampleVector

◆ PermutationMatrix

◆ PulseDecompLDLT

◆ PulseDecompLLT

◆ PulseMatrix

◆ PulseMatrixType

◆ PulseVector

◆ SampleDecompLLT

◆ SampleDecompLLTD

◆ SampleGainVector

◆ SampleMatrix

◆ SampleMatrixD

◆ SampleMatrixGainArray

◆ SamplePulseMatrix

◆ SampleVector

◆ SingleMatrix

◆ SingleVector

Enumeration Type Documentation

◆ MinimizationState

◆ TimeComputationState

Function Documentation

◆ entryPoint()

◆ getLength()

◆ kernel_permute_results()

◆ kernel_prep_1d_and_initialize()

◆ kernel_prep_2d()

◆ kernel_time_computation_init()

◆ kernel_time_compute_ampl()

◆ kernel_time_compute_findamplchi2_and_finish()

◆ kernel_time_compute_fixMGPAslew()

◆ kernel_time_compute_makeratio()

◆ kernel_time_compute_nullhypot()

◆ kernel_time_correction_and_finalize()

Variable Documentation

◆ FullSampleVectorSize

◆ NGains

◆ PulseVectorSize

◆ SampleVectorSize