GPU: Remove all AliRoot code

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

@@ -65,9 +65,7 @@ class PIDResponse
   float mMIP = 50.f;
   float mChargeFactor = 2.299999952316284f;
 
-#ifndef GPUCA_ALIROOT_LIB
   ClassDefNV(PIDResponse, 1);
-#endif
 };
 
 GPUd() void PIDResponse::setBetheBlochParams(const float betheBlochParams[5])

DataFormats/Detectors/TRD/include/DataFormatsTRD/Tracklet64.h

@@ -144,7 +144,7 @@ class Tracklet64
   GPUd() float getPadColFloat(bool applyShift) const { return getPositionFloat() + getMCMCol() * constants::NCOLMCM + 8.f + (applyShift ? 1.f : 0.f); }
 
   // pad column number inside pad row as int can be off by +-1 pad (same function name as for TRD digit)
-  GPUd() int getPadCol(bool applyShift = false) const { return GPUCA_NAMESPACE::gpu::CAMath::Float2IntRn(getPadColFloat(applyShift)); }
+  GPUd() int getPadCol(bool applyShift = false) const { return o2::gpu::CAMath::Float2IntRn(getPadColFloat(applyShift)); }
 
   // translate local position into global y (in cm) not taking into account calibrations (ExB, vDrift, t0)
   GPUd() float getUncalibratedY(bool applyShift = false) const { return (getPadColFloat(applyShift) - (constants::NCOLUMN / 2.f)) * getPadWidth(); }

Detectors/AOD/src/AODProducerWorkflowSpec.cxx

@@ -1470,7 +1470,7 @@ void AODProducerWorkflowDPL::countTPCClusters(const o2::globaltracking::RecoCont
       o2::tpc::TrackTPC::getClusterReference(tpcClusRefs, i, sectorIndex, rowIndex, clusterIndex, track.getClusterRef());
       unsigned int absoluteIndex = tpcClusAcc.clusterOffset[sectorIndex][rowIndex] + clusterIndex;
       clMap[rowIndex] = true;
-      if (tpcClusShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+      if (tpcClusShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
         if (!shMap[rowIndex]) {
           counters.shared++;
         }

Detectors/GlobalTracking/include/GlobalTracking/TrackMethods.h

@@ -47,7 +47,7 @@ class TrackMethods
       o2::tpc::TrackTPC::getClusterReference(tpcClusRefs, i, sectorIndex, rowIndex, clusterIndex, track.getClusterRef());
       unsigned int absoluteIndex = tpcClusAcc.clusterOffset[sectorIndex][rowIndex] + clusterIndex;
       clMap[rowIndex] = true;
-      if (tpcClusShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+      if (tpcClusShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
         if (!shMap[rowIndex]) {
           shared++;
         }

Detectors/GlobalTracking/src/MatchTPCITS.cxx

@@ -2880,7 +2880,7 @@ void MatchTPCITS::dumpTPCOrig(bool acc, int tpcIndex)
   for (int i = 0; i < tpcOrig.getNClusterReferences(); i++) {
     tpcOrig.getClusterReference(mTPCTrackClusIdx, i, clSect, clRow, clIdx);
     unsigned int absoluteIndex = mTPCClusterIdxStruct->clusterOffset[clSect][clRow] + clIdx;
-    if (mTPCRefitterShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+    if (mTPCRefitterShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
       if (!(prevRow == clRow && prevRawShared)) {
         nshared++;
       }

Detectors/GlobalTrackingWorkflow/study/src/TrackingStudy.cxx

@@ -266,7 +266,7 @@ void TrackingStudySpec::process(o2::globaltracking::RecoContainer& recoData)
           clRowP = clRow;
         }
         unsigned int absoluteIndex = tpcClusAcc.clusterOffset[clSect][clRow] + clIdx;
-        if (shMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+        if (shMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
           trExt.nClTPCShared++;
         }
       }

Detectors/ITSMFT/ITS/postprocessing/studies/include/ITSStudies/TrackCuts.h

@@ -76,7 +76,7 @@ class TrackCuts
         o2::tpc::TrackTPC::getClusterReference(tpcClusRefs, i, sectorIndex, rowIndex, clusterIndex, track.getClusterRef());
         unsigned int absoluteIndex = tpcClusAcc.clusterOffset[sectorIndex][rowIndex] + clusterIndex;
         clMap[rowIndex] = true;
-        if (tpcClusShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+        if (tpcClusShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
           if (!shMap[rowIndex]) {
             counters.shared++;
           }

Detectors/ITSMFT/ITS/postprocessing/studies/include/ITSStudies/TrackMethods.h

@@ -53,7 +53,7 @@ class TrackMethods
       o2::tpc::TrackTPC::getClusterReference(tpcClusRefs, i, sectorIndex, rowIndex, clusterIndex, track.getClusterRef());
       unsigned int absoluteIndex = tpcClusAcc.clusterOffset[sectorIndex][rowIndex] + clusterIndex;
       clMap[rowIndex] = true;
-      if (tpcClusShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+      if (tpcClusShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
         if (!shMap[rowIndex]) {
           shared++;
         }

Detectors/TPC/calibration/include/TPCCalibration/CorrectionMapsLoader.h

@@ -75,7 +75,7 @@ class CorrectionMapsLoader : public o2::gpu::CorrectionMapsHelper
 
   float mInstLumiCTPFactor = 1.0; // multiplicative factor for inst. lumi
   int mLumiCTPSource = 0;         // 0: main, 1: alternative CTP lumi source
-  std::unique_ptr<GPUCA_NAMESPACE::gpu::TPCFastTransform> mCorrMapMShape{nullptr};
+  std::unique_ptr<o2::gpu::TPCFastTransform> mCorrMapMShape{nullptr};
 #endif
 };
 

Detectors/TPC/calibration/src/CalculatedEdx.cxx

@@ -159,7 +159,7 @@ void CalculatedEdx::calculatedEdx(o2::tpc::TrackTPC& track, dEdxInfo& output, fl
 
     // check if the cluster is shared
     const unsigned int absoluteIndex = mClusterIndex->clusterOffset[sectorIndex][rowIndex] + clusterIndexNumb;
-    const bool isShared = mRefit ? (mTPCRefitterShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) : 0;
+    const bool isShared = mRefit ? (mTPCRefitterShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) : 0;
 
     // get region, pad, stack and stack ID
     const int region = Mapper::REGION[rowIndex];

Detectors/TPC/calibration/src/DigitAdd.cxx

@@ -23,13 +23,13 @@ int DigitAdd::sector() const
 
 float DigitAdd::lx() const
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
   return gpuGeom.Row2X(mRow);
 }
 
 float DigitAdd::ly() const
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
   return gpuGeom.LinearPad2Y(sector(), mRow, getPad());
 }
 
@@ -49,6 +49,6 @@ float DigitAdd::gy() const
 
 float DigitAdd::cpad() const
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
   return getPad() - gpuGeom.NPads(mRow) / 2.f;
 }

Detectors/TPC/calibration/src/TrackDump.cxx

@@ -47,7 +47,7 @@ void TrackDump::filter(const gsl::span<const TrackTPC> tracks, ClusterNativeAcce
 
   ClExcludes excludes;
 
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
 
   for (const auto& track : tracks) {
     const int nCl = track.getNClusterReferences();
@@ -141,7 +141,7 @@ void TrackDump::finalize()
 
 void TrackDump::fillClNativeAdd(ClusterNativeAccess const& clusterIndex, std::vector<ClusterNativeAdd>& clInfos, ClExcludes* excludes)
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
 
   for (int sector = 0; sector < MAXSECTOR; ++sector) {
     for (int padrow = 0; padrow < MAXGLOBALPADROW; ++padrow) {
@@ -164,19 +164,19 @@ void TrackDump::fillClNativeAdd(ClusterNativeAccess const& clusterIndex, std::ve
 
 float TrackDump::ClusterNativeAdd::cpad() const
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
   return getPad() - gpuGeom.NPads(padrow) / 2.f;
 }
 
 float TrackDump::ClusterNativeAdd::lx() const
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
   return gpuGeom.Row2X(padrow);
 }
 
 float TrackDump::ClusterNativeAdd::ly() const
 {
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
   return gpuGeom.LinearPad2Y(sector, padrow, getPad());
 }
 

Detectors/TPC/monitor/src/SimpleEventDisplayGUI.cxx

@@ -1195,7 +1195,7 @@ void SimpleEventDisplayGUI::showClusters(int roc, int row)
     selFlags += mCheckClFlags[iFlag]->IsDown() << (iFlag - 1);
   }
   const bool fillSingleTB = mCheckSingleTB->IsDown();
-  const GPUCA_NAMESPACE::gpu::GPUTPCGeometry gpuGeom;
+  const o2::gpu::GPUTPCGeometry gpuGeom;
 
   const int rowMin = fillSingleTB ? 0 : row;
   const int rowMax = fillSingleTB ? constants::MAXGLOBALPADROW : row + 1;

Detectors/TPC/workflow/src/TPCRefitter.cxx

@@ -509,7 +509,7 @@ bool TPCRefitterSpec::processTPCTrack(o2::tpc::TrackTPC tr, o2::MCCompLabel lbl,
       unsigned int absoluteIndex = mTPCClusterIdxStruct->clusterOffset[sector][row] + clusterIndex;
       cl = &mTPCClusterIdxStruct->clusters[sector][row][clusterIndex];
       uint8_t clflags = cl->getFlags();
-      if (mTPCRefitterShMap[absoluteIndex] & GPUCA_NAMESPACE::gpu::GPUTPCGMMergedTrackHit::flagShared) {
+      if (mTPCRefitterShMap[absoluteIndex] & o2::gpu::GPUTPCGMMergedTrackHit::flagShared) {
         clflags |= 0x10;
       }
       clData.clSector.emplace_back(sector);

GPU/CMakeLists.txt

@@ -9,14 +9,9 @@
 # granted to it by virtue of its status as an Intergovernmental Organization
 # or submit itself to any jurisdiction.
 
-# Subdirectories will be compiled with O2 / AliRoot / Standalone To simplify the
+# Subdirectories will be compiled with O2 / Standalone To simplify the
 # CMake, variables are defined for Sources / Headers first. Then, the actual
 # CMake build scripts use these variables.
-#
-# SRCS: Common Sources for all builds HDRS_CINT: Headers for ROOT dictionary
-# (always) HDRS_CINT_ALIROOT: Headers for ROOT dictionary (only in AliRoot)
-# HDRS_CINT_O2: Headers for ROOT dictionary (only for O2) HDRS_INSTALL: Headers
-# for installation only
 
 if(NOT DEFINED GPUCA_NO_FAST_MATH)
   set(GPUCA_NO_FAST_MATH 0)

GPU/Common/CMakeLists.txt

@@ -66,26 +66,3 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2")
   endif()
   install(FILES ${HDRS_INSTALL} DESTINATION include/GPU)
 endif()
-
-if(ALIGPU_BUILD_TYPE STREQUAL "ALIROOT")
-  add_definitions(-DGPUCA_ALIROOT_LIB)
-
-  set(SRCS ${SRCS} ../GPUTracking/utils/EmptyFile.cxx)
-
-  # Add a library to the project using the specified source files
-  add_library_tested(Ali${MODULE} SHARED ${SRCS})
-
-  # Additional compilation flags
-  set_target_properties(Ali${MODULE} PROPERTIES COMPILE_FLAGS "")
-
-  # System dependent: Modify the way the library is build
-  if(${CMAKE_SYSTEM} MATCHES Darwin)
-    set_target_properties(Ali${MODULE}
-                          PROPERTIES LINK_FLAGS "-undefined dynamic_lookup")
-  endif(${CMAKE_SYSTEM} MATCHES Darwin)
-
-  # Installation
-  install(TARGETS Ali${MODULE} ARCHIVE DESTINATION lib LIBRARY DESTINATION lib)
-
-  install(FILES ${HDRS_INSTALL} DESTINATION include)
-endif()

GPU/Common/GPUCommonAlgorithm.h

@@ -24,7 +24,7 @@
 
 // ----------------------------- SORTING -----------------------------
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -72,9 +72,9 @@ class GPUCommonAlgorithm
   GPUd() static void IterSwap(I a, I b) noexcept;
 };
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -218,15 +218,15 @@ GPUdi() void GPUCommonAlgorithm::QuickSort(I f, I l) noexcept
 typedef GPUCommonAlgorithm CAAlgo;
 
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
 #if (((defined(__CUDACC__) && !defined(__clang__)) || defined(__HIPCC__))) && !defined(GPUCA_GPUCODE_GENRTC) && !defined(GPUCA_GPUCODE_HOSTONLY)
 
 #include "GPUCommonAlgorithmThrust.h"
 
 #else
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -248,12 +248,12 @@ GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end, const S& co
 }
 
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
 #endif // THRUST
 // sort and sortInBlock below are not taken from Thrust, since our implementations are faster
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -329,7 +329,7 @@ GPUdi() void GPUCommonAlgorithm::swap(T& a, T& b)
 #endif
 
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
 // ----------------------------- WORK GROUP FUNCTIONS -----------------------------
 

GPU/Common/GPUCommonAlgorithmThrust.h

@@ -30,7 +30,7 @@
 #define GPUCA_THRUST_NAMESPACE thrust::hip
 #endif
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -88,6 +88,6 @@ GPUdi() void GPUCommonAlgorithm::sortDeviceDynamic(T* begin, T* end, const S& co
 }
 
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
 #endif

GPU/Common/GPUCommonConstants.h

@@ -17,7 +17,7 @@
 
 #include "GPUCommonDef.h"
 
-namespace GPUCA_NAMESPACE::gpu::gpu_common_constants
+namespace o2::gpu::gpu_common_constants
 {
 static constexpr const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this now that we use only OpenCL CPP
 }

GPU/Common/GPUCommonDef.h

@@ -40,22 +40,16 @@
   #endif
 #endif
 
-// Set AliRoot / O2 namespace
-#if defined(GPUCA_STANDALONE) || (defined(GPUCA_O2_LIB) && !defined(GPUCA_O2_INTERFACE)) || defined(GPUCA_ALIROOT_LIB) || defined (GPUCA_GPUCODE)
+#if defined(GPUCA_STANDALONE) || (defined(GPUCA_O2_LIB) && !defined(GPUCA_O2_INTERFACE)) || defined (GPUCA_GPUCODE)
   #define GPUCA_ALIGPUCODE
 #endif
-#ifdef GPUCA_ALIROOT_LIB
-  #define GPUCA_NAMESPACE AliGPU
-#else
-  #define GPUCA_NAMESPACE o2
-#endif
 
 #if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY))
   #define GPUCA_NO_CONSTANT_MEMORY
 #elif defined(__CUDACC__) || defined(__HIPCC__)
   #define GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM
 #endif
-#if !defined(GPUCA_HAVE_O2HEADERS) && (defined(GPUCA_O2_LIB) || (!defined(GPUCA_ALIROOT_LIB) && !defined(GPUCA_STANDALONE)))
+#if !defined(GPUCA_HAVE_O2HEADERS) && (defined(GPUCA_O2_LIB) || !defined(GPUCA_STANDALONE))
   #define GPUCA_HAVE_O2HEADERS
 #endif
 

GPU/Common/GPUCommonDefAPI.h

@@ -104,12 +104,12 @@
     #define GPUbarrier() barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE)
     #define GPUbarrierWarp()
     #if defined(__OPENCL__) && defined(GPUCA_OPENCL_CLANG_C11_ATOMICS)
-      namespace GPUCA_NAMESPACE { namespace gpu {
+      namespace o2 { namespace gpu {
       template <class T> struct oclAtomic;
       template <> struct oclAtomic<uint32_t> {typedef atomic_uint t;};
       static_assert(sizeof(oclAtomic<uint32_t>::t) == sizeof(uint32_t), "Invalid size of atomic type");
       }}
-      #define GPUAtomic(type) GPUCA_NAMESPACE::gpu::oclAtomic<type>::t
+      #define GPUAtomic(type) o2::gpu::oclAtomic<type>::t
     #else
       #define GPUAtomic(type) volatile type
     #endif

GPU/Common/GPUCommonLogger.h

@@ -45,7 +45,7 @@ struct DummyLogger {
 #define LOGP(...)
 // #define LOGP(...) static_assert(false, "LOGP(...) unsupported in GPU code");
 
-#elif defined(GPUCA_STANDALONE) || defined(GPUCA_ALIROOT_LIB)
+#elif defined(GPUCA_STANDALONE)
 #include <iostream>
 #include <cstdio>
 #define LOG(type) std::cout

GPU/Common/GPUCommonMath.h

@@ -31,7 +31,7 @@
 #include <cstdint>
 #endif
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -552,6 +552,6 @@ GPUdii() void GPUCommonMath::AtomicMinInternal(GPUglobalref() GPUgeneric() GPUAt
 #undef CHOICE
 
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
 #endif // GPUCOMMONMATH_H

GPU/Common/GPUCommonTransform3D.h

@@ -17,7 +17,7 @@
 
 #include "GPUCommonDef.h"
 
-namespace GPUCA_NAMESPACE
+namespace o2
 {
 namespace gpu
 {
@@ -80,6 +80,6 @@ class Transform3D
                                 kDZ = 11 };
 };
 } // namespace gpu
-} // namespace GPUCA_NAMESPACE
+} // namespace o2
 
 #endif