CMS 3D CMS Logo
TimeComputationKernels.h
Go to the documentation of this file.
1 #ifndef RecoLocalCalo_EcalRecProducers_plugins_TimeComputationKernels_h
2 #define RecoLocalCalo_EcalRecProducers_plugins_TimeComputationKernels_h
3 
4 #include <iostream>
5 #include <limits>
6 
7 #include <cuda.h>
8 
9 #include "DataFormats/Math/interface/approx_exp.h"
10 #include "DataFormats/Math/interface/approx_log.h"
11 
12 #include "DeclsForKernels.h"
13 #include "EigenMatrixTypes_gpu.h"
14 
15 //#define DEBUG
16 
17 //#define ECAL_RECO_CUDA_DEBUG
18 
19 namespace ecal {
20  namespace multifit {
21 
22  __global__ void kernel_time_compute_nullhypot(SampleVector::Scalar const* sample_values,
23  SampleVector::Scalar const* sample_value_errors,
24  bool const* useless_sample_values,
25  SampleVector::Scalar* chi2s,
26  SampleVector::Scalar* sum0s,
27  SampleVector::Scalar* sumAAs,
28  int const nchannels);
29  //
30  // launch ctx parameters are
31  // 45 threads per channel, X channels per block, Y blocks
32  // 45 comes from: 10 samples for i <- 0 to 9 and for j <- i+1 to 9
33  // TODO: it might be much beter to use 32 threads per channel instead of 45
34  // to simplify the synchronization
35  //
36  __global__ void kernel_time_compute_makeratio(SampleVector::Scalar const* sample_values,
37  SampleVector::Scalar const* sample_value_errors,
38  uint32_t const* dids_eb,
39  uint32_t const* dids_ee,
40  bool const* useless_sample_values,
41  char const* pedestal_nums,
42  ConfigurationParameters::type const* amplitudeFitParametersEB,
43  ConfigurationParameters::type const* amplitudeFitParametersEE,
44  ConfigurationParameters::type const* timeFitParametersEB,
45  ConfigurationParameters::type const* timeFitParametersEE,
46  SampleVector::Scalar const* sumAAsNullHypot,
47  SampleVector::Scalar const* sum0sNullHypot,
48  SampleVector::Scalar* tMaxAlphaBetas,
49  SampleVector::Scalar* tMaxErrorAlphaBetas,
50  SampleVector::Scalar* g_accTimeMax,
51  SampleVector::Scalar* g_accTimeWgt,
52  TimeComputationState* g_state,
53  unsigned int const timeFitParameters_sizeEB,
54  unsigned int const timeFitParameters_sizeEE,
55  ConfigurationParameters::type const timeFitLimits_firstEB,
56  ConfigurationParameters::type const timeFitLimits_firstEE,
57  ConfigurationParameters::type const timeFitLimits_secondEB,
58  ConfigurationParameters::type const timeFitLimits_secondEE,
59  int const nchannels,
60  uint32_t const offsetForInputs);
61 
65  //#define DEBUG_FINDAMPLCHI2_AND_FINISH
66  __global__ void kernel_time_compute_findamplchi2_and_finish(
67  SampleVector::Scalar const* sample_values,
68  SampleVector::Scalar const* sample_value_errors,
69  uint32_t const* dids_eb,
70  uint32_t const* dids_ee,
71  bool const* useless_samples,
72  SampleVector::Scalar const* g_tMaxAlphaBeta,
73  SampleVector::Scalar const* g_tMaxErrorAlphaBeta,
74  SampleVector::Scalar const* g_accTimeMax,
75  SampleVector::Scalar const* g_accTimeWgt,
76  ConfigurationParameters::type const* amplitudeFitParametersEB,
77  ConfigurationParameters::type const* amplitudeFitParametersEE,
78  SampleVector::Scalar const* sumAAsNullHypot,
79  SampleVector::Scalar const* sum0sNullHypot,
80  SampleVector::Scalar const* chi2sNullHypot,
81  TimeComputationState* g_state,
82  SampleVector::Scalar* g_ampMaxAlphaBeta,
83  SampleVector::Scalar* g_ampMaxError,
84  SampleVector::Scalar* g_timeMax,
85  SampleVector::Scalar* g_timeError,
86  int const nchannels,
87  uint32_t const offsetForInputs);
88 
89  __global__ void kernel_time_compute_fixMGPAslew(uint16_t const* digis_eb,
90  uint16_t const* digis_ee,
91  SampleVector::Scalar* sample_values,
92  SampleVector::Scalar* sample_value_errors,
93  bool* useless_sample_values,
94  unsigned int const sample_mask,
95  int const nchannels,
96  uint32_t const offsetForInputs);
97 
98  __global__ void kernel_time_compute_ampl(SampleVector::Scalar const* sample_values,
99  SampleVector::Scalar const* sample_value_errors,
100  uint32_t const* dids_eb,
101  uint32_t const* dids_ed,
102  bool const* useless_samples,
103  SampleVector::Scalar const* g_timeMax,
104  SampleVector::Scalar const* amplitudeFitParametersEB,
105  SampleVector::Scalar const* amplitudeFitParametersEE,
106  SampleVector::Scalar* g_amplitudeMax,
107  int const nchannels,
108  uint32_t const offsetForInputs);
109 
110  //#define ECAL_RECO_CUDA_TC_INIT_DEBUG
111  __global__ void kernel_time_computation_init(uint16_t const* digis_eb,
112  uint32_t const* dids_eb,
113  uint16_t const* digis_ee,
114  uint32_t const* dids_ee,
115  float const* rms_x12,
116  float const* rms_x6,
117  float const* rms_x1,
118  float const* mean_x12,
119  float const* mean_x6,
120  float const* mean_x1,
121  float const* gain12Over6,
122  float const* gain6Over1,
123  SampleVector::Scalar* sample_values,
124  SampleVector::Scalar* sample_value_errors,
125  SampleVector::Scalar* ampMaxError,
126  bool* useless_sample_values,
127  char* pedestal_nums,
128  uint32_t const offsetForHashes,
129  uint32_t const offsetForInputs,
130  unsigned int const sample_maskEB,
131  unsigned int const sample_maskEE,
132  int nchannels);
133 
137  //#define DEBUG_TIME_CORRECTION
138  __global__ void kernel_time_correction_and_finalize(
139  // SampleVector::Scalar const* g_amplitude,
140  ::ecal::reco::StorageScalarType const* g_amplitudeEB,
141  ::ecal::reco::StorageScalarType const* g_amplitudeEE,
142  uint16_t const* digis_eb,
143  uint32_t const* dids_eb,
144  uint16_t const* digis_ee,
145  uint32_t const* dids_ee,
146  float const* amplitudeBinsEB,
147  float const* amplitudeBinsEE,
148  float const* shiftBinsEB,
149  float const* shiftBinsEE,
150  SampleVector::Scalar const* g_timeMax,
151  SampleVector::Scalar const* g_timeError,
152  float const* g_rms_x12,
153  float const* timeCalibConstant,
154  ::ecal::reco::StorageScalarType* g_jitterEB,
155  ::ecal::reco::StorageScalarType* g_jitterEE,
156  ::ecal::reco::StorageScalarType* g_jitterErrorEB,
157  ::ecal::reco::StorageScalarType* g_jitterErrorEE,
158  uint32_t* flagsEB,
159  uint32_t* flagsEE,
160  int const amplitudeBinsSizeEB,
161  int const amplitudeBinsSizeEE,
162  ConfigurationParameters::type const timeConstantTermEB,
163  ConfigurationParameters::type const timeConstantTermEE,
164  float const offsetTimeValueEB,
165  float const offsetTimeValueEE,
166  ConfigurationParameters::type const timeNconstEB,
167  ConfigurationParameters::type const timeNconstEE,
168  ConfigurationParameters::type const amplitudeThresholdEB,
169  ConfigurationParameters::type const amplitudeThresholdEE,
170  ConfigurationParameters::type const outOfTimeThreshG12pEB,
171  ConfigurationParameters::type const outOfTimeThreshG12pEE,
172  ConfigurationParameters::type const outOfTimeThreshG12mEB,
173  ConfigurationParameters::type const outOfTimeThreshG12mEE,
174  ConfigurationParameters::type const outOfTimeThreshG61pEB,
175  ConfigurationParameters::type const outOfTimeThreshG61pEE,
176  ConfigurationParameters::type const outOfTimeThreshG61mEB,
177  ConfigurationParameters::type const outOfTimeThreshG61mEE,
178  uint32_t const offsetForHashes,
179  uint32_t const offsetForInputs,
180  int const nchannels);
181 
182  } // namespace multifit
183 } // namespace ecal
184 
185 #endif // RecoLocalCalo_EcalRecProducers_plugins_TimeComputationKernels_h
ecal::reco::StorageScalarType
float StorageScalarType
Definition: RecoTypes.h:8
ecal::multifit::kernel_time_computation_init
__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)
ecal
Definition: ElectronicsMappingGPU.h:13
__global__
#define __global__
Definition: cudaCompat.h:19
align::Scalar
double Scalar
Definition: Definitions.h:25
ecal::multifit::kernel_time_compute_nullhypot
__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)
ecal::multifit::ConfigurationParameters::type
double type
Definition: DeclsForKernels.h:68
HLT_FULL_cff.amplitudeThresholdEB
amplitudeThresholdEB
Definition: HLT_FULL_cff.py:8217
approx_log.h
ecal::multifit::kernel_time_correction_and_finalize
__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)
HLT_FULL_cff.amplitudeThresholdEE
amplitudeThresholdEE
Definition: HLT_FULL_cff.py:8219
ecal::multifit::kernel_time_compute_ampl
__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)
ecal::multifit::kernel_time_compute_findamplchi2_and_finish
__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)
DeclsForKernels.h
ecal::multifit::kernel_time_compute_fixMGPAslew
__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)
ecal::multifit::kernel_time_compute_makeratio
__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)
approx_exp.h
EigenMatrixTypes_gpu.h
ecal::multifit::TimeComputationState
TimeComputationState
Definition: DeclsForKernels.h:52
Generated for CMSSW Reference Manual by doxygen 1.8.16