File indexing completed on 2023-03-17 10:49:13
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
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;
0022
0023 using AverageGeometry = typename hitSoA::AverageGeometry;
0024 using ParamsOnGPU = typename hitSoA::ParamsOnGPU;
0025
0026
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
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
0044
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 }
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 }
0073 private:
0074 uint32_t offsetBPIX2_ = 0;
0075 };
0076
0077
0078 using TrackingRecHitSoADevicePhase1 = TrackingRecHitSoADevice<pixelTopology::Phase1>;
0079 using TrackingRecHitSoADevicePhase2 = TrackingRecHitSoADevice<pixelTopology::Phase2>;
0080
0081 #endif