Merge pull request #40575 from PixelTracksAlpaka/portable_hits_uninit…

…_fix TrackingRecHitSoA Classes Clean Up
cms-sw · Feb 4, 2023 · 3e10dc1 · 3e10dc1
2 parents 2c5a031 + 2b5ba63
commit 3e10dc1
Show file tree

Hide file tree

Showing 7 changed files with 16 additions and 41 deletions.
diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h
@@ -22,20 +22,15 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection<Track
 
   using AverageGeometry = typename hitSoA::AverageGeometry;
   using ParamsOnGPU = typename hitSoA::ParamsOnGPU;
-  using PhiBinnerStorageType = typename hitSoA::PhiBinnerStorageType;
-  using PhiBinner = typename hitSoA::PhiBinner;
+
   // Constructor which specifies the SoA size
   explicit TrackingRecHitSoADevice(uint32_t nHits,
                                    int32_t offsetBPIX2,
                                    ParamsOnGPU const* cpeParams,
                                    uint32_t const* hitsModuleStart,
                                    cudaStream_t stream)
       : cms::cuda::PortableDeviceCollection<TrackingRecHitLayout<TrackerTraits>>(nHits, stream),
-        nHits_(nHits),
-        cpeParams_(cpeParams),
-        hitsModuleStart_(hitsModuleStart),
         offsetBPIX2_(offsetBPIX2) {
-    phiBinner_ = &(view().phiBinner());
     cudaCheck(cudaMemcpyAsync(&(view().nHits()), &nHits, sizeof(uint32_t), cudaMemcpyDefault, stream));
     // hitsModuleStart is on Device
     cudaCheck(cudaMemcpyAsync(view().hitsModuleStart().data(),
@@ -50,8 +45,6 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection<Track
     cudaCheck(cudaMemcpyAsync(&(view().cpeParams()), cpeParams, int(sizeof(ParamsOnGPU)), cudaMemcpyDefault, stream));
   }
 
-  uint32_t nHits() const { return nHits_; }  //go to size of view
-
   cms::cuda::host::unique_ptr<float[]> localCoordToHostAsync(cudaStream_t stream) const {
     auto ret = cms::cuda::make_host_unique<float[]>(4 * nHits(), stream);
     size_t rowSize = sizeof(float) * nHits();
@@ -70,21 +63,12 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection<Track
     return ret;
   }
 
-  auto phiBinnerStorage() { return phiBinnerStorage_; }
-  auto hitsModuleStart() const { return hitsModuleStart_; }
-  uint32_t offsetBPIX2() const { return offsetBPIX2_; }
-  auto phiBinner() { return phiBinner_; }
-
+  uint32_t nHits() const { return view().metadata().size(); }
+  uint32_t offsetBPIX2() const {
+    return offsetBPIX2_;
+  }  //offsetBPIX2 is used on host functions so is useful to have it also stored in the class and not only in the layout
 private:
-  uint32_t nHits_;  //Needed for the host SoA size
-
-  //TODO: this is used not that much from the hits (only once in BrokenLineFit), would make sens to remove it from this class.
-  ParamsOnGPU const* cpeParams_;
-  uint32_t const* hitsModuleStart_;
-  uint32_t offsetBPIX2_;
-
-  PhiBinnerStorageType* phiBinnerStorage_;
-  PhiBinner* phiBinner_;
+  uint32_t offsetBPIX2_ = 0;
 };
 
 //Classes definition for Phase1/Phase2, to make the classes_def lighter. Not actually used in the code.

diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h
@@ -37,10 +37,7 @@ class TrackingRecHitSoAHost : public cms::cuda::PortableHostCollection<TrackingR
                                  int32_t offsetBPIX2,
                                  ParamsOnGPU const* cpeParams,
                                  uint32_t const* hitsModuleStart)
-      : cms::cuda::PortableHostCollection<TrackingRecHitLayout<TrackerTraits>>(nHits),
-        nHits_(nHits),
-        cpeParams_(cpeParams),
-        offsetBPIX2_(offsetBPIX2) {
+      : cms::cuda::PortableHostCollection<TrackingRecHitLayout<TrackerTraits>>(nHits), offsetBPIX2_(offsetBPIX2) {
     view().nHits() = nHits;
     std::copy(hitsModuleStart, hitsModuleStart + TrackerTraits::numberOfModules + 1, view().hitsModuleStart().begin());
     memcpy(&(view().cpeParams()), cpeParams, sizeof(ParamsOnGPU));
@@ -53,25 +50,19 @@ class TrackingRecHitSoAHost : public cms::cuda::PortableHostCollection<TrackingR
                                  uint32_t const* hitsModuleStart,
                                  cudaStream_t stream)
       : cms::cuda::PortableHostCollection<TrackingRecHitLayout<TrackerTraits>>(nHits, stream),
-        nHits_(nHits),
-        cpeParams_(cpeParams),
         offsetBPIX2_(offsetBPIX2) {
     view().nHits() = nHits;
     std::copy(hitsModuleStart, hitsModuleStart + TrackerTraits::numberOfModules + 1, view().hitsModuleStart().begin());
     memcpy(&(view().cpeParams()), cpeParams, sizeof(ParamsOnGPU));
     view().offsetBPIX2() = offsetBPIX2;
   }
 
-  uint32_t nHits() const { return nHits_; }
-  uint32_t offsetBPIX2() const { return offsetBPIX2_; }
-  auto phiBinnerStorage() { return phiBinnerStorage_; }
-
+  uint32_t nHits() const { return view().metadata().size(); }
+  uint32_t offsetBPIX2() const {
+    return offsetBPIX2_;
+  }  //offsetBPIX2 is used on host functions so is useful to have it also stored in the class and not only in the layout
 private:
-  uint32_t nHits_;  //Needed for the host SoA size
-  ParamsOnGPU const* cpeParams_;
-  uint32_t offsetBPIX2_;
-
-  PhiBinnerStorageType* phiBinnerStorage_;
+  uint32_t offsetBPIX2_ = 0;
 };
 
 using TrackingRecHitSoAHostPhase1 = TrackingRecHitSoAHost<pixelTopology::Phase1>;

diff --git a/CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu b/CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu
@@ -43,7 +43,7 @@ namespace testTrackingRecHitSoA {
     fill<TrackerTraits><<<10, 100, 0, stream>>>(hits.view());
 
     cudaCheck(cudaDeviceSynchronize());
-    cms::cuda::fillManyFromVector(hits.phiBinner(),
+    cms::cuda::fillManyFromVector(&(hits.view().phiBinner()),
                                   10,
                                   hits.view().iphi(),
                                   hits.view().hitsLayerStart().data(),

diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu
@@ -76,7 +76,7 @@ namespace pixelgpudetails {
             <<<1, 32, 0, stream>>>(clusters_d->clusModuleStart(), cpeParams, hits_d.view().hitsLayerStart().data());
         cudaCheck(cudaGetLastError());
         constexpr auto nLayers = TrackerTraits::numberOfLayers;
-        cms::cuda::fillManyFromVector(hits_d.phiBinner(),
+        cms::cuda::fillManyFromVector(&(hits_d.view().phiBinner()),
                                       nLayers,
                                       hits_d.view().iphi(),
                                       hits_d.view().hitsLayerStart().data(),

diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu
@@ -161,7 +161,6 @@ void CAHitNtupletGeneratorKernelsGPU<TrackerTraits>::buildDoublets(const HitsCon
                                                                    int32_t offsetBPIX2,
                                                                    cudaStream_t stream) {
   int32_t nhits = hh.metadata().size();
-
   using namespace gpuPixelDoublets;
 
   using GPUCACell = GPUCACellT<TrackerTraits>;

diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.h
@@ -236,7 +236,7 @@ class CAHitNtupletGeneratorKernels {
 
   void classifyTuples(const HitsConstView& hh, TkSoAView& track_view, cudaStream_t cudaStream);
 
-  void buildDoublets(const HitsConstView& hh, int32_t offsetBPIX2, cudaStream_t stream);
+  void buildDoublets(const HitsConstView& hh, cudaStream_t stream);
   void allocateOnGPU(int32_t nHits, cudaStream_t stream);
   void cleanup(cudaStream_t cudaStream);
 

diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h
@@ -5,6 +5,7 @@
 
 #include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h"
 #include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h"
+
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h"
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h"
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitsUtilities.h"