Back to home page

Project CMSSW displayed by LXR
 
 

    

File indexing completed on 2024-04-06 12:25:45 

0001 #ifndef RecoLocalCalo_EcalRecProducers_plugins_TimeComputationKernels_h
0002 #define RecoLocalCalo_EcalRecProducers_plugins_TimeComputationKernels_h
0003 
0004 #include <iostream>
0005 #include <limits>
0006 
0007 #include <cuda.h>
0008 
0009 #include "DataFormats/Math/interface/approx_exp.h"
0010 #include "DataFormats/Math/interface/approx_log.h"
0011 #include "RecoLocalCalo/EcalRecProducers/interface/EigenMatrixTypes_gpu.h"
0012 
0013 #include "DeclsForKernels.h"
0014 
0015 //#define DEBUG
0016 
0017 //#define ECAL_RECO_CUDA_DEBUG
0018 
0019 namespace ecal {
0020   namespace multifit {
0021 
0022     __global__ void kernel_time_compute_nullhypot(SampleVector::Scalar const* sample_values,
0023                                                   SampleVector::Scalar const* sample_value_errors,
0024                                                   bool const* useless_sample_values,
0025                                                   SampleVector::Scalar* chi2s,
0026                                                   SampleVector::Scalar* sum0s,
0027                                                   SampleVector::Scalar* sumAAs,
0028                                                   int const nchannels);
0029     //
0030     // launch ctx parameters are
0031     // 45 threads per channel, X channels per block, Y blocks
0032     // 45 comes from: 10 samples for i <- 0 to 9 and for j <- i+1 to 9
0033     // TODO: it might be much beter to use 32 threads per channel instead of 45
0034     // to simplify the synchronization
0035     //
0036     __global__ void kernel_time_compute_makeratio(SampleVector::Scalar const* sample_values,
0037                                                   SampleVector::Scalar const* sample_value_errors,
0038                                                   uint32_t const* dids_eb,
0039                                                   uint32_t const* dids_ee,
0040                                                   bool const* useless_sample_values,
0041                                                   char const* pedestal_nums,
0042                                                   ConfigurationParameters::type const* amplitudeFitParametersEB,
0043                                                   ConfigurationParameters::type const* amplitudeFitParametersEE,
0044                                                   ConfigurationParameters::type const* timeFitParametersEB,
0045                                                   ConfigurationParameters::type const* timeFitParametersEE,
0046                                                   SampleVector::Scalar const* sumAAsNullHypot,
0047                                                   SampleVector::Scalar const* sum0sNullHypot,
0048                                                   SampleVector::Scalar* tMaxAlphaBetas,
0049                                                   SampleVector::Scalar* tMaxErrorAlphaBetas,
0050                                                   SampleVector::Scalar* g_accTimeMax,
0051                                                   SampleVector::Scalar* g_accTimeWgt,
0052                                                   TimeComputationState* g_state,
0053                                                   unsigned int const timeFitParameters_sizeEB,
0054                                                   unsigned int const timeFitParameters_sizeEE,
0055                                                   ConfigurationParameters::type const timeFitLimits_firstEB,
0056                                                   ConfigurationParameters::type const timeFitLimits_firstEE,
0057                                                   ConfigurationParameters::type const timeFitLimits_secondEB,
0058                                                   ConfigurationParameters::type const timeFitLimits_secondEE,
0059                                                   int const nchannels,
0060                                                   uint32_t const offsetForInputs);
0061 
0062     /// launch ctx parameters are
0063     /// 10 threads per channel, N channels per block, Y blocks
0064     /// TODO: do we need to keep the state around or can be removed?!
0065     //#define DEBUG_FINDAMPLCHI2_AND_FINISH
0066     __global__ void kernel_time_compute_findamplchi2_and_finish(
0067         SampleVector::Scalar const* sample_values,
0068         SampleVector::Scalar const* sample_value_errors,
0069         uint32_t const* dids_eb,
0070         uint32_t const* dids_ee,
0071         bool const* useless_samples,
0072         SampleVector::Scalar const* g_tMaxAlphaBeta,
0073         SampleVector::Scalar const* g_tMaxErrorAlphaBeta,
0074         SampleVector::Scalar const* g_accTimeMax,
0075         SampleVector::Scalar const* g_accTimeWgt,
0076         ConfigurationParameters::type const* amplitudeFitParametersEB,
0077         ConfigurationParameters::type const* amplitudeFitParametersEE,
0078         SampleVector::Scalar const* sumAAsNullHypot,
0079         SampleVector::Scalar const* sum0sNullHypot,
0080         SampleVector::Scalar const* chi2sNullHypot,
0081         TimeComputationState* g_state,
0082         SampleVector::Scalar* g_ampMaxAlphaBeta,
0083         SampleVector::Scalar* g_ampMaxError,
0084         SampleVector::Scalar* g_timeMax,
0085         SampleVector::Scalar* g_timeError,
0086         int const nchannels,
0087         uint32_t const offsetForInputs);
0088 
0089     __global__ void kernel_time_compute_fixMGPAslew(uint16_t const* digis_eb,
0090                                                     uint16_t const* digis_ee,
0091                                                     SampleVector::Scalar* sample_values,
0092                                                     SampleVector::Scalar* sample_value_errors,
0093                                                     bool* useless_sample_values,
0094                                                     unsigned int const sample_mask,
0095                                                     int const nchannels,
0096                                                     uint32_t const offsetForInputs);
0097 
0098     __global__ void kernel_time_compute_ampl(SampleVector::Scalar const* sample_values,
0099                                              SampleVector::Scalar const* sample_value_errors,
0100                                              uint32_t const* dids_eb,
0101                                              uint32_t const* dids_ed,
0102                                              bool const* useless_samples,
0103                                              SampleVector::Scalar const* g_timeMax,
0104                                              SampleVector::Scalar const* amplitudeFitParametersEB,
0105                                              SampleVector::Scalar const* amplitudeFitParametersEE,
0106                                              SampleVector::Scalar* g_amplitudeMax,
0107                                              int const nchannels,
0108                                              uint32_t const offsetForInputs);
0109 
0110     //#define ECAL_RECO_CUDA_TC_INIT_DEBUG
0111     __global__ void kernel_time_computation_init(uint16_t const* digis_eb,
0112                                                  uint32_t const* dids_eb,
0113                                                  uint16_t const* digis_ee,
0114                                                  uint32_t const* dids_ee,
0115                                                  float const* rms_x12,
0116                                                  float const* rms_x6,
0117                                                  float const* rms_x1,
0118                                                  float const* mean_x12,
0119                                                  float const* mean_x6,
0120                                                  float const* mean_x1,
0121                                                  float const* gain12Over6,
0122                                                  float const* gain6Over1,
0123                                                  SampleVector::Scalar* sample_values,
0124                                                  SampleVector::Scalar* sample_value_errors,
0125                                                  SampleVector::Scalar* ampMaxError,
0126                                                  bool* useless_sample_values,
0127                                                  char* pedestal_nums,
0128                                                  uint32_t const offsetForHashes,
0129                                                  uint32_t const offsetForInputs,
0130                                                  unsigned int const sample_maskEB,
0131                                                  unsigned int const sample_maskEE,
0132                                                  int nchannels);
0133 
0134     ///
0135     /// launch context parameters: 1 thread per channel
0136     ///
0137     //#define DEBUG_TIME_CORRECTION
0138     __global__ void kernel_time_correction_and_finalize(
0139         //        SampleVector::Scalar const* g_amplitude,
0140         ::ecal::reco::StorageScalarType const* g_amplitudeEB,
0141         ::ecal::reco::StorageScalarType const* g_amplitudeEE,
0142         uint16_t const* digis_eb,
0143         uint32_t const* dids_eb,
0144         uint16_t const* digis_ee,
0145         uint32_t const* dids_ee,
0146         float const* amplitudeBinsEB,
0147         float const* amplitudeBinsEE,
0148         float const* shiftBinsEB,
0149         float const* shiftBinsEE,
0150         SampleVector::Scalar const* g_timeMax,
0151         SampleVector::Scalar const* g_timeError,
0152         float const* g_rms_x12,
0153         float const* timeCalibConstant,
0154         ::ecal::reco::StorageScalarType* g_jitterEB,
0155         ::ecal::reco::StorageScalarType* g_jitterEE,
0156         ::ecal::reco::StorageScalarType* g_jitterErrorEB,
0157         ::ecal::reco::StorageScalarType* g_jitterErrorEE,
0158         uint32_t* flagsEB,
0159         uint32_t* flagsEE,
0160         int const amplitudeBinsSizeEB,
0161         int const amplitudeBinsSizeEE,
0162         ConfigurationParameters::type const timeConstantTermEB,
0163         ConfigurationParameters::type const timeConstantTermEE,
0164         float const offsetTimeValueEB,
0165         float const offsetTimeValueEE,
0166         ConfigurationParameters::type const timeNconstEB,
0167         ConfigurationParameters::type const timeNconstEE,
0168         ConfigurationParameters::type const amplitudeThresholdEB,
0169         ConfigurationParameters::type const amplitudeThresholdEE,
0170         ConfigurationParameters::type const outOfTimeThreshG12pEB,
0171         ConfigurationParameters::type const outOfTimeThreshG12pEE,
0172         ConfigurationParameters::type const outOfTimeThreshG12mEB,
0173         ConfigurationParameters::type const outOfTimeThreshG12mEE,
0174         ConfigurationParameters::type const outOfTimeThreshG61pEB,
0175         ConfigurationParameters::type const outOfTimeThreshG61pEE,
0176         ConfigurationParameters::type const outOfTimeThreshG61mEB,
0177         ConfigurationParameters::type const outOfTimeThreshG61mEE,
0178         uint32_t const offsetForHashes,
0179         uint32_t const offsetForInputs,
0180         int const nchannels);
0181 
0182   }  // namespace multifit
0183 }  // namespace ecal
0184 
0185 #endif  // RecoLocalCalo_EcalRecProducers_plugins_TimeComputationKernels_h
[ Source navigation ] [ Diff markup ] [ Identifier search ] [ general search ]

This page was automatically generated by the 2.2.1 LXR engine.

The LXR team
Valid CSS 2.1! Valid HTML 4.01!