ITS: template Vertexer&Traits, IndexTableUtils

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

@@ -28,6 +28,7 @@ template <int nLayers = 7>
 class TimeFrameGPU : public TimeFrame<nLayers>
 {
   using typename TimeFrame<nLayers>::CellSeedN;
+  using typename TimeFrame<nLayers>::IndexTableUtilsN;
 
  public:
   TimeFrameGPU();
@@ -36,8 +37,8 @@ class TimeFrameGPU : public TimeFrame<nLayers>
   /// Most relevant operations
   void registerHostMemory(const int);
   void unregisterHostMemory(const int);
-  void initialise(const int, const TrackingParameters&, const int, IndexTableUtils* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr);
-  void initDevice(IndexTableUtils*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int);
+  void initialise(const int, const TrackingParameters&, const int, IndexTableUtilsN* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr);
+  void initDevice(IndexTableUtilsN*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int);
   void initDeviceSAFitting();
   void loadIndexTableUtils(const int);
   void loadTrackingFrameInfoDevice(const int, const int);
@@ -98,7 +99,7 @@ class TimeFrameGPU : public TimeFrame<nLayers>
 
   /// interface
   int getNClustersInRofSpan(const int, const int, const int) const;
-  IndexTableUtils* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; }
+  IndexTableUtilsN* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; }
   int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; }
   auto& getTrackITSExt() { return mTrackITSExt; }
   Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; }
@@ -165,7 +166,7 @@ class TimeFrameGPU : public TimeFrame<nLayers>
   std::array<int, nLayers - 3> mNNeighbours;
 
   // Device pointers
-  IndexTableUtils* mIndexTableUtilsDevice;
+  IndexTableUtilsN* mIndexTableUtilsDevice;
 
   // Hybrid pref
   uint8_t* mMultMaskDevice;

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

@@ -22,6 +22,8 @@ namespace o2::its
 template <int nLayers = 7>
 class TrackerTraitsGPU final : public TrackerTraits<nLayers>
 {
+  using typename TrackerTraits<nLayers>::IndexTableUtilsN;
+
  public:
   TrackerTraitsGPU() = default;
   ~TrackerTraitsGPU() final = default;
@@ -48,7 +50,7 @@ class TrackerTraitsGPU final : public TrackerTraits<nLayers>
   int getTFNumberOfCells() const override;
 
  private:
-  IndexTableUtils* mDeviceIndexTableUtils;
+  IndexTableUtilsN* mDeviceIndexTableUtils;
   gpu::TimeFrameGPU<nLayers>* mTimeFrameGPU;
 };
 

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

@@ -27,13 +27,14 @@ template <int>
 class CellSeed;
 class TrackingFrameInfo;
 class Tracklet;
+template <int>
 class IndexTableUtils;
 class Cluster;
 class TrackITSExt;
 class ExternalAllocator;
 
 template <int nLayers = 7>
-void countTrackletsInROFsHandler(const IndexTableUtils* utils,
+void countTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
                                  const uint8_t* multMask,
                                  const int layer,
                                  const int startROF,
@@ -66,7 +67,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils,
                                  gpu::Streams& streams);
 
 template <int nLayers = 7>
-void computeTrackletsInROFsHandler(const IndexTableUtils* utils,
+void computeTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
                                    const uint8_t* multMask,
                                    const int layer,
                                    const int startROF,

Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt

@@ -22,8 +22,8 @@ if(CUDA_ENABLED)
                            TimeFrameGPU.cu
                            TracerGPU.cu
                            TrackingKernels.cu
-                           VertexingKernels.cu
-                           VertexerTraitsGPU.cxx
+                           # VertexingKernels.cu
+                           # VertexerTraitsGPU.cxx
                    PUBLIC_INCLUDE_DIRECTORIES ../
                    PUBLIC_LINK_LIBRARIES O2::ITStracking
                                          O2::SimConfig

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

@@ -61,11 +61,11 @@ void TimeFrameGPU<nLayers>::loadIndexTableUtils(const int iteration)
 {
   GPUTimer timer("loading indextable utils");
   if (!iteration) {
-    GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtils) / constants::MB);
-    allocMem(reinterpret_cast<void**>(&mIndexTableUtilsDevice), sizeof(IndexTableUtils), this->getExtAllocator());
+    GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB);
+    allocMem(reinterpret_cast<void**>(&mIndexTableUtilsDevice), sizeof(IndexTableUtilsN), this->getExtAllocator());
   }
-  GPULog("gpu-transfer: loading IndexTableUtils object, for {:.2f} MB.", sizeof(IndexTableUtils) / constants::MB);
-  GPUChkErrS(cudaMemcpy(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtils), cudaMemcpyHostToDevice));
+  GPULog("gpu-transfer: loading IndexTableUtils object, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB);
+  GPUChkErrS(cudaMemcpy(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtilsN), cudaMemcpyHostToDevice));
 }
 
 template <int nLayers>
@@ -547,7 +547,7 @@ template <int nLayers>
 void TimeFrameGPU<nLayers>::initialise(const int iteration,
                                        const TrackingParameters& trkParam,
                                        const int maxLayers,
-                                       IndexTableUtils* utils,
+                                       IndexTableUtilsN* utils,
                                        const TimeFrameGPUParameters* gpuParam)
 {
   mGpuStreams.resize(nLayers);

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

@@ -95,8 +95,9 @@ GPUdii() int4 getEmptyBinsRect()
   return int4{0, 0, 0, 0};
 }
 
+template <int nLayers>
 GPUdii() const int4 getBinsRect(const Cluster& currentCluster, const int layerIndex,
-                                const o2::its::IndexTableUtils& utils,
+                                const IndexTableUtils<nLayers>& utils,
                                 const float z1, const float z2, float maxdeltaz, float maxdeltaphi)
 {
   const float zRangeMin = o2::gpu::CAMath::Min(z1, z2) - maxdeltaz;
@@ -331,7 +332,7 @@ GPUg() void fitTrackSeedsKernel(
     temporaryTrack.resetCovariance();
     temporaryTrack.setChi2(0);
     auto& clusters = seed.getClusters();
-    for (int iL{0}; iL < 7; ++iL) {
+    for (int iL{0}; iL < nLayers; ++iL) {
       temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::UnusedIndex);
     }
     bool fitSuccess = fitTrack(temporaryTrack,               // TrackITSExt& track,
@@ -523,9 +524,9 @@ GPUg() void computeLayerCellsKernel(
   }
 }
 
-template <bool initRun>
+template <bool initRun, int nLayers>
 GPUg() void computeLayerTrackletsMultiROFKernel(
-  const IndexTableUtils* utils,
+  const IndexTableUtils<nLayers>* utils,
   const uint8_t* multMask,
   const int layerIndex,
   const int startROF,
@@ -601,7 +602,7 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
         const float zAtRmax{tanLambda * (maxR - currentCluster.radius) + currentCluster.zCoordinate};
         const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + constants::Tolerance)}; /// protecting from overflows adding the detector resolution
         const float sigmaZ{o2::gpu::CAMath::Sqrt(math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInverseDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * MSAngle))};
-        const int4 selectedBinsRect{getBinsRect(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)};
+        const int4 selectedBinsRect{getBinsRect<nLayers>(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)};
         if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {
           continue;
         }
@@ -769,7 +770,7 @@ GPUhi() void deallocateMemory(void* p, size_t bytes, cudaStream_t stream = nullp
 } // namespace gpu
 
 template <int nLayers>
-void countTrackletsInROFsHandler(const IndexTableUtils* utils,
+void countTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
                                  const uint8_t* multMask,
                                  const int layer,
                                  const int startROF,
@@ -833,7 +834,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils,
 }
 
 template <int nLayers>
-void computeTrackletsInROFsHandler(const IndexTableUtils* utils,
+void computeTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
                                    const uint8_t* multMask,
                                    const int layer,
                                    const int startROF,
@@ -1241,7 +1242,7 @@ void trackSeedHandler(CellSeed<nLayers>* trackSeeds,
 }
 
 /// Explicit instantiation of ITS2 handlers
-template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
+template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
                                              const uint8_t* multMask,
                                              const int layer,
                                              const int startROF,
@@ -1273,7 +1274,7 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
                                              const int nThreads,
                                              gpu::Streams& streams);
 
-template void computeTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
+template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
                                                const uint8_t* multMask,
                                                const int layer,
                                                const int startROF,

Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt

@@ -19,8 +19,8 @@ if(HIP_ENABLED)
                          ../cuda/TrackerTraitsGPU.cxx
                          ../cuda/TracerGPU.cu
                          ../cuda/TrackingKernels.cu
-                         ../cuda/VertexingKernels.cu
-                         ../cuda/VertexerTraitsGPU.cxx
+                         # ../cuda/VertexingKernels.cu
+                         # ../cuda/VertexerTraitsGPU.cxx
                  PUBLIC_INCLUDE_DIRECTORIES ../
                  PUBLIC_LINK_LIBRARIES O2::ITStracking
                                        O2::GPUTracking

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

@@ -22,32 +22,23 @@
 namespace o2::its
 {
 
+template <int>
 class IndexTableUtils;
 
 struct Cluster final {
   GPUhdDefault() Cluster() = default;
   GPUhd() Cluster(const float x, const float y, const float z, const int idx);
-  GPUhd() Cluster(const int, const IndexTableUtils& utils, const Cluster&);
-  GPUhd() Cluster(const int, const float3&, const IndexTableUtils& utils, const Cluster&);
+  template <int nLayers>
+  GPUhd() Cluster(const int, const IndexTableUtils<nLayers>& utils, const Cluster&);
+  template <int nLayers>
+  GPUhd() Cluster(const int, const float3&, const IndexTableUtils<nLayers>& utils, const Cluster&);
   GPUhdDefault() Cluster(const Cluster&) = default;
   GPUhdDefault() Cluster(Cluster&&) noexcept = default;
   GPUhdDefault() ~Cluster() = default;
 
   GPUhdDefault() Cluster& operator=(const Cluster&) = default;
   GPUhdDefault() Cluster& operator=(Cluster&&) noexcept = default;
-
-  // TODO
-  /*GPUhdDefault() bool operator==(const Cluster&) const = default;*/
-  GPUhd() bool operator==(const Cluster& other) const
-  {
-    return xCoordinate == other.xCoordinate &&
-           yCoordinate == other.yCoordinate &&
-           zCoordinate == other.zCoordinate &&
-           phi == other.phi &&
-           radius == other.radius &&
-           clusterId == other.clusterId &&
-           indexTableBinIndex == other.indexTableBinIndex;
-  }
+  GPUhdDefault() bool operator==(const Cluster&) const = default;
 
   GPUhd() void print() const;
 

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

@@ -16,16 +16,19 @@
 #ifndef TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_
 #define TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_
 
+#include <array>
+
+#include "ITStracking/Constants.h"
 #include "ITStracking/Configuration.h"
 #include "ITStracking/Definitions.h"
 #include "CommonConstants/MathConstants.h"
 #include "GPUCommonMath.h"
 #include "GPUCommonDef.h"
 
-namespace o2
-{
-namespace its
+namespace o2::its
 {
+
+template <int nLayers>
 class IndexTableUtils
 {
  public:
@@ -48,12 +51,13 @@ class IndexTableUtils
   int mNzBins = 0;
   int mNphiBins = 0;
   float mInversePhiBinSize = 0.f;
-  float mLayerZ[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
-  float mInverseZBinSize[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
+  std::array<float, nLayers> mLayerZ{};
+  std::array<float, nLayers> mInverseZBinSize{};
 };
 
+template <int nLayers>
 template <class T>
-inline void IndexTableUtils::setTrackingParameters(const T& params)
+inline void IndexTableUtils<nLayers>::setTrackingParameters(const T& params)
 {
   mInversePhiBinSize = params.PhiBins / o2::constants::math::TwoPI;
   mNzBins = params.ZBins;
@@ -66,43 +70,48 @@ inline void IndexTableUtils::setTrackingParameters(const T& params)
   }
 }
 
-inline float IndexTableUtils::getInverseZCoordinate(const int layerIndex) const
+template <int nLayers>
+inline float IndexTableUtils<nLayers>::getInverseZCoordinate(const int layerIndex) const
 {
   return 0.5f * mNzBins / mLayerZ[layerIndex];
 }
 
-GPUhdi() int IndexTableUtils::getZBinIndex(const int layerIndex, const float zCoordinate) const
+template <int nLayers>
+GPUhdi() int IndexTableUtils<nLayers>::getZBinIndex(const int layerIndex, const float zCoordinate) const
 {
   return (zCoordinate + mLayerZ[layerIndex]) * mInverseZBinSize[layerIndex];
 }
 
-GPUhdi() int IndexTableUtils::getPhiBinIndex(const float currentPhi) const
+template <int nLayers>
+GPUhdi() int IndexTableUtils<nLayers>::getPhiBinIndex(const float currentPhi) const
 {
   return (currentPhi * mInversePhiBinSize);
 }
 
-GPUhdi() int IndexTableUtils::getBinIndex(const int zIndex, const int phiIndex) const
+template <int nLayers>
+GPUhdi() int IndexTableUtils<nLayers>::getBinIndex(const int zIndex, const int phiIndex) const
 {
   return o2::gpu::GPUCommonMath::Min(phiIndex * mNzBins + zIndex, mNzBins * mNphiBins - 1);
 }
 
-GPUhdi() int IndexTableUtils::countRowSelectedBins(const int* indexTable, const int phiBinIndex,
-                                                   const int minZBinIndex, const int maxZBinIndex) const
+template <int nLayers>
+GPUhdi() int IndexTableUtils<nLayers>::countRowSelectedBins(const int* indexTable, const int phiBinIndex,
+                                                            const int minZBinIndex, const int maxZBinIndex) const
 {
   const int firstBinIndex{getBinIndex(minZBinIndex, phiBinIndex)};
   const int maxBinIndex{firstBinIndex + maxZBinIndex - minZBinIndex + 1};
 
   return indexTable[maxBinIndex] - indexTable[firstBinIndex];
 }
 
-GPUhdi() void IndexTableUtils::print() const
+template <int nLayers>
+GPUhdi() void IndexTableUtils<nLayers>::print() const
 {
   printf("NzBins: %d, NphiBins: %d, InversePhiBinSize: %f\n", mNzBins, mNphiBins, mInversePhiBinSize);
-  for (int iLayer{0}; iLayer < 7; ++iLayer) {
+  for (int iLayer{0}; iLayer < nLayers; ++iLayer) {
     printf("Layer %d: Z: %f, InverseZBinSize: %f\n", iLayer, mLayerZ[iLayer], mInverseZBinSize[iLayer]);
   }
 }
-} // namespace its
-} // namespace o2
 
+} // namespace o2::its
 #endif /* TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_ */

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

@@ -65,6 +65,7 @@ class TimeFrameGPU;
 
 template <int nLayers = 7>
 struct TimeFrame {
+  using IndexTableUtilsN = IndexTableUtils<nLayers>;
   using CellSeedN = CellSeed<nLayers>;
   friend class gpu::TimeFrameGPU<nLayers>;
 
@@ -273,7 +274,7 @@ struct TimeFrame {
   void printCellLUTs();
   void printSliceInfo(const int, const int);
 
-  IndexTableUtils mIndexTableUtils;
+  IndexTableUtilsN mIndexTableUtils;
 
   bool mIsGPU = false;
 

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

@@ -21,6 +21,7 @@
 #include "DetectorsBase/Propagator.h"
 #include "ITStracking/Configuration.h"
 #include "ITStracking/MathUtils.h"
+#include "ITStracking/IndexTableUtils.h"
 #include "ITStracking/TimeFrame.h"
 #include "ITStracking/Cell.h"
 #include "ITStracking/BoundedAllocator.h"
@@ -40,9 +41,10 @@ class TrackITSExt;
 template <int nLayers = 7>
 class TrackerTraits
 {
+ public:
+  using IndexTableUtilsN = IndexTableUtils<nLayers>;
   using CellSeedN = CellSeed<nLayers>;
 
- public:
   virtual ~TrackerTraits() = default;
   virtual void adoptTimeFrame(TimeFrame<nLayers>* tf) { mTimeFrame = tf; }
   virtual void initialiseTimeFrame(const int iteration) { mTimeFrame->initialise(iteration, mTrkParams[iteration], mTrkParams[iteration].NLayers); }
@@ -119,7 +121,7 @@ inline const int4 TrackerTraits<nLayers>::getBinsRect(const int layerIndex, floa
     return getEmptyBinsRect();
   }
 
-  const IndexTableUtils& utils{mTimeFrame->mIndexTableUtils};
+  const IndexTableUtilsN& utils{mTimeFrame->mIndexTableUtils};
   return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex, zRangeMin)),
               utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMin)),
               o2::gpu::GPUCommonMath::Min(mTrkParams[0].ZBins - 1, utils.getZBinIndex(layerIndex, zRangeMax)), // /!\ trkParams can potentially change across iterations