ALICE3: factor GPU tracking into dynamically loaded CUDA/HIP backend

Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt

@@ -13,24 +13,9 @@ if(Acts_FOUND)
   set(actsTarget Acts::Core)
 endif()
 
-set(alice3GlobalRecoGpuSources "")
-set(alice3GlobalRecoGpuTargets "")
-set(alice3GlobalRecoGpuPrivateTargets "")
-if(CUDA_ENABLED)
-  find_package(CUDAToolkit REQUIRED)
-  list(APPEND alice3GlobalRecoGpuSources src/TimeFrameGPU.cxx src/GPUExternalAllocator.cxx)
-  list(APPEND alice3GlobalRecoGpuTargets O2::ITStrackingCUDA)
-  list(APPEND alice3GlobalRecoGpuPrivateTargets CUDA::cudart)
-elseif(HIP_ENABLED)
-  list(APPEND alice3GlobalRecoGpuSources src/TimeFrameGPU.cxx src/GPUExternalAllocator.cxx)
-  list(APPEND alice3GlobalRecoGpuTargets O2::ITStrackingHIP)
-  list(APPEND alice3GlobalRecoGpuPrivateTargets hip::host)
-endif()
-
 o2_add_library(ALICE3GlobalReconstruction
                TARGETVARNAME targetName
                SOURCES src/TimeFrame.cxx
-                       ${alice3GlobalRecoGpuSources}
                        $<$<BOOL:${Acts_FOUND}>:src/TrackerACTS.cxx>
                PUBLIC_LINK_LIBRARIES
                        O2::ITStracking
@@ -48,26 +33,10 @@ o2_add_library(ALICE3GlobalReconstruction
                        O2::TRKReconstruction
                        O2::TRKSimulation
                        nlohmann_json::nlohmann_json
-                       ${alice3GlobalRecoGpuTargets}
                        ${actsTarget}
                PRIVATE_LINK_LIBRARIES
                        O2::Steer
-                       TBB::tbb
-                       ${alice3GlobalRecoGpuPrivateTargets})
-
-if(alice3GlobalRecoGpuTargets)
-  target_compile_definitions(${targetName} PUBLIC TRK_HAS_GPU_TRACKING)
-endif()
-
-if(CUDA_ENABLED)
-  target_include_directories(${targetName} PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
-endif()
-
-if(CUDA_ENABLED)
-  target_compile_definitions(${targetName} PUBLIC TRK_HAS_CUDA_TRACKING)
-elseif(HIP_ENABLED)
-  target_compile_definitions(${targetName} PUBLIC TRK_HAS_HIP_TRACKING)
-endif()
+                       TBB::tbb)
 
 if(Acts_FOUND)
   target_compile_definitions(${targetName} PUBLIC O2_WITH_ACTS)

Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx → Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cu

@@ -9,11 +9,9 @@
 // granted to it by virtue of its status as an Intergovernmental Organization
 // or submit itself to any jurisdiction.
 
-#if defined(TRK_HAS_CUDA_TRACKING)
+#define GPUCA_GPUCODE_HOSTONLY
+
 #include <cuda_runtime.h>
-#elif defined(TRK_HAS_HIP_TRACKING)
-#include <hip/hip_runtime.h>
-#endif
 
 #include "ALICE3GlobalReconstruction/GPUExternalAllocator.h"
 
@@ -23,21 +21,12 @@
 
 namespace
 {
-#if defined(TRK_HAS_CUDA_TRACKING)
 void checkGpuError(cudaError_t error, const char* call)
 {
   if (error != cudaSuccess) {
     throw std::runtime_error(std::string(call) + ": " + cudaGetErrorString(error));
   }
 }
-#elif defined(TRK_HAS_HIP_TRACKING)
-void checkGpuError(hipError_t error, const char* call)
-{
-  if (error != hipSuccess) {
-    throw std::runtime_error(std::string(call) + ": " + hipGetErrorString(error));
-  }
-}
-#endif
 } // namespace
 
 namespace o2::trk
@@ -147,26 +136,14 @@ void GPUExternalAllocator::releaseAll()
 void* GPUExternalAllocator::allocateHost(size_t size)
 {
   void* ptr = nullptr;
-#if defined(TRK_HAS_CUDA_TRACKING)
   checkGpuError(cudaHostAlloc(&ptr, size, cudaHostAllocPortable), "cudaHostAlloc");
-#elif defined(TRK_HAS_HIP_TRACKING)
-  checkGpuError(hipHostMalloc(&ptr, size, hipHostMallocPortable), "hipHostMalloc");
-#else
-  throw std::runtime_error("GPUExternalAllocator built without a GPU backend");
-#endif
   return ptr;
 }
 
 void* GPUExternalAllocator::allocateDevice(size_t size)
 {
   void* ptr = nullptr;
-#if defined(TRK_HAS_CUDA_TRACKING)
   checkGpuError(cudaMalloc(&ptr, size), "cudaMalloc");
-#elif defined(TRK_HAS_HIP_TRACKING)
-  checkGpuError(hipMalloc(&ptr, size), "hipMalloc");
-#else
-  throw std::runtime_error("GPUExternalAllocator built without a GPU backend");
-#endif
   return ptr;
 }
 
@@ -176,21 +153,11 @@ void GPUExternalAllocator::freeAllocation(void* ptr, AllocationSpace space)
     return;
   }
 
-#if defined(TRK_HAS_CUDA_TRACKING)
   if (space == AllocationSpace::Host) {
     checkGpuError(cudaFreeHost(ptr), "cudaFreeHost");
   } else {
     checkGpuError(cudaFree(ptr), "cudaFree");
   }
-#elif defined(TRK_HAS_HIP_TRACKING)
-  if (space == AllocationSpace::Host) {
-    checkGpuError(hipHostFree(ptr), "hipHostFree");
-  } else {
-    checkGpuError(hipFree(ptr), "hipFree");
-  }
-#else
-  (void)space;
-#endif
 }
 
 void GPUExternalAllocator::removeFromTagLocked(uint64_t tag, void* ptr)

Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt

@@ -24,8 +24,42 @@ o2_add_library(ALICE3GlobalReconstructionWorkflow
                                      O2::TRKBase
                                      O2::TRKSimulation
                                      O2::ALICE3GlobalReconstruction
+                                     O2::CommonUtils
                                      nlohmann_json::nlohmann_json)
 
+if(CUDA_ENABLED OR HIP_ENABLED)
+  target_compile_definitions(${targetName} PUBLIC TRK_HAS_GPU_TRACKING)
+endif()
+
+if(CUDA_ENABLED)
+  find_package(CUDAToolkit REQUIRED)
+  target_compile_definitions(${targetName} PUBLIC TRK_HAS_CUDA_TRACKING)
+  o2_add_library(ALICE3GlobalReconstructionWorkflowCUDA
+                 TARGETVARNAME cudaTargetName
+                 SOURCES src/TrackerSpecGPU.cxx
+                         ../reconstruction/src/TimeFrameGPU.cxx
+                         ../reconstruction/src/GPUExternalAllocator.cu
+                 PUBLIC_LINK_LIBRARIES
+                         O2::ALICE3GlobalReconstructionWorkflow
+                         O2::ITStrackingCUDA
+                 PRIVATE_LINK_LIBRARIES
+                         CUDA::cudart)
+  target_include_directories(${cudaTargetName} PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
+endif()
+
+if(HIP_ENABLED)
+  target_compile_definitions(${targetName} PUBLIC TRK_HAS_HIP_TRACKING)
+  o2_add_hipified_library(ALICE3GlobalReconstructionWorkflowHIP
+                          SOURCES src/TrackerSpecGPU.cxx
+                                  ../reconstruction/src/TimeFrameGPU.cxx
+                                  ../reconstruction/src/GPUExternalAllocator.cu
+                          PUBLIC_LINK_LIBRARIES
+                                  O2::ALICE3GlobalReconstructionWorkflow
+                                  O2::ITStrackingHIP
+                          PRIVATE_LINK_LIBRARIES
+                                  hip::host)
+endif()
+
 o2_add_executable(reco-workflow
                   SOURCES src/alice3-global-reconstruction-workflow.cxx
                   COMPONENT_NAME alice3-global-reconstruction

Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h

@@ -32,6 +32,10 @@
 
 #include <nlohmann/json.hpp>
 
+#include <memory>
+#include <utility>
+#include <vector>
+
 namespace o2::trk
 {
 class TrackerDPL : public framework::Task
@@ -48,10 +52,15 @@ class TrackerDPL : public framework::Task
   void endOfStream(framework::EndOfStreamContext& ec) final;
   //   void finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj) final;
   void stop() final;
+  template <typename TimeFrameT, typename TrackerTraitsT>
+  void runTracking(framework::ProcessingContext& pc, TimeFrameT& timeFrame, TrackerTraitsT& trackerTraits);
+  const std::shared_ptr<its::ExternalAllocator>& getGPUAllocator() const noexcept { return mGPUAllocator; }
+  void setGPUAllocator(std::shared_ptr<its::ExternalAllocator> allocator) { mGPUAllocator = std::move(allocator); }
 
  private:
   void updateTimeDependentParams(framework::ProcessingContext& pc);
   std::vector<o2::its::TrackingParameters> createTrackingParamsFromConfig();
+  void runGPUTracking(framework::ProcessingContext& pc);
   //   std::unique_ptr<o2::gpu::GPUReconstruction> mRecChain = nullptr;
   //   std::unique_ptr<o2::gpu::GPUChainITS> mChainITS = nullptr;
   //   std::shared_ptr<o2::base::GRPGeomRequest> mGGCCDBRequest;
@@ -61,6 +70,7 @@ class TrackerDPL : public framework::Task
   std::shared_ptr<its::BoundedMemoryResource> mMemoryPool;
   std::shared_ptr<its::ExternalAllocator> mGPUAllocator;
   std::shared_ptr<tbb::task_arena> mTaskArena;
+  std::vector<o2::its::TrackingParameters> mTrackingParams;
   nlohmann::json mHitRecoConfig;
   nlohmann::json mClusterRecoConfig;
   TStopwatch mTimer;