EcalRecProducers/plugins/AmplitudeComputationCommonKernels.h

0001 #ifndef RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationCommonKernels_h
0002 #define RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationCommonKernels_h
0003
0004 #include "RecoLocalCalo/EcalRecProducers/interface/EigenMatrixTypes_gpu.h"
0005 #include "DeclsForKernels.h"
0006
0007 class EcalPulseShape;
0008 // this flag setting is applied to all of the cases
0009 class EcalPulseCovariance;
0010 class EcalUncalibratedRecHit;
0011
0012 namespace ecal {
0013   namespace multifit {
0014
0015     ///
0016     /// assume kernel launch configuration is
0017     /// (MAXSAMPLES * nchannels, blocks)
0018     /// TODO: is there a point to split this kernel further to separate reductions
0019     ///
0020     __global__ void kernel_prep_1d_and_initialize(EcalPulseShape const* shapes_in,
0021                                                   uint16_t const* digis_in_eb,
0022                                                   uint32_t const* dids_eb,
0023                                                   uint16_t const* digis_in_ee,
0024                                                   uint32_t const* dids_ee,
0025                                                   SampleVector* amplitudes,
0026                                                   SampleVector* amplitudesForMinimizationEB,
0027                                                   SampleVector* amplitudesForMinimizationEE,
0028                                                   SampleGainVector* gainsNoise,
0029                                                   float const* mean_x1,
0030                                                   float const* mean_x12,
0031                                                   float const* rms_x12,
0032                                                   float const* mean_x6,
0033                                                   float const* gain6Over1,
0034                                                   float const* gain12Over6,
0035                                                   bool* hasSwitchToGain6,
0036                                                   bool* hasSwitchToGain1,
0037                                                   bool* isSaturated,
0038                                                   ::ecal::reco::StorageScalarType* energiesEB,
0039                                                   ::ecal::reco::StorageScalarType* energiesEE,
0040                                                   ::ecal::reco::StorageScalarType* chi2EB,
0041                                                   ::ecal::reco::StorageScalarType* chi2EE,
0042                                                   ::ecal::reco::StorageScalarType* pedestalEB,
0043                                                   ::ecal::reco::StorageScalarType* pedestalEE,
0044                                                   uint32_t* dids_outEB,
0045                                                   uint32_t* dids_outEE,
0046                                                   uint32_t* flagsEB,
0047                                                   uint32_t* flagsEE,
0048                                                   char* acState,
0049                                                   BXVectorType* bxs,
0050                                                   uint32_t const offsetForHashes,
0051                                                   uint32_t const offsetForInputs,
0052                                                   bool const gainSwitchUseMaxSampleEB,
0053                                                   bool const gainSwitchUseMaxSampleEE,
0054                                                   int const nchannels);
0055
0056     ///
0057     /// assume kernel launch configuration is
0058     /// ([MAXSAMPLES, MAXSAMPLES], nchannels)
0059     ///
0060     __global__ void kernel_prep_2d(SampleGainVector const* gainNoise,
0061                                    uint32_t const* dids_eb,
0062                                    uint32_t const* dids_ee,
0063                                    float const* rms_x12,
0064                                    float const* rms_x6,
0065                                    float const* rms_x1,
0066                                    float const* gain12Over6,
0067                                    float const* gain6Over1,
0068                                    double const* G12SamplesCorrelationEB,
0069                                    double const* G6SamplesCorrelationEB,
0070                                    double const* G1SamplesCorrelationEB,
0071                                    double const* G12SamplesCorrelationEE,
0072                                    double const* G6SamplesCorrelationEE,
0073                                    double const* G1SamplesCorrelationEE,
0074                                    SampleMatrix* noisecov,
0075                                    PulseMatrixType* pulse_matrix,
0076                                    EcalPulseShape const* pulse_shape,
0077                                    bool const* hasSwitchToGain6,
0078                                    bool const* hasSwitchToGain1,
0079                                    bool const* isSaturated,
0080                                    uint32_t const offsetForHashes,
0081                                    uint32_t const offsetForInputs);
0082
0083     __global__ void kernel_permute_results(SampleVector* amplitudes,
0084                                            BXVectorType const* activeBXs,
0085                                            ::ecal::reco::StorageScalarType* energies,
0086                                            char const* acState,
0087                                            int const nchannels);
0088
0089 ///
0090 /// Build an Ecal RecHit.
0091 /// TODO: Use SoA data structures on the host directly
0092 /// the reason for removing this from minimize kernel is to isolate the minimize +
0093 /// again, building an aos rec hit involves strides... -> bad memory access pattern
0094 ///
0095 #ifdef RUN_BUILD_AOS_RECHIT
0096     __global__ void kernel_build_rechit(
0097         float const* energies, float const* chi2s, uint32_t* dids, EcalUncalibratedRecHit* rechits, int nchannels);
0098 #endif  // RUN_BUILD_AOS_RECHIT
0099
0100   }  // namespace multifit
0101 }  // namespace ecal
0102
0103 #endif  // RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationCommonKernels_h