GPU: Renaming and removal of obsolete macros / defines

Author: davidrohr

DataFormats/Detectors/TPC/include/DataFormatsTPC/Constants.h

@@ -28,7 +28,7 @@ namespace constants
 constexpr int MAXSECTOR = 36;
 
 // the number of global pad rows
-#if defined(GPUCA_STANDALONE) && !defined(GPUCA_O2_LIB) && !defined(GPUCA_TPC_GEOMETRY_O2)
+#if defined(GPUCA_STANDALONE) && defined(GPUCA_RUN2)
 constexpr int MAXGLOBALPADROW = 159; // Number of pad rows in Run 2, used for GPU TPC tests with Run 2 data
 #else
 constexpr int MAXGLOBALPADROW = 152; // Correct number of pad rows in Run 3

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_ROW_COUNT - 1);
+  const int row1 = ((padrow == mTPCGeometry.EndIROC() - 1) || (padrow == mTPCGeometry.EndOROC1() - 1) || (padrow == mTPCGeometry.EndOROC2() - 1)) ? padrow : std::clamp(padrow + 1, 0, GPUCA_NROWS - 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_ROW_COUNT - 1);
+  const int row0 = ((padrow == mTPCGeometry.EndIROC()) || (padrow == mTPCGeometry.EndOROC1()) || (padrow == mTPCGeometry.EndOROC2())) ? padrow : std::clamp(padrow - 1, 0, GPUCA_NROWS - 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/workflow/src/EntropyEncoderSpec.cxx

@@ -205,23 +205,23 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
       }
     }
     offset = 0;
-    unsigned int offsets[GPUCA_NSECTORS][GPUCA_ROW_COUNT];
+    unsigned int offsets[GPUCA_NSECTORS][GPUCA_NROWS];
     for (unsigned int i = 0; i < GPUCA_NSECTORS; i++) {
-      for (unsigned int j = 0; j < GPUCA_ROW_COUNT; j++) {
-        if (i * GPUCA_ROW_COUNT + j >= clusters.nSliceRows) {
+      for (unsigned int j = 0; j < GPUCA_NROWS; j++) {
+        if (i * GPUCA_NROWS + j >= clusters.nSliceRows) {
           break;
         }
         offsets[i][j] = offset;
-        offset += (i * GPUCA_ROW_COUNT + j >= clusters.nSliceRows) ? 0 : clusters.nSliceRowClusters[i * GPUCA_ROW_COUNT + j];
+        offset += (i * GPUCA_NROWS + j >= clusters.nSliceRows) ? 0 : clusters.nSliceRowClusters[i * GPUCA_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
 #endif
     for (unsigned int ii = 0; ii < clusters.nSliceRows; ii++) {
-      unsigned int i = ii / GPUCA_ROW_COUNT;
-      unsigned int j = ii % GPUCA_ROW_COUNT;
+      unsigned int i = ii / GPUCA_NROWS;
+      unsigned int j = ii % GPUCA_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_ROW_COUNT + j]--;
+          clustersFiltered.nSliceRowClusters[i * GPUCA_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_ROW_COUNT + j];
+      unsigned int end = offsets[i][j] + clusters.nSliceRowClusters[i * GPUCA_NROWS + j];
       o2::gpu::TPCClusterDecompressionCore::decompressHits(clusters, offsets[i][j], end, checker);
     }
     tmpBuffer[0].first.reserve(clustersFiltered.nUnattachedClusters);

GPU/GPUTracking/Base/GPUParam.cxx

@@ -38,7 +38,7 @@ void GPUParam::SetDefaults(float solenoidBz, bool assumeConstantBz)
   occupancyMapSize = 0;
   occupancyTotal = 0;
 
-#ifdef GPUCA_TPC_GEOMETRY_O2
+#ifndef GPUCA_RUN2
   const float kErrorsY[4] = {0.06, 0.24, 0.12, 0.1};
   const float kErrorsZ[4] = {0.06, 0.24, 0.15, 0.1};
 
@@ -161,7 +161,7 @@ void GPUParam::SetDefaults(const GPUSettingsGRP* g, const GPUSettingsRec* r, con
 
 void GPUParam::UpdateRun3ClusterErrors(const float* yErrorParam, const float* zErrorParam)
 {
-#ifdef GPUCA_TPC_GEOMETRY_O2
+#ifndef GPUCA_RUN2
   for (int32_t yz = 0; yz < 2; yz++) {
     const float* param = yz ? zErrorParam : yErrorParam;
     for (int32_t rowType = 0; rowType < 4; rowType++) {

GPU/GPUTracking/Base/GPUParam.h

@@ -68,7 +68,7 @@ struct GPUParam_t {
   GPUParamSector SectorParam[GPUCA_NSECTORS];
 
  protected:
-#ifdef GPUCA_TPC_GEOMETRY_O2
+#ifndef GPUCA_RUN2
   float ParamErrors[2][4][4]; // cluster error parameterization used during seeding and fit
 #else
   float ParamErrorsSeeding0[2][3][4]; // cluster error parameterization used during seeding

GPU/GPUTracking/Base/GPUParam.inc

@@ -39,7 +39,7 @@ GPUdi() void GPUParam::Global2Sector(int32_t iSector, float X, float Y, float Z,
   *z = Z;
 }
 
-#ifdef GPUCA_TPC_GEOMETRY_O2
+#ifndef GPUCA_RUN2
 
 GPUdi() void GPUParam::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
 {
@@ -121,7 +121,7 @@ GPUdi() float GPUParam::GetSystematicClusterErrorC122(float x, float y, uint8_t
   return errC12 * errC12;
 }
 
-#else // GPUCA_TPC_GEOMETRY_O2
+#else // !GPUCA_RUN2
 
 GPUdi() float GPUParam::GetClusterErrorSeeding(int32_t yz, int32_t type, float zDiff, float angle2, float scaledMult) const
 {
@@ -169,7 +169,7 @@ GPUdi() float GPUParam::GetSystematicClusterErrorC122(float trackX, float trackY
   return 0;
 }
 
-#endif // !GPUCA_TPC_GEOMETRY_O2
+#endif // GPUCA_RUN2
 
 GPUdi() void GPUParam::GetClusterErrors2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float avgInvCharge, float invCharge, float& ErrY2, float& ErrZ2) const
 {

GPU/GPUTracking/Base/GPUReconstruction.cxx

@@ -42,7 +42,7 @@
 #include "utils/strtag.h"
 #include "utils/stdspinlock.h"
 
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
 #include "GPUO2InterfaceConfiguration.h"
 #endif
 
@@ -1249,15 +1249,15 @@ int32_t GPUReconstruction::ReadSettings(const char* dir)
 
 void GPUReconstruction::SetSettings(float solenoidBzNominalGPU, const GPURecoStepConfiguration* workflow)
 {
-#ifdef GPUCA_O2_LIB
+#ifdef GPUCA_STANDALONE
+  GPUSettingsGRP grp;
+  grp.solenoidBzNominalGPU = solenoidBzNominalGPU;
+  SetSettings(&grp, nullptr, nullptr, workflow);
+#else
   GPUO2InterfaceConfiguration config;
   config.ReadConfigurableParam(config);
   config.configGRP.solenoidBzNominalGPU = solenoidBzNominalGPU;
   SetSettings(&config.configGRP, &config.configReconstruction, &config.configProcessing, workflow);
-#else
-  GPUSettingsGRP grp;
-  grp.solenoidBzNominalGPU = solenoidBzNominalGPU;
-  SetSettings(&grp, nullptr, nullptr, workflow);
 #endif
 }
 

GPU/GPUTracking/Base/GPUReconstruction.h

@@ -92,10 +92,10 @@ class GPUReconstruction
   using InOutTypeField = gpudatatypes::InOutTypeField;
 
   static constexpr const char* const GEOMETRY_TYPE_NAMES[] = {"INVALID", "ALIROOT", "O2"};
-#ifdef GPUCA_TPC_GEOMETRY_O2
-  static constexpr GeometryType geometryType = GeometryType::O2;
-#else
+#ifdef GPUCA_RUN2
   static constexpr GeometryType geometryType = GeometryType::ALIROOT;
+#else
+  static constexpr GeometryType geometryType = GeometryType::O2;
 #endif
 
   static DeviceType GetDeviceType(const char* type);

GPU/GPUTracking/Base/GPUReconstructionConvert.cxx

@@ -12,7 +12,7 @@
 /// \file GPUReconstructionConvert.cxx
 /// \author David Rohr
 
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
 #include "DetectorsRaw/RawFileWriter.h"
 #include "TPCBase/Sector.h"
 #include "DataFormatsTPC/Digit.h"
@@ -54,13 +54,13 @@ void GPUReconstructionConvert::ConvertNativeToClusterData(o2::tpc::ClusterNative
   uint32_t offset = 0;
   for (uint32_t i = 0; i < NSECTORS; i++) {
     uint32_t nClSector = 0;
-    for (int32_t j = 0; j < GPUCA_ROW_COUNT; j++) {
+    for (int32_t j = 0; j < GPUCA_NROWS; j++) {
       nClSector += native->nClusters[i][j];
     }
     nClusters[i] = nClSector;
     clusters[i].reset(new GPUTPCClusterData[nClSector]);
     nClSector = 0;
-    for (int32_t j = 0; j < GPUCA_ROW_COUNT; j++) {
+    for (int32_t j = 0; j < GPUCA_NROWS; j++) {
       for (uint32_t k = 0; k < native->nClusters[i][j]; k++) {
         const auto& clin = native->clusters[i][j][k];
         float x = 0, y = 0, z = 0;
@@ -98,7 +98,7 @@ void GPUReconstructionConvert::ConvertRun2RawToNative(o2::tpc::ClusterNativeAcce
   native.clustersLinear = nativeBuffer.get();
   native.setOffsetPtrs();
   for (uint32_t i = 0; i < NSECTORS; i++) {
-    for (uint32_t j = 0; j < GPUCA_ROW_COUNT; j++) {
+    for (uint32_t j = 0; j < GPUCA_NROWS; j++) {
       native.nClusters[i][j] = 0;
     }
     for (uint32_t j = 0; j < nRawClusters[i]; j++) {
@@ -119,7 +119,7 @@ int32_t GPUReconstructionConvert::GetMaxTimeBin(const ClusterNativeAccess& nativ
 {
   float retVal = 0;
   for (uint32_t i = 0; i < NSECTORS; i++) {
-    for (uint32_t j = 0; j < GPUCA_ROW_COUNT; j++) {
+    for (uint32_t j = 0; j < GPUCA_NROWS; j++) {
       for (uint32_t k = 0; k < native.nClusters[i][j]; k++) {
         if (native.clusters[i][j][k].getTime() > retVal) {
           retVal = native.clusters[i][j][k].getTime();
@@ -174,7 +174,7 @@ int32_t GPUReconstructionConvert::GetMaxTimeBin(const GPUTrackingInOutZS& zspage
 
 // ------------------------------------------------- TPC ZS -------------------------------------------------
 
-#ifdef GPUCA_TPC_GEOMETRY_O2
+#ifndef GPUCA_RUN2
 namespace o2::gpu
 {
 namespace // anonymous
@@ -193,7 +193,7 @@ struct zsEncoder {
   const o2::InteractionRecord* ir = nullptr;
   const GPUParam* param = nullptr;
   bool padding = false;
-  int32_t lastEndpoint = -2, lastTime = -1, lastRow = GPUCA_ROW_COUNT;
+  int32_t lastEndpoint = -2, lastTime = -1, lastRow = GPUCA_NROWS;
   int32_t endpoint = 0, outputEndpoint = 0;
   int64_t hbf = -1, nexthbf = 0;
   zsPage* page = nullptr;
@@ -244,11 +244,11 @@ inline void zsEncoder::ZSstreamOut(uint16_t* bufIn, uint32_t& lenIn, uint8_t* bu
 
 static inline auto ZSEncoderGetDigits(const GPUTrackingInOutDigits& in, int32_t i) { return in.tpcDigits[i]; }
 static inline auto ZSEncoderGetNDigits(const GPUTrackingInOutDigits& in, int32_t i) { return in.nTPCDigits[i]; }
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
 using DigitArray = std::array<gsl::span<const o2::tpc::Digit>, o2::tpc::Sector::MAXSECTOR>;
 static inline auto ZSEncoderGetDigits(const DigitArray& in, int32_t i) { return in[i].data(); }
 static inline auto ZSEncoderGetNDigits(const DigitArray& in, int32_t i) { return in[i].size(); }
-#endif // GPUCA_O2_LIB
+#endif
 
 // ------------------------------------------------- TPC ZS Original Row-based ZS -------------------------------------------------
 
@@ -364,7 +364,7 @@ uint32_t zsEncoderRow::encodeSequence(std::vector<o2::tpc::Digit>& tmpBuffer, ui
     curTBHdr = reinterpret_cast<TPCZSTBHDR*>(pagePtr);
     curTBHdr->rowMask |= (endpoint & 1) << 15;
     nRowsInTB = 0;
-    lastRow = GPUCA_ROW_COUNT;
+    lastRow = GPUCA_NROWS;
   }
   if (tmpBuffer[k].getRow() != lastRow) {
     curTBHdr->rowMask |= 1 << (tmpBuffer[k].getRow() - endpointStart);
@@ -460,7 +460,7 @@ void zsEncoderRow::decodePage(std::vector<o2::tpc::Digit>& outputBuffer, const z
 
 // ------------------------------------------------- TPC ZS Link Based ZS -------------------------------------------------
 
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
 struct zsEncoderLinkBased : public zsEncoder {
   TPCZSHDRV2* hdr = nullptr;
   TPCZSHDRV2 hdrBuffer;
@@ -1045,7 +1045,7 @@ void zsEncoderDenseLinkBased::amendPageErrorMessage(std::ostringstream& oss, con
   oss << "Meta header of page: " << dumpBuffer << "\n";
 }
 
-#endif // GPUCA_O2_LIB
+#endif // !GPUCA_STANDALONE
 
 // ------------------------------------------------- TPC ZS Main Encoder -------------------------------------------------
 
@@ -1095,7 +1095,7 @@ inline uint32_t zsEncoderRun<T>::run(std::vector<zsPage>* buffer, std::vector<o2
 {
   uint32_t totalPages = 0;
   zsPage singleBuffer;
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
   int32_t maxhbf = 0;
   int32_t minhbf = o2::constants::lhc::LHCMaxBunches;
 #endif
@@ -1178,7 +1178,7 @@ inline uint32_t zsEncoderRun<T>::run(std::vector<zsPage>* buffer, std::vector<o2
         }
         size_t size = !std::is_same_v<T, struct zsEncoderDenseLinkBased> && (padding || lastEndpoint == -1 || hbf == nexthbf) ? TPCZSHDR::TPC_ZS_PAGE_SIZE : (pagePtr - (uint8_t*)page);
         size = CAMath::nextMultipleOf<o2::raw::RDHUtils::GBTWord128>(size);
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
         if (raw) {
           raw->addData(rawfeeid, rawcru, 0, rawendpoint, *ir + hbf * o2::constants::lhc::LHCMaxBunches, gsl::span<char>((char*)page + sizeof(o2::header::RAWDataHeader), (char*)page + size), true, 0, 2);
           maxhbf = std::max<int32_t>(maxhbf, hbf);
@@ -1257,7 +1257,7 @@ inline uint32_t zsEncoderRun<T>::run(std::vector<zsPage>* buffer, std::vector<o2
     k += nEncoded;
   }
   if (raw) {
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
     if (iSector == 0) {
       for (int32_t i = minhbf; i <= maxhbf; i++) {
         raw->addData(46208, 360, rdh_utils::SACLinkID, 0, *ir + i * o2::constants::lhc::LHCMaxBunches, gsl::span<char>((char*)&singleBuffer, (char*)&singleBuffer), true, 0, 4);
@@ -1310,7 +1310,7 @@ size_t zsEncoderRun<T>::compare(std::vector<zsPage>* buffer, std::vector<o2::tpc
 
 } // anonymous namespace
 } // namespace o2::gpu
-#endif // GPUCA_TPC_GEOMETRY_O2
+#endif // !GPUCA_RUN2
 
 template <class S>
 void GPUReconstructionConvert::RunZSEncoder(const S& in, std::unique_ptr<uint64_t[]>* outBuffer, uint32_t* outSizes, o2::raw::RawFileWriter* raw, const o2::InteractionRecord* ir, const GPUParam& param, int32_t version, bool verify, float threshold, bool padding, std::function<void(std::vector<o2::tpc::Digit>&)> digitsFilter)
@@ -1320,7 +1320,7 @@ void GPUReconstructionConvert::RunZSEncoder(const S& in, std::unique_ptr<uint64_
   if (((outBuffer == nullptr) ^ (outSizes == nullptr)) || ((raw != nullptr) && (ir == nullptr)) || !((outBuffer == nullptr) ^ (raw == nullptr)) || (raw && verify)) {
     throw std::runtime_error("Invalid parameters");
   }
-#ifdef GPUCA_TPC_GEOMETRY_O2
+#ifndef GPUCA_RUN2
   std::vector<zsPage> buffer[NSECTORS][GPUTrackingInOutZS::NENDPOINTS];
   struct tmpReductionResult {
     uint32_t totalPages = 0;
@@ -1365,7 +1365,7 @@ void GPUReconstructionConvert::RunZSEncoder(const S& in, std::unique_ptr<uint64_
         zsEncoderRun<zsEncoderRow> enc{{{.iSector = i, .raw = raw, .ir = ir, .param = &param, .padding = padding}}};
         runZS(enc);
       } else if (version >= ZSVersion::ZSVersionLinkBasedWithMeta && version <= ZSVersion::ZSVersionDenseLinkBasedV2) {
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
         if (version == ZSVersion::ZSVersionLinkBasedWithMeta) {
           zsEncoderRun<zsEncoderImprovedLinkBased> enc{{{{.iSector = i, .raw = raw, .ir = ir, .param = &param, .padding = padding}}}};
           runZS(enc);
@@ -1410,7 +1410,7 @@ void GPUReconstructionConvert::RunZSEncoder(const S& in, std::unique_ptr<uint64_
 }
 
 template void GPUReconstructionConvert::RunZSEncoder<GPUTrackingInOutDigits>(const GPUTrackingInOutDigits&, std::unique_ptr<uint64_t[]>*, uint32_t*, o2::raw::RawFileWriter*, const o2::InteractionRecord*, const GPUParam&, int32_t, bool, float, bool, std::function<void(std::vector<o2::tpc::Digit>&)> digitsFilter);
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
 template void GPUReconstructionConvert::RunZSEncoder<DigitArray>(const DigitArray&, std::unique_ptr<uint64_t[]>*, uint32_t*, o2::raw::RawFileWriter*, const o2::InteractionRecord*, const GPUParam&, int32_t, bool, float, bool, std::function<void(std::vector<o2::tpc::Digit>&)> digitsFilter);
 #endif
 
@@ -1454,7 +1454,7 @@ void GPUReconstructionConvert::RunZSFilter(std::unique_ptr<o2::tpc::Digit[]>* bu
   }
 }
 
-#ifdef GPUCA_O2_LIB
+#ifndef GPUCA_STANDALONE
 namespace o2::gpu::internal
 {
 template <class T>

GPU/GPUTracking/CMakeLists.txt

@@ -355,7 +355,7 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2")
                                        O2::TPCFastTransformation
                  PRIVATE_LINK_LIBRARIES O2::DataFormatsTPC
                  SOURCES ${SRCS_DATATYPES})
-  target_compile_definitions(${targetName} PRIVATE GPUCA_O2_LIB GPUCA_TPC_GEOMETRY_O2)
+  target_compile_definitions(${targetName} PRIVATE GPUCA_O2_LIB)
 
   o2_target_root_dictionary(GPUDataTypes
                             HEADERS ${HDRS_CINT_DATATYPES} ${HDRS_CINT_O2_ADDITIONAL}
@@ -384,7 +384,7 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2")
     ${targetName}
     PRIVATE $<TARGET_PROPERTY:O2::Framework,INTERFACE_INCLUDE_DIRECTORIES>)
 
-  target_compile_definitions(${targetName} PRIVATE GPUCA_O2_LIB GPUCA_TPC_GEOMETRY_O2)
+  target_compile_definitions(${targetName} PRIVATE GPUCA_O2_LIB)
 
   o2_target_root_dictionary(${MODULE}
                             HEADERS ${HDRS_CINT_O2} ${HDRS_CINT_O2_ADDITIONAL}