GPU: Change some defines to constexpr variables, some renaming

Author: davidrohr

Detectors/TPC/calibration/src/CorrectdEdxDistortions.cxx

@@ -87,7 +87,7 @@ float o2::tpc::CorrectdEdxDistortions::getCorrection(const float time, unsigned
   const float ly = mTPCGeometry.LinearPad2Y(sector, padrow, pad);
 
   // get correction at "pad + 0.5*padlength" pos1 and dont extrapolate/interpolate across GEM gaps
-  const int row1 = ((padrow == mTPCGeometry.EndIROC() - 1) || (padrow == mTPCGeometry.EndOROC1() - 1) || (padrow == mTPCGeometry.EndOROC2() - 1)) ? padrow : std::clamp(padrow + 1, 0, GPUCA_NROWS - 1);
+  const int row1 = ((padrow == mTPCGeometry.EndIROC() - 1) || (padrow == mTPCGeometry.EndOROC1() - 1) || (padrow == mTPCGeometry.EndOROC2() - 1)) ? padrow : std::clamp(padrow + 1, 0, o2::tpc::constants::MAXGLOBALPADROW - 1);
 
   float lxT_1 = 0;
   float lyT_1 = 0;
@@ -101,7 +101,7 @@ float o2::tpc::CorrectdEdxDistortions::getCorrection(const float time, unsigned
   const float r_1_f = std::sqrt(lxT_1 * lxT_1 + lyT_1 * lyT_1);
 
   // get correction at "pad - 0.5*padlength" pos0 and dont extrapolate/interpolate across GEM gaps
-  const int row0 = ((padrow == mTPCGeometry.EndIROC()) || (padrow == mTPCGeometry.EndOROC1()) || (padrow == mTPCGeometry.EndOROC2())) ? padrow : std::clamp(padrow - 1, 0, GPUCA_NROWS - 1);
+  const int row0 = ((padrow == mTPCGeometry.EndIROC()) || (padrow == mTPCGeometry.EndOROC1()) || (padrow == mTPCGeometry.EndOROC2())) ? padrow : std::clamp(padrow - 1, 0, o2::tpc::constants::MAXGLOBALPADROW - 1);
 
   // check if previous pad row has enough pads
   const unsigned char pad0 = std::clamp(static_cast<int>(pad), 0, mTPCGeometry.NPads(row0) - 1);

Detectors/TPC/monitor/src/SimpleEventDisplayGUI.cxx

@@ -1227,7 +1227,7 @@ void SimpleEventDisplayGUI::showClusters(int roc, int row)
         }
         if (fillSingleTB && std::abs(cl.getTime() - timeBin) < 2) {
           const auto ly = gpuGeom.LinearPad2Y(sector, irow, cl.getPad() + 0.5);
-          mClustersRowPad->SetNextPoint(gpuGeom.Row2X(irow), (sector >= GPUCA_NSECTORS / 2) ? -ly : ly);
+          mClustersRowPad->SetNextPoint(gpuGeom.Row2X(irow), (sector >= gpuGeom.NSECTORS / 2) ? -ly : ly);
         }
       }
       // fmt::print("\n");

Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx

@@ -158,7 +158,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
 
     const auto& tinfo = pc.services().get<o2::framework::TimingInfo>();
     const auto firstIR = o2::InteractionRecord(0, tinfo.firstTForbit);
-    const float totalT = std::max(mFastTransform->getMaxDriftTime(0), mFastTransform->getMaxDriftTime(GPUCA_NSECTORS / 2));
+    const float totalT = std::max(mFastTransform->getMaxDriftTime(0), mFastTransform->getMaxDriftTime(GPUTPCGeometry::NSECTORS / 2));
 
     unsigned int offset = 0, lasti = 0;
     const unsigned int maxTime = (mParam->continuousMaxTimeBin + 1) * o2::tpc::ClusterNative::scaleTimePacked - 1;
@@ -205,23 +205,23 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
       }
     }
     offset = 0;
-    unsigned int offsets[GPUCA_NSECTORS][GPUCA_NROWS];
-    for (unsigned int i = 0; i < GPUCA_NSECTORS; i++) {
-      for (unsigned int j = 0; j < GPUCA_NROWS; j++) {
-        if (i * GPUCA_NROWS + j >= clusters.nSliceRows) {
+    unsigned int offsets[GPUTPCGeometry::NSECTORS][GPUTPCGeometry::NROWS];
+    for (unsigned int i = 0; i < GPUTPCGeometry::NSECTORS; i++) {
+      for (unsigned int j = 0; j < GPUTPCGeometry::NROWS; j++) {
+        if (i * GPUTPCGeometry::NROWS + j >= clusters.nSliceRows) {
           break;
         }
         offsets[i][j] = offset;
-        offset += (i * GPUCA_NROWS + j >= clusters.nSliceRows) ? 0 : clusters.nSliceRowClusters[i * GPUCA_NROWS + j];
+        offset += (i * GPUTPCGeometry::NROWS + j >= clusters.nSliceRows) ? 0 : clusters.nSliceRowClusters[i * GPUTPCGeometry::NROWS + j];
       }
     }
 
 #ifdef WITH_OPENMP
-#pragma omp parallel for num_threads(mNThreads) schedule(static, (GPUCA_NSECTORS + mNThreads - 1) / mNThreads) // Static round-robin scheduling with one chunk per thread to ensure correct order of the final vector
+#pragma omp parallel for num_threads(mNThreads) schedule(static, (GPUTPCGeometry::NSECTORS + mNThreads - 1) / mNThreads) // Static round-robin scheduling with one chunk per thread to ensure correct order of the final vector
 #endif
     for (unsigned int ii = 0; ii < clusters.nSliceRows; ii++) {
-      unsigned int i = ii / GPUCA_NROWS;
-      unsigned int j = ii % GPUCA_NROWS;
+      unsigned int i = ii / GPUTPCGeometry::NROWS;
+      unsigned int j = ii % GPUTPCGeometry::NROWS;
       o2::tpc::ClusterNative preCl;
 #ifdef WITH_OPENMP
       int myThread = omp_get_thread_num();
@@ -240,7 +240,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
         const bool reject = mCTFCoder.getIRFramesSelector().check(o2::dataformats::IRFrame(chkVal, chkVal + 1), chkExt, 0) < 0;
         if (reject) {
           rejectHits[k] = true;
-          clustersFiltered.nSliceRowClusters[i * GPUCA_NROWS + j]--;
+          clustersFiltered.nSliceRowClusters[i * GPUTPCGeometry::NROWS + j]--;
           static std::atomic_flag lock = ATOMIC_FLAG_INIT;
           while (lock.test_and_set(std::memory_order_acquire)) {
           }
@@ -253,7 +253,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
           preCl = cl;
         }
       };
-      unsigned int end = offsets[i][j] + clusters.nSliceRowClusters[i * GPUCA_NROWS + j];
+      unsigned int end = offsets[i][j] + clusters.nSliceRowClusters[i * GPUTPCGeometry::NROWS + j];
       o2::gpu::TPCClusterDecompressionCore::decompressHits(clusters, offsets[i][j], end, checker);
     }
     tmpBuffer[0].first.reserve(clustersFiltered.nUnattachedClusters);

GPU/GPUTracking/Base/GPUConstantMem.h

@@ -40,13 +40,13 @@ namespace o2::gpu
 {
 struct GPUConstantMem {
   GPUParam param;
-  GPUTPCTracker tpcTrackers[GPUCA_NSECTORS];
+  GPUTPCTracker tpcTrackers[GPUTPCGeometry::NSECTORS];
   GPUTPCCompression tpcCompressor;
   GPUTPCDecompression tpcDecompressor;
   GPUTPCGMMerger tpcMerger;
   GPUTRDTrackerGPU trdTrackerGPU;
   GPUTRDTracker trdTrackerO2;
-  GPUTPCClusterFinder tpcClusterer[GPUCA_NSECTORS];
+  GPUTPCClusterFinder tpcClusterer[GPUTPCGeometry::NSECTORS];
   GPUTrackingRefitProcessor trackingRefit;
   GPUTrackingInOutPointers ioPtrs;
   GPUCalibObjectsConst calibObjects;
@@ -55,7 +55,7 @@ struct GPUConstantMem {
   GPUKernelDebugOutput debugOutput;
 #endif
 #ifdef GPUCA_HAS_ONNX
-  GPUTPCNNClusterizer tpcNNClusterer[GPUCA_NSECTORS];
+  GPUTPCNNClusterizer tpcNNClusterer[GPUTPCGeometry::NSECTORS];
 #endif
   template <int32_t I>
   GPUd() auto& getTRDTracker();

GPU/GPUTracking/Base/GPUParam.cxx

@@ -91,16 +91,16 @@ void GPUParam::SetDefaults(float solenoidBz, bool assumeConstantBz)
   constexpr float plusZmax = 249.778;
   constexpr float minusZmin = -249.645;
   constexpr float minusZmax = -0.0799937;
-  for (int32_t i = 0; i < GPUCA_NSECTORS; i++) {
-    const bool zPlus = (i < GPUCA_NSECTORS / 2);
+  for (uint32_t i = 0; i < GPUTPCGeometry::NSECTORS; i++) {
+    const bool zPlus = (i < GPUTPCGeometry::NSECTORS / 2);
     SectorParam[i].ZMin = zPlus ? plusZmin : minusZmin;
     SectorParam[i].ZMax = zPlus ? plusZmax : minusZmax;
     int32_t tmp = i;
-    if (tmp >= GPUCA_NSECTORS / 2) {
-      tmp -= GPUCA_NSECTORS / 2;
+    if (tmp >= (int32_t)GPUTPCGeometry::NSECTORS / 2) {
+      tmp -= GPUTPCGeometry::NSECTORS / 2;
     }
-    if (tmp >= GPUCA_NSECTORS / 4) {
-      tmp -= GPUCA_NSECTORS / 2;
+    if (tmp >= (int32_t)GPUTPCGeometry::NSECTORS / 4) {
+      tmp -= GPUTPCGeometry::NSECTORS / 2;
     }
     SectorParam[i].Alpha = 0.174533f + dAlpha * tmp;
     SectorParam[i].CosAlpha = CAMath::Cos(SectorParam[i].Alpha);

GPU/GPUTracking/Base/GPUParam.h

@@ -20,6 +20,7 @@
 #include "GPUDef.h"
 #include "GPUSettings.h"
 #include "GPUTPCGMPolynomialField.h"
+#include "GPUTPCGeometry.h"
 
 #if !defined(GPUCA_GPUCODE)
 namespace o2::base
@@ -65,7 +66,7 @@ struct GPUParam_t {
   uint32_t occupancyTotal;                 // Total occupancy in the TPC (nCl / nHbf)
   uint32_t occupancyMapSize;               // Size of occupancy map
 
-  GPUParamSector SectorParam[GPUCA_NSECTORS];
+  GPUParamSector SectorParam[GPUTPCGeometry::NSECTORS];
 
  protected:
 #ifndef GPUCA_RUN2
@@ -87,13 +88,14 @@ struct GPUParam : public internal::GPUParam_t<GPUSettingsRec, GPUSettingsParam>
   void UpdateRun3ClusterErrors(const float* yErrorParam, const float* zErrorParam);
 #endif
 
-  GPUd() float Alpha(int32_t iSector) const
+  GPUd() constexpr uint32_t tpcMinHitsB5(float qPtB5) const { return CAMath::Abs(qPtB5) > 10 ? 10 : (CAMath::Abs(qPtB5) > 5 ? 15 : 29); } // Minimum hits should depend on Pt, low Pt tracks can have few hits. 29 Hits default, 15 for < 200 mev, 10 for < 100 mev
+  GPUd() constexpr float Alpha(int32_t iSector) const
   {
-    if (iSector >= GPUCA_NSECTORS / 2) {
-      iSector -= GPUCA_NSECTORS / 2;
+    if (iSector >= (int32_t)GPUTPCGeometry::NSECTORS / 2) {
+      iSector -= GPUTPCGeometry::NSECTORS / 2;
     }
-    if (iSector >= GPUCA_NSECTORS / 4) {
-      iSector -= GPUCA_NSECTORS / 2;
+    if (iSector >= (int32_t)GPUTPCGeometry::NSECTORS / 4) {
+      iSector -= GPUTPCGeometry::NSECTORS / 2;
     }
     return 0.174533f + dAlpha * iSector;
   }

GPU/GPUTracking/Base/GPUParam.inc

@@ -116,7 +116,7 @@ GPUdi() float GPUParam::GetSystematicClusterErrorC122(float x, float y, uint8_t
     return 0.f;
   }
   constexpr float dEdgeInv = 18.f / CAMath::Pi();
-  const float dy = (sector == (GPUCA_NSECTORS / 2 + 1) ? 0.5f : -0.5f) * (y / x) * dEdgeInv + 0.5f;
+  const float dy = (sector == (GPUTPCGeometry::NSECTORS / 2 + 1) ? 0.5f : -0.5f) * (y / x) * dEdgeInv + 0.5f;
   const float errC12 = rec.tpc.sysClusErrorC12Norm * occupancyTotal * dy;
   return errC12 * errC12;
 }

GPU/GPUTracking/Base/GPUProcessor.h

@@ -62,7 +62,7 @@ class GPUProcessor
     return *(T*)(mGPUProcessorType == PROCESSOR_TYPE_DEVICE ? mLinkedProcessor : this);
   }
 
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT>
   static constexpr inline size_t getAlignmentMod(size_t addr)
   {
     static_assert((alignment & (alignment - 1)) == 0, "Invalid alignment, not power of 2");
@@ -71,7 +71,7 @@ class GPUProcessor
     }
     return addr & (alignment - 1);
   }
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT>
   static constexpr inline size_t getAlignment(size_t addr)
   {
     size_t mod = getAlignmentMod<alignment>(addr);
@@ -80,7 +80,7 @@ class GPUProcessor
     }
     return (alignment - mod);
   }
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT>
   static constexpr inline size_t nextMultipleOf(size_t size)
   {
     return size + getAlignment<alignment>(size);
@@ -97,22 +97,22 @@ class GPUProcessor
       return (size + alignment - 1) & ~(alignment - 1);
     }
   }
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT>
   static inline void* alignPointer(void* ptr)
   {
     return (reinterpret_cast<void*>(nextMultipleOf<alignment>(reinterpret_cast<size_t>(ptr))));
   }
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT>
   static inline size_t getAlignmentMod(void* addr)
   {
     return (getAlignmentMod<alignment>(reinterpret_cast<size_t>(addr)));
   }
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT>
   static inline size_t getAlignment(void* addr)
   {
     return (getAlignment<alignment>(reinterpret_cast<size_t>(addr)));
   }
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT, class S>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT, class S>
   static inline S* getPointerWithAlignment(size_t& basePtr, size_t nEntries = 1)
   {
     if (basePtr == 0) {
@@ -125,7 +125,7 @@ class GPUProcessor
     return retVal;
   }
 
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT, class S>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT, class S>
   static inline S* getPointerWithAlignment(void*& basePtr, size_t nEntries = 1)
   {
     size_t tmp = (size_t)basePtr;
@@ -134,7 +134,7 @@ class GPUProcessor
     return retVal;
   }
 
-  template <size_t alignment = GPUCA_BUFFER_ALIGNMENT, class T, class S>
+  template <size_t alignment = constants::GPU_BUFFER_ALIGNMENT, class T, class S>
   static inline void computePointerWithAlignment(T*& basePtr, S*& objPtr, size_t nEntries = 1)
   {
     size_t tmp = (size_t)basePtr;
@@ -145,8 +145,8 @@ class GPUProcessor
   template <class T, class S>
   static inline void computePointerWithoutAlignment(T*& basePtr, S*& objPtr, size_t nEntries = 1)
   {
-    if ((size_t)basePtr < GPUCA_BUFFER_ALIGNMENT) {
-      basePtr = (T*)GPUCA_BUFFER_ALIGNMENT;
+    if ((size_t)basePtr < constants::GPU_BUFFER_ALIGNMENT) {
+      basePtr = (T*)constants::GPU_BUFFER_ALIGNMENT;
     }
     size_t tmp = (size_t)basePtr;
     objPtr = reinterpret_cast<S*>(getPointerWithAlignment<1, char>(tmp, nEntries * sizeof(S)));