Move multiplicity mask to a vector<uint8_t>

Author: mconcas

Detectors/ITSMFT/ITS/reconstruction/include/ITSReconstruction/FastMultEst.h

@@ -45,7 +45,7 @@ struct FastMultEst {
 
   static uint32_t getCurrentRandomSeed();
   int selectROFs(const gsl::span<const o2::itsmft::ROFRecord> rofs, const gsl::span<const o2::itsmft::CompClusterExt> clus,
-                 const gsl::span<const o2::itsmft::PhysTrigger> trig, std::vector<bool>& sel);
+                 const gsl::span<const o2::itsmft::PhysTrigger> trig, std::vector<uint8_t>& sel);
 
   void fillNClPerLayer(const gsl::span<const o2::itsmft::CompClusterExt>& clusters);
   float process(const std::array<int, NLayers> ncl)

Detectors/ITSMFT/ITS/reconstruction/src/FastMultEst.cxx

@@ -125,7 +125,7 @@ float FastMultEst::processNoiseImposed(const std::array<int, NLayers> ncl)
 }
 
 int FastMultEst::selectROFs(const gsl::span<const o2::itsmft::ROFRecord> rofs, const gsl::span<const o2::itsmft::CompClusterExt> clus,
-                            const gsl::span<const o2::itsmft::PhysTrigger> trig, std::vector<bool>& sel)
+                            const gsl::span<const o2::itsmft::PhysTrigger> trig, std::vector<uint8_t>& sel)
 {
   int nrof = rofs.size(), nsel = 0;
   const auto& multEstConf = FastMultEstConfig::Instance(); // parameters for mult estimation and cuts

Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h

@@ -54,6 +54,8 @@ class TimeFrameGPU : public TimeFrame
   void loadTrackingFrameInfoDevice(const int);
   void loadUnsortedClustersDevice(const int);
   void loadClustersDevice(const int);
+  void loadMultiplicityCutMask(const int);
+  ///
   void loadTrackletsDevice();
   void loadTrackletsLUTDevice();
   void loadCellsDevice();
@@ -116,6 +118,7 @@ class TimeFrameGPU : public TimeFrame
   o2::track::TrackParCovF** getDeviceArrayTrackSeeds() { return mCellSeedsDeviceArray; }
   float** getDeviceArrayTrackSeedsChi2() { return mCellSeedsChi2DeviceArray; }
   int* getDeviceNeighboursIndexTables(const int layer) { return mNeighboursIndexTablesDevice[layer]; }
+  bool* getDevicemMultMask() { return mMultMaskDevice; }
 
   void setDevicePropagator(const o2::base::PropagatorImpl<float>*) override;
 
@@ -147,6 +150,7 @@ class TimeFrameGPU : public TimeFrame
   int* mROFramesPVDevice;
 
   // Hybrid pref
+  bool* mMultMaskDevice;
   std::array<Cluster*, nLayers> mClustersDevice;
   std::array<Cluster*, nLayers> mUnsortedClustersDevice;
   const Cluster** mClustersDeviceArray;
@@ -186,10 +190,6 @@ class TimeFrameGPU : public TimeFrame
   std::vector<std::vector<int>> mNVerticesInChunks;
   std::vector<std::vector<o2::MCCompLabel>> mLabelsInChunks;
 
-  // Host memory used only in GPU tracking
-  std::vector<int> mHostNTracklets;
-  std::vector<int> mHostNCells;
-
   // Temporary buffer for storing output tracks from GPU tracking
   std::vector<TrackITSExt> mTrackITSExt;
 };

Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h

@@ -52,7 +52,10 @@ GPUg() void fitTrackSeedsKernel(
 
 template <int nLayers = 7>
 void computeTrackletsInRofsHandler(const int startROF,
-                                   const int endROF);
+                                   const int endROF,
+                                   const int vertexId,
+                                   const int nBlocks,
+                                   const int nThreads);
 
 void countCellsHandler(const Cluster** sortedClusters,
                        const Cluster** unsortedClusters,

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu

@@ -146,6 +146,17 @@ void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice(const int iteration)
   STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
 }
 
+template <int nLayers>
+void TimeFrameGPU<nLayers>::loadMultiplicityCutMask(const int iteration)
+{
+  START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "loading multiplicity cut mask");
+  if (!iteration) {
+    LOGP(info, "gpu-transfer: loading multiplicity cut mask with {} elements, for {} MB.", mMultiplicityCutMask.size(), mMultiplicityCutMask.size() * sizeof(bool) / MB);
+    allocMemAsync(reinterpret_cast<void**>(&mMultMaskDevice), mMultiplicityCutMask.size() * sizeof(bool), nullptr, getExtAllocator());
+    checkGPUError(cudaMemcpyAsync(mMultMaskDevice, mMultiplicityCutMask.data(), mMultiplicityCutMask.size() * sizeof(uint8_t), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
+  }
+  STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
+}
 template <int nLayers>
 void TimeFrameGPU<nLayers>::loadTrackletsDevice()
 {

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx

@@ -32,6 +32,7 @@ void TrackerTraitsGPU<nLayers>::initialiseTimeFrame(const int iteration)
   mTimeFrameGPU->loadClustersDevice(iteration);
   mTimeFrameGPU->loadUnsortedClustersDevice(iteration);
   mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration);
+  mTimeFrameGPU->loadMultiplicityCutMask(iteration);
 }
 
 template <int nLayers>
@@ -84,13 +85,18 @@ int TrackerTraitsGPU<nLayers>::getTFNumberOfCells() const
 template <int nLayers>
 void TrackerTraitsGPU<nLayers>::computeTrackletsHybrid(const int iteration, int iROFslice, int iVertex)
 {
+  auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
   TrackerTraits::computeLayerTracklets(iteration, iROFslice, iVertex);
 
   const Vertex diamondVert({mTrkParams[iteration].Diamond[0], mTrkParams[iteration].Diamond[1], mTrkParams[iteration].Diamond[2]}, {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f}, 1, 1.f);
   gsl::span<const Vertex> diamondSpan(&diamondVert, 1);
   int startROF{mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * mTrkParams[iteration].nROFsPerIterations : 0};
   int endROF{mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * mTrkParams[iteration].nROFsPerIterations + mTrkParams[iteration].DeltaROF : mTimeFrameGPU->getNrof()};
-  computeTrackletsInRofsHandler<nLayers>(startROF, endROF);
+  computeTrackletsInRofsHandler<nLayers>(startROF,
+                                         endROF,
+                                         iVertex,
+                                         conf.nBlocks,
+                                         conf.nThreads);
 }
 
 template <int nLayers>