Adapt GPU code to the new mult mask

Author: mpuccio

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

@@ -143,7 +143,6 @@ class TimeFrameGPU final : public TimeFrame<NLayers>
   o2::track::TrackParCovF** getDeviceArrayTrackSeeds() { return mCellSeedsDeviceArray; }
   float** getDeviceArrayTrackSeedsChi2() { return mCellSeedsChi2DeviceArray; }
   int* getDeviceNeighboursIndexTables(const int layer) { return mNeighboursIndexTablesDevice[layer]; }
-  uint8_t* getDeviceMultCutMask() { return mMultMaskDevice; }
 
   void setDevicePropagator(const o2::base::PropagatorImpl<float>* p) final { this->mPropagatorDevice = p; }
 

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

@@ -36,7 +36,7 @@ class ExternalAllocator;
 
 template <int NLayers = 7>
 void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
-                                 const uint8_t* multMask,
+                                 const typename ROFMaskTable<NLayers>::View& multMask,
                                  const int layer,
                                  const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
                                  const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,
@@ -66,7 +66,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
 
 template <int NLayers = 7>
 void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
-                                   const uint8_t* multMask,
+                                   const typename ROFMaskTable<NLayers>::View& multMask,
                                    const int layer,
                                    const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
                                    const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,

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

@@ -85,7 +85,7 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
     mTimeFrameGPU->createTrackletsLUTDevice(iteration, iLayer);
     mTimeFrameGPU->waitEvent(iLayer, iLayer + 1); // wait stream until all data is available
     countTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
-                                         mTimeFrameGPU->getDeviceMultCutMask(),
+                                         mTimeFrameGPU->getDeviceROFMaskTableView(),
                                          iLayer,
                                          mTimeFrameGPU->getDeviceROFOverlapTableView(),
                                          mTimeFrameGPU->getDeviceROFVertexLookupTableView(),
@@ -117,7 +117,7 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
       continue;
     }
     computeTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
-                                           mTimeFrameGPU->getDeviceMultCutMask(),
+                                           mTimeFrameGPU->getDeviceROFMaskTableView(),
                                            iLayer,
                                            mTimeFrameGPU->getDeviceROFOverlapTableView(),
                                            mTimeFrameGPU->getDeviceROFVertexLookupTableView(),

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

@@ -536,7 +536,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerCellsKernel(
 template <bool initRun, int NLayers>
 GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel(
   const IndexTableUtils<NLayers>* utils,
-  const uint8_t* multMask,
+  const typename ROFMaskTable<NLayers>::View multMask,
   const int layerIndex,
   const typename ROFOverlapTable<NLayers>::View rofOverlaps,
   const typename ROFVertexLookupTable<NLayers>::View vertexLUT,
@@ -565,6 +565,10 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel(
   const int totalROFs0 = rofOverlaps.getLayer(layerIndex).mNROFsTF;
   const int totalROFs1 = rofOverlaps.getLayer(layerIndex + 1).mNROFsTF;
   for (unsigned int pivotROF{blockIdx.x}; pivotROF < totalROFs0; pivotROF += gridDim.x) {
+    if (!multMask.isROFEnabled(layerIndex, pivotROF)) {
+      continue;
+    }
+
     const auto& pvs = vertexLUT.getVertices(layerIndex, pivotROF);
     auto primaryVertices = gpuSpan<const Vertex>(&vertices[pvs.getFirstEntry()], pvs.getEntries());
     if (primaryVertices.empty()) {
@@ -782,7 +786,7 @@ GPUg() void __launch_bounds__(256, 1) processNeighboursKernel(
 
 template <int NLayers>
 void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
-                                 const uint8_t* multMask,
+                                 const typename ROFMaskTable<NLayers>::View& multMask,
                                  const int layer,
                                  const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
                                  const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,
@@ -840,7 +844,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
 
 template <int NLayers>
 void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
-                                   const uint8_t* multMask,
+                                   const typename ROFMaskTable<NLayers>::View& multMask,
                                    const int layer,
                                    const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
                                    const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,
@@ -1300,7 +1304,7 @@ void computeTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
 
 /// Explicit instantiation of ITS2 handlers
 template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
-                                             const uint8_t* multMask,
+                                             const ROFMaskTable<7>::View& multMask,
                                              const int layer,
                                              const ROFOverlapTable<7>::View& rofOverlaps,
                                              const ROFVertexLookupTable<7>::View& vertexLUT,
@@ -1329,7 +1333,7 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
                                              gpu::Streams& streams);
 
 template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
-                                               const uint8_t* multMask,
+                                               const ROFMaskTable<7>::View& multMask,
                                                const int layer,
                                                const ROFOverlapTable<7>::View& rofOverlaps,
                                                const ROFVertexLookupTable<7>::View& vertexLUT,

Detectors/ITSMFT/ITS/tracking/include/ITStracking/FastMultEst.h

@@ -37,12 +37,12 @@ struct FastMultEst {
   using ROFOverlapTableN = ROFOverlapTable<NLayers>;
   using ROFMaskTableN = ROFMaskTable<NLayers>;
 
-  float mult = 0.;                         /// estimated signal clusters multiplicity on the selected multiplicity layer
-  float noisePerChip = 0.;                 /// imposed noise per chip (when enabled by configuration)
-  float cov[3] = {0.};                     /// retained for compatibility; set to zero in single-layer mode
-  float chi2 = 0.;                         /// retained for compatibility; set to zero in single-layer mode
-  int nLayersUsed = 0;                     /// number of layers used by estimator (0/1 in single-layer mode)
-  uint32_t lastRandomSeed = 0;             /// state of the gRandom before
+  float mult = 0.;             /// estimated signal clusters multiplicity on the selected multiplicity layer
+  float noisePerChip = 0.;     /// imposed noise per chip (when enabled by configuration)
+  float cov[3] = {0.};         /// retained for compatibility; set to zero in single-layer mode
+  float chi2 = 0.;             /// retained for compatibility; set to zero in single-layer mode
+  int nLayersUsed = 0;         /// number of layers used by estimator (0/1 in single-layer mode)
+  uint32_t lastRandomSeed = 0; /// state of the gRandom before
   FastMultEst();
 
   static uint32_t getCurrentRandomSeed();

Detectors/ITSMFT/ITS/tracking/include/ITStracking/FastMultEstConfig.h

@@ -34,14 +34,14 @@ struct FastMultEstConfig : public o2::conf::ConfigurableParamHelper<FastMultEstC
   float imposeNoisePerChip = 1.e-9 * 1024 * 512; // assumed noise, free parameter if<0
 
   // cuts to reject to low or too high mult events
-  int cutMultClusLayer = NLayers - 1;   /// layer used for cluster multiplicity selection (by default the outermost one)
-  float cutMultClusLow = 0;   /// reject ROF with estimated cluster mult. below this value (no cut if <0)
-  float cutMultClusHigh = -1; /// reject ROF with estimated cluster mult. above this value (no cut if <0)
-  float cutMultVtxLow = -1;   /// reject seed vertex if its multiplicity below this value (no cut if <0)
-  float cutMultVtxHigh = -1;  /// reject seed vertex if its multiplicity above this value (no cut if <0)
-  float cutRandomFraction = -1.; /// apply random cut rejecting requested fraction
-  int randomSeed = 0;            /// 0 - do not seet seed, >0 : set as is, <0 : use current time
-  bool preferTriggered = true;   /// prefer ROFs with highest number of physics triggers
+  int cutMultClusLayer = NLayers - 1; /// layer used for cluster multiplicity selection (by default the outermost one)
+  float cutMultClusLow = 0;           /// reject ROF with estimated cluster mult. below this value (no cut if <0)
+  float cutMultClusHigh = -1;         /// reject ROF with estimated cluster mult. above this value (no cut if <0)
+  float cutMultVtxLow = -1;           /// reject seed vertex if its multiplicity below this value (no cut if <0)
+  float cutMultVtxHigh = -1;          /// reject seed vertex if its multiplicity above this value (no cut if <0)
+  float cutRandomFraction = -1.;      /// apply random cut rejecting requested fraction
+  int randomSeed = 0;                 /// 0 - do not seet seed, >0 : set as is, <0 : use current time
+  bool preferTriggered = true;        /// prefer ROFs with highest number of physics triggers
 
   bool isMultCutRequested() const { return cutMultClusLow >= 0.f && cutMultClusHigh > 0.f; };
   bool isVtxMultCutRequested() const { return cutMultVtxLow >= 0.f && cutMultVtxHigh > 0.f; };

Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx

@@ -214,7 +214,7 @@ void ITSTrackingInterface::run(framework::ProcessingContext& pc)
           if (vtxSpan[0].isFlagSet(Vertex::UPCMode) == 1) { // at least one vertex in this ROF and it is from second vertex iteration
             LOGP(debug, "ROF {} rejected as vertices are from the UPC iteration", iRof);
             processUPCMask.selectROF({clockTiming.getROFStartInBC(iRof), clockTiming.getROFEndInBC(iRof)});
-              vtxROF.setFlag(o2::itsmft::ROFRecord::VtxUPCMode);
+            vtxROF.setFlag(o2::itsmft::ROFRecord::VtxUPCMode);
           } else { // in all cases except if as standard mode vertex was found, the ROF was processed with UPC settings
             vtxROF.setFlag(o2::itsmft::ROFRecord::VtxStdMode);
           }