SiStripObjects/interface/SiStripClusterizerConditionsGPU.h

0001 #ifndef CalibFormats_SiStripObjects_SiStripClusterizerConditionsGPU_h
0002 #define CalibFormats_SiStripObjects_SiStripClusterizerConditionsGPU_h
0003
0004 #include "DataFormats/SiStripCluster/interface/SiStripTypes.h"
0005 #include "DataFormats/SiStripCommon/interface/ConstantsForHardwareSystems.h"
0006
0007 #include "HeterogeneousCore/CUDACore/interface/ESProduct.h"
0008 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
0009 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
0010 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
0011 #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
0012
0013 class SiStripQuality;
0014 class SiStripGain;
0015 class SiStripNoises;
0016
0017 namespace stripgpu {
0018   __host__ __device__ inline fedId_t fedIndex(fedId_t fed) { return fed - sistrip::FED_ID_MIN; }
0019   __host__ __device__ inline std::uint32_t stripIndex(fedId_t fed, fedCh_t channel, stripId_t strip) {
0020     return fedIndex(fed) * sistrip::FEDCH_PER_FED * sistrip::STRIPS_PER_FEDCH + channel * sistrip::STRIPS_PER_FEDCH +
0021            (strip % sistrip::STRIPS_PER_FEDCH);
0022   }
0023   __host__ __device__ inline std::uint32_t apvIndex(fedId_t fed, fedCh_t channel, stripId_t strip) {
0024     return fedIndex(fed) * sistrip::APVS_PER_FEDCH * sistrip::FEDCH_PER_FED + sistrip::APVS_PER_CHAN * channel +
0025            (strip % sistrip::STRIPS_PER_FEDCH) / sistrip::STRIPS_PER_APV;
0026   }
0027   __host__ __device__ inline std::uint32_t channelIndex(fedId_t fed, fedCh_t channel) {
0028     return fedIndex(fed) * sistrip::FEDCH_PER_FED + channel;
0029   }
0030
0031   class SiStripClusterizerConditionsGPU {
0032   public:
0033     class DetToFed {
0034     public:
0035       DetToFed(detId_t detid, apvPair_t ipair, fedId_t fedid, fedCh_t fedch)
0036           : detid_(detid), ipair_(ipair), fedid_(fedid), fedch_(fedch) {}
0037       detId_t detID() const { return detid_; }
0038       apvPair_t pair() const { return ipair_; }
0039       fedId_t fedID() const { return fedid_; }
0040       fedCh_t fedCh() const { return fedch_; }
0041
0042     private:
0043       detId_t detid_;
0044       apvPair_t ipair_;
0045       fedId_t fedid_;
0046       fedCh_t fedch_;
0047     };
0048     using DetToFeds = std::vector<DetToFed>;
0049
0050     static constexpr std::uint16_t badBit = 1 << 15;
0051
0052     class Data {
0053     public:
0054       struct DeviceView {
0055         __device__ inline detId_t detID(fedId_t fed, fedCh_t channel) const {
0056           return detID_[channelIndex(fed, channel)];
0057         }
0058
0059         __device__ inline apvPair_t iPair(fedId_t fed, fedCh_t channel) const {
0060           return iPair_[channelIndex(fed, channel)];
0061         }
0062
0063         __device__ inline float invthick(fedId_t fed, fedCh_t channel) const {
0064           return invthick_[channelIndex(fed, channel)];
0065         }
0066
0067         __device__ inline float noise(fedId_t fed, fedCh_t channel, stripId_t strip) const {
0068           // noise is stored as 9 bits with a fixed point scale factor of 0.1
0069           return 0.1f * (noise_[stripIndex(fed, channel, strip)] & ~badBit);
0070         }
0071
0072         __device__ inline float gain(fedId_t fed, fedCh_t channel, stripId_t strip) const {
0073           return gain_[apvIndex(fed, channel, strip)];
0074         }
0075
0076         __device__ inline bool bad(fedId_t fed, fedCh_t channel, stripId_t strip) const {
0077           return badBit == (noise_[stripIndex(fed, channel, strip)] & badBit);
0078         }
0079         const std::uint16_t* noise_;  //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED * sistrip::STRIPS_PER_FEDCH];
0080         const float* invthick_;       //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED];
0081         const detId_t* detID_;        //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED];
0082         const apvPair_t* iPair_;      //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED];
0083         const float* gain_;           //[sistrip::NUMBER_OF_FEDS*sistrip::APVS_PER_FEDCH * sistrip::FEDCH_PER_FED];
0084       };
0085
0086       const DeviceView* deviceView() const { return deviceView_.get(); }
0087
0088       cms::cuda::device::unique_ptr<DeviceView> deviceView_;
0089       cms::cuda::host::unique_ptr<DeviceView> hostView_;
0090
0091       cms::cuda::device::unique_ptr<std::uint16_t[]>
0092           noise_;  //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED * sistrip::STRIPS_PER_FEDCH];
0093       cms::cuda::device::unique_ptr<float[]> invthick_;   //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED];
0094       cms::cuda::device::unique_ptr<detId_t[]> detID_;    //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED];
0095       cms::cuda::device::unique_ptr<apvPair_t[]> iPair_;  //[sistrip::NUMBER_OF_FEDS*sistrip::FEDCH_PER_FED];
0096       cms::cuda::device::unique_ptr<float[]>
0097           gain_;  //[sistrip::NUMBER_OF_FEDS*sistrip::APVS_PER_FEDCH * sistrip::FEDCH_PER_FED];
0098     };
0099
0100     SiStripClusterizerConditionsGPU(const SiStripQuality& quality,
0101                                     const SiStripGain* gains,
0102                                     const SiStripNoises& noises);
0103     ~SiStripClusterizerConditionsGPU() = default;
0104
0105     // Function to return the actual payload on the memory of the current device
0106     Data const& getGPUProductAsync(cudaStream_t stream) const;
0107
0108     const DetToFeds& detToFeds() const { return detToFeds_; }
0109
0110   private:
0111     void setStrip(fedId_t fed, fedCh_t channel, stripId_t strip, std::uint16_t noise, float gain, bool bad) {
0112       gain_[apvIndex(fed, channel, strip)] = gain;
0113       noise_[stripIndex(fed, channel, strip)] = noise;
0114       if (bad) {
0115         noise_[stripIndex(fed, channel, strip)] |= badBit;
0116       }
0117     }
0118
0119     void setInvThickness(fedId_t fed, fedCh_t channel, float invthick) {
0120       invthick_[channelIndex(fed, channel)] = invthick;
0121     }
0122
0123     // Holds the data in pinned CPU memory
0124     std::vector<std::uint16_t, cms::cuda::HostAllocator<std::uint16_t>> noise_;
0125     std::vector<float, cms::cuda::HostAllocator<float>> invthick_;
0126     std::vector<detId_t, cms::cuda::HostAllocator<detId_t>> detID_;
0127     std::vector<apvPair_t, cms::cuda::HostAllocator<apvPair_t>> iPair_;
0128     std::vector<float, cms::cuda::HostAllocator<float>> gain_;
0129
0130     // Helper that takes care of complexity of transferring the data to
0131     // multiple devices
0132     cms::cuda::ESProduct<Data> gpuData_;
0133     DetToFeds detToFeds_;
0134   };
0135 }  // namespace stripgpu
0136
0137 #endif