SiStripClusterizer/plugins/ChannelLocsGPU.h

0001 #ifndef RecoLocalTracker_SiStripClusterizer_plugins_ChannelLocsGPU_h
0002 #define RecoLocalTracker_SiStripClusterizer_plugins_ChannelLocsGPU_h
0003
0004 #include <memory>
0005 #include <vector>
0006
0007 #include <cuda_runtime.h>
0008
0009 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
0010 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
0011 #include "DataFormats/SiStripCluster/interface/SiStripTypes.h"
0012
0013 class ChannelLocsGPU;
0014
0015 template <template <typename> class T>
0016 class ChannelLocsBase {
0017 public:
0018   ChannelLocsBase(size_t size) : size_(size) {}
0019   virtual ~ChannelLocsBase() = default;
0020
0021   ChannelLocsBase(ChannelLocsBase&& arg)
0022       : input_(std::move(arg.input_)),
0023         inoff_(std::move(arg.inoff_)),
0024         offset_(std::move(arg.offset_)),
0025         length_(std::move(arg.length_)),
0026         fedID_(std::move(arg.fedID_)),
0027         fedCh_(std::move(arg.fedCh_)),
0028         detID_(std::move(arg.detID_)),
0029         size_(arg.size_) {}
0030
0031   void setChannelLoc(uint32_t index,
0032                      const uint8_t* input,
0033                      size_t inoff,
0034                      size_t offset,
0035                      uint16_t length,
0036                      stripgpu::fedId_t fedID,
0037                      stripgpu::fedCh_t fedCh,
0038                      stripgpu::detId_t detID) {
0039     input_[index] = input;
0040     inoff_[index] = inoff;
0041     offset_[index] = offset;
0042     length_[index] = length;
0043     fedID_[index] = fedID;
0044     fedCh_[index] = fedCh;
0045     detID_[index] = detID;
0046   }
0047
0048   size_t size() const { return size_; }
0049
0050   const uint8_t* input(uint32_t index) const { return input_[index]; }
0051   size_t inoff(uint32_t index) const { return inoff_[index]; }
0052   size_t offset(uint32_t index) const { return offset_[index]; }
0053   uint16_t length(uint32_t index) const { return length_[index]; }
0054   stripgpu::fedId_t fedID(uint32_t index) const { return fedID_[index]; }
0055   stripgpu::fedCh_t fedCh(uint32_t index) const { return fedCh_[index]; }
0056   stripgpu::detId_t detID(uint32_t index) const { return detID_[index]; }
0057
0058   const uint8_t* const* input() const { return input_.get(); }
0059   size_t* inoff() const { return inoff_.get(); }
0060   size_t* offset() const { return offset_.get(); }
0061   uint16_t* length() const { return length_.get(); }
0062   stripgpu::fedId_t* fedID() const { return fedID_.get(); }
0063   stripgpu::fedCh_t* fedCh() const { return fedCh_.get(); }
0064   stripgpu::detId_t* detID() const { return detID_.get(); }
0065
0066 protected:
0067   T<const uint8_t*[]> input_;  // input raw data for channel
0068   T<size_t[]> inoff_;          // offset in input raw data
0069   T<size_t[]> offset_;         // global offset in alldata
0070   T<uint16_t[]> length_;       // length of channel data
0071   T<stripgpu::fedId_t[]> fedID_;
0072   T<stripgpu::fedCh_t[]> fedCh_;
0073   T<stripgpu::detId_t[]> detID_;
0074   size_t size_ = 0;
0075 };
0076
0077 class ChannelLocs : public ChannelLocsBase<cms::cuda::host::unique_ptr> {
0078   friend class ChannelLocsGPU;
0079
0080 public:
0081   ChannelLocs(size_t size, cudaStream_t stream);
0082   ChannelLocs(ChannelLocs&& arg) : ChannelLocsBase(std::move(arg)) {}
0083
0084   ChannelLocs(ChannelLocs&) = delete;
0085   ChannelLocs(const ChannelLocs&) = delete;
0086   ChannelLocs& operator=(const ChannelLocs&) = delete;
0087   ChannelLocs& operator=(ChannelLocs&&) = delete;
0088
0089   ~ChannelLocs() override = default;
0090 };
0091
0092 class ChannelLocsView {
0093 public:
0094   void fill(const ChannelLocsGPU& c);
0095
0096   __device__ size_t size() const { return size_; }
0097
0098   __device__ const uint8_t* input(uint32_t index) const { return input_[index]; }
0099   __device__ size_t inoff(uint32_t index) const { return inoff_[index]; }
0100   __device__ size_t offset(uint32_t index) const { return offset_[index]; }
0101   __device__ uint16_t length(uint32_t index) const { return length_[index]; }
0102   __device__ stripgpu::fedId_t fedID(uint32_t index) const { return fedID_[index]; }
0103   __device__ stripgpu::fedCh_t fedCh(uint32_t index) const { return fedCh_[index]; }
0104   __device__ stripgpu::detId_t detID(uint32_t index) const { return detID_[index]; }
0105
0106 private:
0107   const uint8_t* const* input_;  // input raw data for channel
0108   size_t* inoff_;                // offset in input raw data
0109   size_t* offset_;               // global offset in alldata
0110   uint16_t* length_;             // length of channel data
0111   stripgpu::fedId_t* fedID_;
0112   stripgpu::fedCh_t* fedCh_;
0113   stripgpu::detId_t* detID_;
0114   size_t size_;
0115 };
0116
0117 class ChannelLocsGPU : public ChannelLocsBase<cms::cuda::device::unique_ptr> {
0118 public:
0119   //using Base = ChannelLocsBase<cms::cuda::device::unique_ptr>;
0120   ChannelLocsGPU(size_t size, cudaStream_t stream);
0121   ChannelLocsGPU(ChannelLocsGPU&& arg)
0122       : ChannelLocsBase(std::move(arg)), channelLocsViewGPU_(std::move(arg.channelLocsViewGPU_)) {}
0123
0124   ChannelLocsGPU(ChannelLocsGPU&) = delete;
0125   ChannelLocsGPU(const ChannelLocsGPU&) = delete;
0126   ChannelLocsGPU& operator=(const ChannelLocsGPU&) = delete;
0127   ChannelLocsGPU& operator=(ChannelLocsGPU&&) = delete;
0128
0129   ~ChannelLocsGPU() override = default;
0130
0131   void setVals(const ChannelLocs* c, cms::cuda::host::unique_ptr<const uint8_t*[]> inputGPU, cudaStream_t stream);
0132   const ChannelLocsView* channelLocsView() const { return channelLocsViewGPU_.get(); }
0133
0134 private:
0135   cms::cuda::device::unique_ptr<ChannelLocsView> channelLocsViewGPU_;
0136 };
0137
0138 #endif