New fastmultEst and related reworking for UPC

Detectors/ITSMFT/ITS/reconstruction/CMakeLists.txt

@@ -11,15 +11,11 @@
 
 o2_add_library(ITSReconstruction
         SOURCES src/RecoGeomHelper.cxx
-                src/FastMultEstConfig.cxx
-                src/FastMultEst.cxx
         PUBLIC_LINK_LIBRARIES O2::ITSBase
                               O2::ITSMFTReconstruction
                               O2::DataFormatsITS
                               O2::CommonUtils)
 
 o2_target_root_dictionary(
   ITSReconstruction
-  HEADERS include/ITSReconstruction/RecoGeomHelper.h
-          include/ITSReconstruction/FastMultEst.h
-          include/ITSReconstruction/FastMultEstConfig.h)
+        HEADERS include/ITSReconstruction/RecoGeomHelper.h)

Detectors/ITSMFT/ITS/reconstruction/src/ITSReconstructionLinkDef.h

@@ -16,8 +16,5 @@
 #pragma link off all functions;
 
 #pragma link C++ class o2::its::RecoGeomHelper + ;
-#pragma link C++ class o2::its::FastMultEst + ;
-#pragma link C++ class o2::its::FastMultEstConfig + ;
-#pragma link C++ class o2::conf::ConfigurableParamHelper < o2::its::FastMultEstConfig> + ;
 
 #endif

Detectors/ITSMFT/ITS/tracking/CMakeLists.txt

@@ -14,6 +14,8 @@ o2_add_library(ITStracking
                SOURCES src/ClusterLines.cxx
                        src/Cluster.cxx
                        src/Configuration.cxx
+                       src/FastMultEstConfig.cxx
+                       src/FastMultEst.cxx
                        src/TimeFrame.cxx
                        src/IOUtils.cxx
                        src/Tracker.cxx
@@ -28,6 +30,7 @@ o2_add_library(ITStracking
                        O2::DataFormatsITSMFT
                        O2::SimulationDataFormat
                        O2::ITSBase
+                       O2::CommonUtils
                        O2::ITSReconstruction
                        O2::ITSMFTReconstruction
                        O2::DataFormatsITS
@@ -49,6 +52,8 @@ o2_target_root_dictionary(ITStracking
                                   include/ITStracking/Tracklet.h
                                   include/ITStracking/Cluster.h
                                   include/ITStracking/Definitions.h
+                                  include/ITStracking/FastMultEst.h
+                                  include/ITStracking/FastMultEstConfig.h
                                   include/ITStracking/TrackingConfigParam.h
                           LINKDEF src/TrackingLinkDef.h)
 

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

@@ -31,6 +31,7 @@ class TimeFrameGPU final : public TimeFrame<NLayers>
   using typename TimeFrame<NLayers>::IndexTableUtilsN;
   using typename TimeFrame<NLayers>::ROFOverlapTableN;
   using typename TimeFrame<NLayers>::ROFVertexLookupTableN;
+  using typename TimeFrame<NLayers>::ROFMaskTableN;
 
  public:
   TimeFrameGPU() = default;
@@ -107,6 +108,7 @@ class TimeFrameGPU final : public TimeFrame<NLayers>
   IndexTableUtilsN* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; }
   const auto getDeviceROFOverlapTableView() { return mDeviceROFOverlapTableView; }
   const auto getDeviceROFVertexLookupTableView() { return mDeviceROFVertexLookupTableView; }
+  const auto getDeviceROFMaskTableView() { return mDeviceROFMaskTableView; }
   int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; }
   auto& getTrackITSExt() { return mTrackITSExt; }
   Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; }
@@ -177,9 +179,11 @@ class TimeFrameGPU final : public TimeFrame<NLayers>
   // device navigation views
   ROFOverlapTableN::View mDeviceROFOverlapTableView;
   ROFVertexLookupTableN::View mDeviceROFVertexLookupTableView;
+  ROFMaskTableN::View mDeviceROFMaskTableView;
 
   // Hybrid pref
   uint8_t* mMultMaskDevice;
+  int32_t* mMultMaskOffsetsDevice;
   Vertex* mPrimaryVerticesDevice;
   int* mROFramesPVDevice;
   std::array<Cluster*, NLayers> mClustersDevice;

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

@@ -252,11 +252,18 @@ void TimeFrameGPU<NLayers>::loadMultiplicityCutMask(const int iteration)
 {
   if (!iteration || iteration == 3) { // we need to re-load the swapped mult-mask in upc iteration
     GPUTimer timer("loading multiplicity cut mask");
-    GPULog("gpu-transfer: iteration {} loading multiplicity cut mask with {} elements, for {:.2f} MB.", iteration, this->mMultiplicityCutMask.size(), this->mMultiplicityCutMask.size() * sizeof(uint8_t) / constants::MB);
-    if (!iteration) { // only allocate on first call
-      allocMem(reinterpret_cast<void**>(&mMultMaskDevice), this->mMultiplicityCutMask.size() * sizeof(uint8_t), this->hasFrameworkAllocator());
+    const auto& hostTable = this->mMultiplicityCutMask;
+    const auto hostView = hostTable.getView();
+    GPULog("gpu-transfer: iteration {} loading multiplicity cut mask with {} elements, for {:.2f} MB.",
+           iteration, hostTable.getFlatMaskSize(), hostTable.getFlatMaskSize() * sizeof(uint8_t) / constants::MB);
+    if (!iteration) { // only allocate on first call; offsets are stable across iterations
+      allocMem(reinterpret_cast<void**>(&mMultMaskDevice), hostTable.getFlatMaskSize() * sizeof(uint8_t), this->hasFrameworkAllocator());
+      allocMem(reinterpret_cast<void**>(&mMultMaskOffsetsDevice), NLayers * sizeof(int32_t), this->hasFrameworkAllocator());
+      GPUChkErrS(cudaMemcpy(mMultMaskOffsetsDevice, hostView.mLayerROFOffsets, NLayers * sizeof(int32_t), cudaMemcpyHostToDevice));
     }
-    GPUChkErrS(cudaMemcpy(mMultMaskDevice, this->mMultiplicityCutMask.data(), this->mMultiplicityCutMask.size() * sizeof(uint8_t), cudaMemcpyHostToDevice));
+    // Re-copy the flat mask on every qualifying iteration (e.g. after swapMasks() for UPC)
+    GPUChkErrS(cudaMemcpy(mMultMaskDevice, hostView.mFlatMask, hostTable.getFlatMaskSize() * sizeof(uint8_t), cudaMemcpyHostToDevice));
+    mDeviceROFMaskTableView = hostTable.getDeviceView(mMultMaskDevice, mMultMaskOffsetsDevice);
   }
 }