diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt b/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt index 8805c1885b079..1dfcb7a22f725 100644 --- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt +++ b/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} $<$: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) diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx b/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cu similarity index 81% rename from Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx rename to Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cu index df2a2c30b037a..c7b5f1cee50f5 100644 --- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx +++ b/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 -#elif defined(TRK_HAS_HIP_TRACKING) -#include -#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) diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt index be6add9c03483..6a4994e11467b 100644 --- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt +++ b/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 diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h index 006bb4cbf5260..c1e7e051fb3f1 100644 --- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h +++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h @@ -32,6 +32,10 @@ #include +#include +#include +#include + 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 + void runTracking(framework::ProcessingContext& pc, TimeFrameT& timeFrame, TrackerTraitsT& trackerTraits); + const std::shared_ptr& getGPUAllocator() const noexcept { return mGPUAllocator; } + void setGPUAllocator(std::shared_ptr allocator) { mGPUAllocator = std::move(allocator); } private: void updateTimeDependentParams(framework::ProcessingContext& pc); std::vector createTrackingParamsFromConfig(); + void runGPUTracking(framework::ProcessingContext& pc); // std::unique_ptr mRecChain = nullptr; // std::unique_ptr mChainITS = nullptr; // std::shared_ptr mGGCCDBRequest; @@ -61,6 +70,7 @@ class TrackerDPL : public framework::Task std::shared_ptr mMemoryPool; std::shared_ptr mGPUAllocator; std::shared_ptr mTaskArena; + std::vector mTrackingParams; nlohmann::json mHitRecoConfig; nlohmann::json mClusterRecoConfig; TStopwatch mTimer; diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h new file mode 100644 index 0000000000000..f6221e485f369 --- /dev/null +++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h @@ -0,0 +1,226 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#ifndef O2_TRK_TRACKERSPECIMPL_H +#define O2_TRK_TRACKERSPECIMPL_H + +#include "ALICE3GlobalReconstructionWorkflow/TrackerSpec.h" + +#include "CommonDataFormat/IRFrame.h" +#include "DataFormatsTRK/Cluster.h" +#include "DataFormatsTRK/ROFRecord.h" +#include "DetectorsBase/GeometryManager.h" +#include "Field/MagFieldParam.h" +#include "Field/MagneticField.h" +#include "Framework/ControlService.h" +#include "ITStracking/Tracker.h" +#include "SimulationDataFormat/MCCompLabel.h" +#include "SimulationDataFormat/MCEventHeader.h" +#include "SimulationDataFormat/MCTruthContainer.h" +#include "TRKBase/GeometryTGeo.h" +#include "TRKSimulation/Hit.h" + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace o2::trk +{ + +template +void TrackerDPL::runTracking(framework::ProcessingContext& pc, TimeFrameT& timeFrame, TrackerTraitsT& trackerTraits) +{ + o2::its::Tracker<11> itsTracker(&trackerTraits); + timeFrame.setMemoryPool(mMemoryPool); + trackerTraits.setMemoryPool(mMemoryPool); + trackerTraits.setNThreads(mTaskArena->max_concurrency(), mTaskArena); + trackerTraits.adoptTimeFrame(static_cast*>(&timeFrame)); + itsTracker.adoptTimeFrame(timeFrame); + trackerTraits.updateTrackingParameters(mTrackingParams); + timeFrame.initTrackerTopologies(mTrackingParams, 11); + + int nRofs{0}; + if (!mHitRecoConfig.empty()) { + TFile hitsFile(mHitRecoConfig["inputfiles"]["hits"].get().c_str(), "READ"); + TFile mcHeaderFile(mHitRecoConfig["inputfiles"]["mcHeader"].get().c_str(), "READ"); + TTree* hitsTree = hitsFile.Get("o2sim"); + std::vector* trkHit = nullptr; + hitsTree->SetBranchAddress("TRKHit", &trkHit); + + TTree* mcHeaderTree = mcHeaderFile.Get("o2sim"); + auto mcheader = new o2::dataformats::MCEventHeader; + mcHeaderTree->SetBranchAddress("MCEventHeader.", &mcheader); + + o2::base::GeometryManager::loadGeometry(mHitRecoConfig["inputfiles"]["geometry"].get().c_str(), false, true); + auto* gman = o2::trk::GeometryTGeo::Instance(); + + const Long64_t nEvents{hitsTree->GetEntries()}; + LOGP(info, "Starting {} reconstruction from hits for {} events", trackerTraits.getName(), nEvents); + + trackerTraits.setBz(mHitRecoConfig["geometry"]["bz"].get()); + auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mHitRecoConfig["geometry"]["bz"].get() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform); + TGeoGlobalMagField::Instance()->SetField(field); + TGeoGlobalMagField::Instance()->Lock(); + + nRofs = timeFrame.loadROFsFromHitTree(hitsTree, gman, mHitRecoConfig); + const int inROFpileup{mHitRecoConfig.contains("inROFpileup") ? mHitRecoConfig["inROFpileup"].get() : 1}; + timeFrame.getPrimaryVerticesFromMC(mcHeaderTree, nRofs, nEvents, inROFpileup); + } else if (!mClusterRecoConfig.empty()) { + LOGP(info, "Starting {} reconstruction from clusters", trackerTraits.getName()); + + o2::base::GeometryManager::loadGeometry(mClusterRecoConfig["inputfiles"]["geometry"].get().c_str(), false, true); + o2::trk::GeometryTGeo::Instance(); + + trackerTraits.setBz(mClusterRecoConfig["geometry"]["bz"].get()); + auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mClusterRecoConfig["geometry"]["bz"].get() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform); + TGeoGlobalMagField::Instance()->SetField(field); + TGeoGlobalMagField::Instance()->Lock(); + + constexpr int nLayers{11}; + std::array, nLayers> layerClusters; + std::array, nLayers> layerPatterns; + std::array, nLayers> layerROFs; + std::array*, nLayers> layerLabels{}; + + size_t nInputRofs{0}; + for (int iLayer = 0; iLayer < nLayers; ++iLayer) { + layerClusters[iLayer] = pc.inputs().get>(std::format("compClusters_{}", iLayer)); + layerPatterns[iLayer] = pc.inputs().get>(std::format("patterns_{}", iLayer)); + layerROFs[iLayer] = pc.inputs().get>(std::format("ROframes_{}", iLayer)); + nInputRofs = std::max(nInputRofs, layerROFs[iLayer].size()); + if (mIsMC) { + layerLabels[iLayer] = pc.inputs().get*>(std::format("trkmclabels_{}", iLayer)).release(); + } + } + + timeFrame.deriveAndInitTiming(layerROFs); + + const float yPlaneMLOT = 0.0010f; + nRofs = timeFrame.loadROFrameData(layerROFs, layerClusters, layerPatterns, mIsMC ? &layerLabels : nullptr, yPlaneMLOT); + timeFrame.addTruthSeedingVertices(); + } + + const auto trackingLoopStart = std::chrono::steady_clock::now(); + for (size_t iter{0}; iter < mTrackingParams.size(); ++iter) { + LOGP(info, "{}", mTrackingParams[iter].asString()); + trackerTraits.initialiseTimeFrame(iter); + trackerTraits.computeLayerTracklets(iter, -1); + LOGP(info, "Number of tracklets in iteration {}: {}", iter, timeFrame.getNumberOfTracklets()); + trackerTraits.computeLayerCells(iter); + LOGP(info, "Number of cells in iteration {}: {}", iter, timeFrame.getNumberOfCells()); + trackerTraits.findCellsNeighbours(iter); + LOGP(info, "Number of cell neighbours in iteration {}: {}", iter, timeFrame.getNumberOfNeighbours()); + trackerTraits.findRoads(iter); + LOGP(info, "Number of roads in iteration {}: {}", iter, timeFrame.getNumberOfTracks()); + } + const auto trackingLoopElapsedMs = std::chrono::duration_cast(std::chrono::steady_clock::now() - trackingLoopStart).count(); + LOGP(info, "Tracking iterations block took {} ms", trackingLoopElapsedMs); + + if (mIsMC) { + itsTracker.computeTracksMClabels(); + } + + const auto& tracks = timeFrame.getTracks(); + const auto& labels = timeFrame.getTracksLabel(); + std::vector allTracks(tracks.begin(), tracks.end()); + std::vector allLabels; + + int totalTracks = allTracks.size(); + int goodTracks = 0; + int fakeTracks = 0; + + if (mIsMC) { + allLabels.assign(labels.begin(), labels.end()); + for (const auto& label : allLabels) { + if (label.isFake()) { + ++fakeTracks; + } else { + ++goodTracks; + } + } + } + + LOGP(info, "=== Tracking Summary ==="); + LOGP(info, "Total tracks reconstructed: {}", totalTracks); + LOGP(info, "Good tracks: {} ({:.1f}%)", goodTracks, totalTracks > 0 ? 100.0 * goodTracks / totalTracks : 0); + LOGP(info, "Fake tracks: {} ({:.1f}%)", fakeTracks, totalTracks > 0 ? 100.0 * fakeTracks / totalTracks : 0); + + const auto& rofView = timeFrame.getROFOverlapTableView(); + const auto& clockLayer = rofView.getClockLayer(); + const int clockLayerId = rofView.getClock(); + const int64_t anchorBC = timeFrame.getTFAnchorIR().toLong(); + + int highestROF = static_cast(clockLayer.mNROFsTF); + for (const auto& trc : allTracks) { + highestROF = std::max(highestROF, static_cast(clockLayer.getROF(trc.getTimeStamp()))); + } + for (const auto& vtx : timeFrame.getPrimaryVertices()) { + highestROF = std::max(highestROF, static_cast(clockLayer.getROF(vtx.getTimeStamp().lower()))); + } + + std::vector allTrackROFs(highestROF); + for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) { + auto& rof = allTrackROFs[iROF]; + o2::InteractionRecord ir; + ir.setFromLong(anchorBC + static_cast(clockLayer.getROFStartInBC(iROF))); + rof.setBCData(ir); + rof.setROFrame(iROF); + rof.setFirstEntry(0); + rof.setNEntries(0); + } + + std::vector rofEntries(highestROF + 1, 0); + for (const auto& trc : allTracks) { + const int rof = static_cast(clockLayer.getROF(trc.getTimeStamp())); + if (rof >= 0 && rof < highestROF) { + ++rofEntries[rof]; + } + } + std::exclusive_scan(rofEntries.begin(), rofEntries.end(), rofEntries.begin(), 0); + + std::vector irFrames; + irFrames.reserve(allTrackROFs.size()); + const auto& maskView = timeFrame.getROFMaskView(); + const auto rofLenMinus1 = clockLayer.mROFLength > 0 ? clockLayer.mROFLength - 1 : 0; + for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) { + allTrackROFs[iROF].setFirstEntry(rofEntries[iROF]); + allTrackROFs[iROF].setNEntries(rofEntries[iROF + 1] - rofEntries[iROF]); + if (maskView.isROFEnabled(clockLayerId, static_cast(iROF))) { + const auto& bcStart = allTrackROFs[iROF].getBCData(); + auto& irFrame = irFrames.emplace_back(bcStart, bcStart + rofLenMinus1); + irFrame.info = allTrackROFs[iROF].getNEntries(); + } + } + + pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKS", 0}, allTracks); + pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSROF", 0}, allTrackROFs); + pc.outputs().snapshot(o2::framework::Output{"TRK", "IRFRAMES", 0}, irFrames); + if (mIsMC) { + pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSMCTR", 0}, allLabels); + } + + LOGP(info, "TRK pushed {} tracks in {} ROFs and {} IR frames{}", + allTracks.size(), allTrackROFs.size(), irFrames.size(), + mIsMC ? " (with MC labels)" : ""); + + timeFrame.wipe(); +} + +} // namespace o2::trk + +#endif // O2_TRK_TRACKERSPECIMPL_H diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx index 9fb2899ab3ef5..6f9f5561a5ef6 100644 --- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx +++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx @@ -17,6 +17,7 @@ #include #include +#include "CommonUtils/DLLoaderBase.h" #include "CommonDataFormat/IRFrame.h" #include "DataFormatsTRK/Cluster.h" #include "DataFormatsTRK/ROFRecord.h" @@ -36,12 +37,8 @@ #include "TRKBase/SegmentationChip.h" #include "TRKSimulation/Hit.h" #include "ALICE3GlobalReconstruction/TimeFrame.h" -#ifdef TRK_HAS_GPU_TRACKING -#include "ALICE3GlobalReconstruction/TimeFrameGPU.h" -#include "ALICE3GlobalReconstruction/GPUExternalAllocator.h" -#include "ITStrackingGPU/TrackerTraitsGPU.h" -#endif #include "ALICE3GlobalReconstructionWorkflow/TrackerSpec.h" +#include "ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h" #include #ifdef O2_WITH_ACTS @@ -58,6 +55,18 @@ namespace trk { using Vertex = o2::dataformats::Vertex>; +namespace +{ +class ALICE3TrackingBackendLoader : public o2::utils::DLLoaderBase +{ + O2DLLoaderDef(ALICE3TrackingBackendLoader) +}; + +O2DLLoaderImpl(ALICE3TrackingBackendLoader) + + constexpr const char* kGPUBackendFunction = "runALICE3GPUTracking"; +} // namespace + TrackerDPL::TrackerDPL(std::shared_ptr gr, bool isMC, const std::string& hitRecoConfigFileName, @@ -249,220 +258,20 @@ void TrackerDPL::run(ProcessingContext& pc) mTaskArena = std::make_shared(1); /// TODO: make it configurable } - auto trackingParams = createTrackingParamsFromConfig(); + mTrackingParams = createTrackingParamsFromConfig(); auto cput = mTimer.CpuTime(); auto realt = mTimer.RealTime(); mTimer.Start(false); const bool useGPU = mDeviceType != o2::gpu::gpudatatypes::DeviceType::CPU; -#ifndef TRK_HAS_GPU_TRACKING - if (useGPU) { - LOGP(fatal, "TRK GPU tracking was requested but this build has no TRK GPU tracking backend"); - } -#else -#ifdef TRK_HAS_CUDA_TRACKING - if (useGPU && mDeviceType != o2::gpu::gpudatatypes::DeviceType::CUDA) { - LOGP(fatal, "This build provides the CUDA TRK tracking backend only, but device type {} was requested", static_cast(mDeviceType)); - } -#elif defined(TRK_HAS_HIP_TRACKING) - if (useGPU && mDeviceType != o2::gpu::gpudatatypes::DeviceType::HIP) { - LOGP(fatal, "This build provides the HIP TRK tracking backend only, but device type {} was requested", static_cast(mDeviceType)); - } -#endif -#endif - - auto runTracking = [&](auto& timeFrame, auto& trackerTraits) { - o2::its::Tracker<11> itsTracker(&trackerTraits); - timeFrame.setMemoryPool(mMemoryPool); - trackerTraits.setMemoryPool(mMemoryPool); - trackerTraits.setNThreads(mTaskArena->max_concurrency(), mTaskArena); - trackerTraits.adoptTimeFrame(static_cast*>(&timeFrame)); - itsTracker.adoptTimeFrame(timeFrame); - trackerTraits.updateTrackingParameters(trackingParams); - - int nRofs{0}; - if (!mHitRecoConfig.empty()) { - TFile hitsFile(mHitRecoConfig["inputfiles"]["hits"].get().c_str(), "READ"); - TFile mcHeaderFile(mHitRecoConfig["inputfiles"]["mcHeader"].get().c_str(), "READ"); - TTree* hitsTree = hitsFile.Get("o2sim"); - std::vector* trkHit = nullptr; - hitsTree->SetBranchAddress("TRKHit", &trkHit); - - TTree* mcHeaderTree = mcHeaderFile.Get("o2sim"); - auto mcheader = new o2::dataformats::MCEventHeader; - mcHeaderTree->SetBranchAddress("MCEventHeader.", &mcheader); - - o2::base::GeometryManager::loadGeometry(mHitRecoConfig["inputfiles"]["geometry"].get().c_str(), false, true); - auto* gman = o2::trk::GeometryTGeo::Instance(); - - const Long64_t nEvents{hitsTree->GetEntries()}; - LOGP(info, "Starting {} reconstruction from hits for {} events", trackerTraits.getName(), nEvents); - - trackerTraits.setBz(mHitRecoConfig["geometry"]["bz"].get()); - auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mHitRecoConfig["geometry"]["bz"].get() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform); - TGeoGlobalMagField::Instance()->SetField(field); - TGeoGlobalMagField::Instance()->Lock(); - - nRofs = timeFrame.loadROFsFromHitTree(hitsTree, gman, mHitRecoConfig); - const int inROFpileup{mHitRecoConfig.contains("inROFpileup") ? mHitRecoConfig["inROFpileup"].get() : 1}; - timeFrame.getPrimaryVerticesFromMC(mcHeaderTree, nRofs, nEvents, inROFpileup); - } else if (!mClusterRecoConfig.empty()) { - LOGP(info, "Starting {} reconstruction from clusters", trackerTraits.getName()); - - o2::base::GeometryManager::loadGeometry(mClusterRecoConfig["inputfiles"]["geometry"].get().c_str(), false, true); - o2::trk::GeometryTGeo::Instance(); - - trackerTraits.setBz(mClusterRecoConfig["geometry"]["bz"].get()); - auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mClusterRecoConfig["geometry"]["bz"].get() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform); - TGeoGlobalMagField::Instance()->SetField(field); - TGeoGlobalMagField::Instance()->Lock(); - - constexpr int nLayers{11}; - std::array, nLayers> layerClusters; - std::array, nLayers> layerPatterns; - std::array, nLayers> layerROFs; - std::array*, nLayers> layerLabels{}; - - size_t nInputRofs{0}; - for (int iLayer = 0; iLayer < nLayers; ++iLayer) { - layerClusters[iLayer] = pc.inputs().get>(std::format("compClusters_{}", iLayer)); - layerPatterns[iLayer] = pc.inputs().get>(std::format("patterns_{}", iLayer)); - layerROFs[iLayer] = pc.inputs().get>(std::format("ROframes_{}", iLayer)); - nInputRofs = std::max(nInputRofs, layerROFs[iLayer].size()); - if (mIsMC) { - layerLabels[iLayer] = pc.inputs().get*>(std::format("trkmclabels_{}", iLayer)).release(); - } - } - - timeFrame.deriveAndInitTiming(layerROFs); - - const float yPlaneMLOT = 0.0010f; - nRofs = timeFrame.loadROFrameData(layerROFs, layerClusters, layerPatterns, mIsMC ? &layerLabels : nullptr, yPlaneMLOT); - timeFrame.addTruthSeedingVertices(); - } - - const auto trackingLoopStart = std::chrono::steady_clock::now(); - for (size_t iter{0}; iter < trackingParams.size(); ++iter) { - LOGP(info, "{}", trackingParams[iter].asString()); - trackerTraits.initialiseTimeFrame(iter); - trackerTraits.computeLayerTracklets(iter, -1); - LOGP(info, "Number of tracklets in iteration {}: {}", iter, timeFrame.getNumberOfTracklets()); - trackerTraits.computeLayerCells(iter); - LOGP(info, "Number of cells in iteration {}: {}", iter, timeFrame.getNumberOfCells()); - trackerTraits.findCellsNeighbours(iter); - LOGP(info, "Number of cell neighbours in iteration {}: {}", iter, timeFrame.getNumberOfNeighbours()); - trackerTraits.findRoads(iter); - LOGP(info, "Number of roads in iteration {}: {}", iter, timeFrame.getNumberOfTracks()); - } - const auto trackingLoopElapsedMs = std::chrono::duration_cast(std::chrono::steady_clock::now() - trackingLoopStart).count(); - LOGP(info, "Tracking iterations block took {} ms", trackingLoopElapsedMs); - - if (mIsMC) { - itsTracker.computeTracksMClabels(); - } - - const auto& tracks = timeFrame.getTracks(); - const auto& labels = timeFrame.getTracksLabel(); - std::vector allTracks(tracks.begin(), tracks.end()); - std::vector allLabels; - - int totalTracks = allTracks.size(); - int goodTracks = 0; - int fakeTracks = 0; - - if (mIsMC) { - allLabels.assign(labels.begin(), labels.end()); - for (const auto& label : allLabels) { - if (label.isFake()) { - ++fakeTracks; - } else { - ++goodTracks; - } - } - } - - LOGP(info, "=== Tracking Summary ==="); - LOGP(info, "Total tracks reconstructed: {}", totalTracks); - LOGP(info, "Good tracks: {} ({:.1f}%)", goodTracks, totalTracks > 0 ? 100.0 * goodTracks / totalTracks : 0); - LOGP(info, "Fake tracks: {} ({:.1f}%)", fakeTracks, totalTracks > 0 ? 100.0 * fakeTracks / totalTracks : 0); - - const auto& rofView = timeFrame.getROFOverlapTableView(); - const auto& clockLayer = rofView.getClockLayer(); - const int clockLayerId = rofView.getClock(); - const int64_t anchorBC = timeFrame.getTFAnchorIR().toLong(); - int highestROF = static_cast(clockLayer.mNROFsTF); - for (const auto& trc : allTracks) { - highestROF = std::max(highestROF, static_cast(clockLayer.getROF(trc.getTimeStamp()))); - } - for (const auto& vtx : timeFrame.getPrimaryVertices()) { - highestROF = std::max(highestROF, static_cast(clockLayer.getROF(vtx.getTimeStamp().lower()))); - } - - std::vector allTrackROFs(highestROF); - for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) { - auto& rof = allTrackROFs[iROF]; - o2::InteractionRecord ir; - ir.setFromLong(anchorBC + static_cast(clockLayer.getROFStartInBC(iROF))); - rof.setBCData(ir); - rof.setROFrame(iROF); - rof.setFirstEntry(0); - rof.setNEntries(0); - } - - std::vector rofEntries(highestROF + 1, 0); - for (const auto& trc : allTracks) { - const int rof = static_cast(clockLayer.getROF(trc.getTimeStamp())); - if (rof >= 0 && rof < highestROF) { - ++rofEntries[rof]; - } - } - std::exclusive_scan(rofEntries.begin(), rofEntries.end(), rofEntries.begin(), 0); - - std::vector irFrames; - irFrames.reserve(allTrackROFs.size()); - const auto& maskView = timeFrame.getROFMaskView(); - const auto rofLenMinus1 = clockLayer.mROFLength > 0 ? clockLayer.mROFLength - 1 : 0; - for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) { - allTrackROFs[iROF].setFirstEntry(rofEntries[iROF]); - allTrackROFs[iROF].setNEntries(rofEntries[iROF + 1] - rofEntries[iROF]); - if (maskView.isROFEnabled(clockLayerId, static_cast(iROF))) { - const auto& bcStart = allTrackROFs[iROF].getBCData(); - auto& irFrame = irFrames.emplace_back(bcStart, bcStart + rofLenMinus1); - irFrame.info = allTrackROFs[iROF].getNEntries(); - } - } - - pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKS", 0}, allTracks); - pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSROF", 0}, allTrackROFs); - pc.outputs().snapshot(o2::framework::Output{"TRK", "IRFRAMES", 0}, irFrames); - if (mIsMC) { - pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSMCTR", 0}, allLabels); - } - - LOGP(info, "TRK pushed {} tracks in {} ROFs and {} IR frames{}", - allTracks.size(), allTrackROFs.size(), irFrames.size(), - mIsMC ? " (with MC labels)" : ""); - - timeFrame.wipe(); - }; - -#ifdef TRK_HAS_GPU_TRACKING if (useGPU) { - o2::trk::TimeFrameGPU<11> timeFrame; - o2::its::TrackerTraitsGPU<11> itsTrackerTraits; - if (!mGPUAllocator) { - mGPUAllocator = std::make_shared(); - } - timeFrame.setFrameworkAllocator(mGPUAllocator.get()); - runTracking(timeFrame, itsTrackerTraits); - } else -#endif - { + runGPUTracking(pc); + } else { o2::trk::TimeFrame<11> timeFrame; o2::its::TrackerTraits<11> itsTrackerTraits; - runTracking(timeFrame, itsTrackerTraits); + runTracking(pc, timeFrame, itsTrackerTraits); } pc.services().get().endOfStream(); @@ -472,6 +281,29 @@ void TrackerDPL::run(ProcessingContext& pc) LOGP(info, "CPU Reconstruction time for this TF {} s (cpu), {} s (wall)", mTimer.CpuTime() - cput, mTimer.RealTime() - realt); } +void TrackerDPL::runGPUTracking(ProcessingContext& pc) +{ + auto& loader = ALICE3TrackingBackendLoader::Instance(); + switch (mDeviceType) { + case o2::gpu::gpudatatypes::DeviceType::CUDA: +#ifdef TRK_HAS_CUDA_TRACKING + loader.executeFunctionAlias("O2ALICE3GlobalReconstructionWorkflowCUDA", kGPUBackendFunction, this, &pc); + return; +#else + LOGP(fatal, "CUDA TRK GPU tracking was requested but this build has no CUDA TRK GPU tracking backend"); +#endif + case o2::gpu::gpudatatypes::DeviceType::HIP: +#ifdef TRK_HAS_HIP_TRACKING + loader.executeFunctionAlias("O2ALICE3GlobalReconstructionWorkflowHIP", kGPUBackendFunction, this, &pc); + return; +#else + LOGP(fatal, "HIP TRK GPU tracking was requested but this build has no HIP TRK GPU tracking backend"); +#endif + default: + LOGP(fatal, "Unsupported TRK GPU device type {}", static_cast(mDeviceType)); + } +} + void TrackerDPL::endOfStream(EndOfStreamContext& ec) { LOGF(info, "TRK CA-Tracker total timing: Cpu: %.3e Real: %.3e s in %d slots", mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1); diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx new file mode 100644 index 0000000000000..ea98ab3f852e5 --- /dev/null +++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx @@ -0,0 +1,28 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#include "ALICE3GlobalReconstruction/GPUExternalAllocator.h" +#include "ALICE3GlobalReconstruction/TimeFrameGPU.h" +#include "ALICE3GlobalReconstructionWorkflow/TrackerSpec.h" +#include "ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h" +#include "ITStrackingGPU/TrackerTraitsGPU.h" + +extern "C" int runALICE3GPUTracking(o2::trk::TrackerDPL* tracker, o2::framework::ProcessingContext* pc) +{ + o2::trk::TimeFrameGPU<11> timeFrame; + o2::its::TrackerTraitsGPU<11> itsTrackerTraits; + if (!tracker->getGPUAllocator()) { + tracker->setGPUAllocator(std::make_shared()); + } + timeFrame.setFrameworkAllocator(tracker->getGPUAllocator().get()); + tracker->runTracking(*pc, timeFrame, itsTrackerTraits); + return 0; +}