TrackingRecHit/interface/TrackingRecHitSoADevice.h

0001 #ifndef CUDADataFormats_RecHits_TrackingRecHitsDevice_h
0002 #define CUDADataFormats_RecHits_TrackingRecHitsDevice_h
0003
0004 #include <cstdint>
0005
0006 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitsUtilities.h"
0007 #include "CUDADataFormats/Common/interface/PortableDeviceCollection.h"
0008 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
0009 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
0010
0011 template <typename TrackerTraits>
0012 class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>> {
0013 public:
0014   using hitSoA = TrackingRecHitSoA<TrackerTraits>;
0015   //Need to decorate the class with the inherited portable accessors being now a template
0016   using cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>>::view;
0017   using cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>>::const_view;
0018   using cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>>::buffer;
0019   using cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>>::bufferSize;
0020
0021   TrackingRecHitSoADevice() = default;  // cms::cuda::Product needs this
0022
0023   using AverageGeometry = typename hitSoA::AverageGeometry;
0024   using ParamsOnGPU = typename hitSoA::ParamsOnGPU;
0025
0026   // Constructor which specifies the SoA size
0027   explicit TrackingRecHitSoADevice(uint32_t nHits,
0028                                    int32_t offsetBPIX2,
0029                                    ParamsOnGPU const* cpeParams,
0030                                    uint32_t const* hitsModuleStart,
0031                                    cudaStream_t stream)
0032       : cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>>(nHits, stream),
0033         offsetBPIX2_(offsetBPIX2) {
0034     cudaCheck(cudaMemcpyAsync(&(view().nHits()), &nHits, sizeof(uint32_t), cudaMemcpyDefault, stream));
0035     // hitsModuleStart is on Device
0036     cudaCheck(cudaMemcpyAsync(view().hitsModuleStart().data(),
0037                               hitsModuleStart,
0038                               sizeof(uint32_t) * int(TrackerTraits::numberOfModules + 1),
0039                               cudaMemcpyDefault,
0040                               stream));
0041     cudaCheck(cudaMemcpyAsync(&(view().offsetBPIX2()), &offsetBPIX2, sizeof(int32_t), cudaMemcpyDefault, stream));
0042
0043     // cpeParams argument is a pointer to device memory, copy
0044     // its contents into the Layout.
0045     cudaCheck(cudaMemcpyAsync(&(view().cpeParams()), cpeParams, int(sizeof(ParamsOnGPU)), cudaMemcpyDefault, stream));
0046   }
0047
0048   cms::cuda::host::unique_ptr<float[]> localCoordToHostAsync(cudaStream_t stream) const {
0049     auto ret = cms::cuda::make_host_unique<float[]>(4 * nHits(), stream);
0050     size_t rowSize = sizeof(float) * nHits();
0051
0052     size_t srcPitch = ptrdiff_t(view().yLocal()) - ptrdiff_t(view().xLocal());
0053     cudaCheck(
0054         cudaMemcpy2DAsync(ret.get(), rowSize, view().xLocal(), srcPitch, rowSize, 4, cudaMemcpyDeviceToHost, stream));
0055
0056     return ret;
0057   }  //move to utilities
0058
0059   cms::cuda::host::unique_ptr<uint32_t[]> hitsModuleStartToHostAsync(cudaStream_t stream) const {
0060     auto ret = cms::cuda::make_host_unique<uint32_t[]>(TrackerTraits::numberOfModules + 1, stream);
0061     cudaCheck(cudaMemcpyAsync(ret.get(),
0062                               view().hitsModuleStart().data(),
0063                               sizeof(uint32_t) * (TrackerTraits::numberOfModules + 1),
0064                               cudaMemcpyDefault,
0065                               stream));
0066     return ret;
0067   }
0068
0069   uint32_t nHits() const { return view().metadata().size(); }
0070   uint32_t offsetBPIX2() const {
0071     return offsetBPIX2_;
0072   }  //offsetBPIX2 is used on host functions so is useful to have it also stored in the class and not only in the layout
0073 private:
0074   uint32_t offsetBPIX2_ = 0;
0075 };
0076
0077 //Classes definition for Phase1/Phase2, to make the classes_def lighter. Not actually used in the code.
0078 using TrackingRecHitSoADevicePhase1 = TrackingRecHitSoADevice<pixelTopology::Phase1>;
0079 using TrackingRecHitSoADevicePhase2 = TrackingRecHitSoADevice<pixelTopology::Phase2>;
0080 using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice<pixelTopology::HIonPhase1>;
0081
0082 #endif  // CUDADataFormats_Track_TrackHeterogeneousT_H