ITS-GPU: Fixup for running on PbPb in async mode (#12489)

mconcas · chiarazampolli · commit f7b582e8961d · 2024-01-15T16:30:24.000+01:00
* Fix tracking mode not correctly forwarded

* Make ITS GPU running on PbPb in async
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h
@@ -190,7 +190,7 @@ class TimeFrameGPU : public TimeFrame
   void initialiseHybrid(const int, const TrackingParameters&, const int, IndexTableUtils* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr);
   void initDevice(const int, IndexTableUtils*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int);
   void initDeviceSAFitting();
-  void loadTrackingFrameInfoDevice();
+  void loadTrackingFrameInfoDevice(const int);
   void loadUnsortedClustersDevice();
   void loadClustersDevice();
   void loadTrackletsDevice();
@@ -199,7 +199,7 @@ class TimeFrameGPU : public TimeFrame
   void loadTrackSeedsChi2Device();
   void loadRoadsDevice();
   void loadTrackSeedsDevice(std::vector<CellSeed>&);
-  void createTrackITSExtDevice(const std::vector<CellSeed>& seeds);
+  void createTrackITSExtDevice(const std::vector<CellSeed>&);
   void createTrackITSExtDevice(); // deprecated
   void downloadTrackITSExtDevice();
   void initDeviceChunks(const int, const int);
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu
@@ -342,7 +342,7 @@ void TimeFrameGPU<nLayers>::allocMemAsync(void** ptr, size_t size, Stream* strPt
   if (extAllocator) {
     *ptr = mAllocator->allocate(size);
   } else {
-    LOGP(debug, "Calling default CUDA allocator");
+    LOGP(info, "Calling default CUDA allocator");
     checkGPUError(cudaMallocAsync(reinterpret_cast<void**>(ptr), size, strPtr->get()));
   }
 }
@@ -373,17 +373,10 @@ void TimeFrameGPU<nLayers>::registerHostMemory(const int maxLayers)
 template <int nLayers>
 void TimeFrameGPU<nLayers>::unregisterHostMemory(const int maxLayers)
 {
-  if (!mHostRegistered) {
-    return;
-  }
-  for (auto iLayer{0}; iLayer < maxLayers; ++iLayer) {
-    checkGPUError(cudaHostUnregister(mClusters[iLayer].data()));
-    checkGPUError(cudaHostUnregister(mNClustersPerROF[iLayer].data()));
-    checkGPUError(cudaHostUnregister(mIndexTables[iLayer].data()));
+  for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {
+    checkGPUError(cudaHostUnregister(mTrackingFrameInfo[iLayer].data()));
   }
-  checkGPUError(cudaHostUnregister(mHostNTracklets.data()));
-  checkGPUError(cudaHostUnregister(mHostNCells.data()));
-  mHostRegistered = false;
+  checkGPUError(cudaHostUnregister(mTrackingFrameInfoDevice.data()));
 }
 
 template <int nLayers>
@@ -522,17 +515,21 @@ void TimeFrameGPU<nLayers>::loadClustersDevice()
 }
 
 template <int nLayers>
-void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice()
+void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice(const int iteration)
 {
   for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {
     LOGP(debug, "gpu-transfer: loading {} tfinfo on layer {}, for {} MB.", mTrackingFrameInfo[iLayer].size(), iLayer, mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo) / MB);
     allocMemAsync(reinterpret_cast<void**>(&mTrackingFrameInfoDevice[iLayer]), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), nullptr, getExtAllocator());
     // Register and move data
-    checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));
+    if (!iteration) {
+      checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));
+    }
     checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDevice[iLayer], mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
   }
   allocMemAsync(reinterpret_cast<void**>(&mTrackingFrameInfoDeviceArray), nLayers * sizeof(TrackingFrameInfo*), nullptr, getExtAllocator());
-  checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));
+  if (!iteration) {
+    checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));
+  }
   checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDeviceArray, mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
 }
 
@@ -570,7 +567,7 @@ template <int nLayers>
 void TimeFrameGPU<nLayers>::loadRoadsDevice()
 {
   LOGP(debug, "gpu-transfer: loading {} roads, for {} MB.", mRoads.size(), mRoads.size() * sizeof(Road<nLayers - 2>) / MB);
-  allocMemAsync(reinterpret_cast<void**>(&mRoadsDevice), mRoads.size() * sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), false);
+  allocMemAsync(reinterpret_cast<void**>(&mRoadsDevice), mRoads.size() * sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), getExtAllocator());
   checkGPUError(cudaHostRegister(mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaHostRegisterPortable));
   checkGPUError(cudaMemcpyAsync(mRoadsDevice, mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
 }
@@ -579,7 +576,7 @@ template <int nLayers>
 void TimeFrameGPU<nLayers>::loadTrackSeedsDevice(std::vector<CellSeed>& seeds)
 {
   LOGP(debug, "gpu-transfer: loading {} track seeds, for {} MB.", seeds.size(), seeds.size() * sizeof(CellSeed) / MB);
-  allocMemAsync(reinterpret_cast<void**>(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), &(mGpuStreams[0]), false);
+  allocMemAsync(reinterpret_cast<void**>(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), &(mGpuStreams[0]), getExtAllocator());
   checkGPUError(cudaHostRegister(seeds.data(), seeds.size() * sizeof(CellSeed), cudaHostRegisterPortable));
   checkGPUError(cudaMemcpyAsync(mTrackSeedsDevice, seeds.data(), seeds.size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
 }
@@ -589,8 +586,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice(const std::vector<CellSeed>&
 {
   mTrackITSExt.clear();
   mTrackITSExt.resize(seeds.size());
-  LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);
-  allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);
+  LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);
+  allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());
   checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));
   checkGPUError(cudaHostRegister(mTrackITSExt.data(), seeds.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));
 }
@@ -600,8 +597,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice()
 {
   mTrackITSExt.clear();
   mTrackITSExt.resize(mRoads.size());
-  LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);
-  allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), mRoads.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);
+  LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);
+  allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), mRoads.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());
   checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, mRoads.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));
   checkGPUError(cudaHostRegister(mTrackITSExt.data(), mRoads.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));
 }
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu
@@ -870,9 +870,7 @@ template <int nLayers>
 void TrackerTraitsGPU<nLayers>::initialiseTimeFrame(const int iteration)
 {
   mTimeFrameGPU->initialiseHybrid(iteration, mTrkParams[iteration], nLayers);
-  mTimeFrameGPU->loadClustersDevice();
-  mTimeFrameGPU->loadUnsortedClustersDevice();
-  mTimeFrameGPU->loadTrackingFrameInfoDevice();
+  mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration);
 }
 
 template <int nLayers>
@@ -1165,14 +1163,12 @@ template <int nLayers>
 void TrackerTraitsGPU<nLayers>::computeTrackletsHybrid(const int iteration)
 {
   TrackerTraits::computeLayerTracklets(iteration);
-  mTimeFrameGPU->loadTrackletsDevice();
 }
 
 template <int nLayers>
 void TrackerTraitsGPU<nLayers>::computeCellsHybrid(const int iteration)
 {
   TrackerTraits::computeLayerCells(iteration);
-  mTimeFrameGPU->loadCellsDevice();
 };
 
 template <int nLayers>
@@ -1203,6 +1199,10 @@ void TrackerTraitsGPU<nLayers>::findRoadsHybrid(const int iteration)
       }
       trackSeeds.insert(trackSeeds.end(), updatedCellSeed.begin(), updatedCellSeed.end());
     }
+    if (!trackSeeds.size()) {
+      LOGP(info, "No track seeds found, skipping track finding");
+      continue;
+    }
     mTimeFrameGPU->createTrackITSExtDevice(trackSeeds);
     mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);
 
@@ -1265,12 +1265,15 @@ void TrackerTraitsGPU<nLayers>::findRoadsHybrid(const int iteration)
       mTimeFrame->getTracks(std::min(rofs[0], rofs[1])).emplace_back(track);
     }
   }
+  if (iteration == 2) {
+    mTimeFrameGPU->unregisterHostMemory(0); // FIXME this needs to work also with sync
+  }
 };
 
 template <int nLayers>
 void TrackerTraitsGPU<nLayers>::findTracksHybrid(const int iteration)
 {
-  // LOGP(info, "propagator device pointer: {}", (void*)mTimeFrameGPU->getDevicePropagator());
+  LOGP(info, "propagator device pointer: {}", (void*)mTimeFrameGPU->getDevicePropagator());
   mTimeFrameGPU->createTrackITSExtDevice();
   gpu::fitTracksKernel<<<20, 512>>>(mTimeFrameGPU->getDeviceArrayClusters(),          // Cluster** foundClusters,
                                     mTimeFrameGPU->getDeviceArrayUnsortedClusters(),  // Cluster** foundUnsortedClusters,
@@ -1287,7 +1290,6 @@ void TrackerTraitsGPU<nLayers>::findTracksHybrid(const int iteration)
                                     mTrkParams[0].MaxChi2NDF,                         // float maxChi2NDF,
                                     mTimeFrameGPU->getDevicePropagator());            // const o2::base::Propagator* propagator
   mTimeFrameGPU->downloadTrackITSExtDevice();
-  discardResult(cudaDeviceSynchronize());
   auto& tracks = mTimeFrameGPU->getTrackITSExt();
   std::sort(tracks.begin(), tracks.end(),
             [](TrackITSExt& track1, TrackITSExt& track2) { return track1.isBetter(track2, 1.e6f); });
diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx
@@ -106,6 +106,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,
   double total{0.};
   mTraits->UpdateTrackingParameters(mTrkParams);
   for (int iteration = 0; iteration < (int)mTrkParams.size(); ++iteration) {
+    LOGP(info, "Iteration {}", iteration);
     total += evaluateTask(&Tracker::initialiseTimeFrameHybrid, "Hybrid Timeframe initialisation", logger, iteration);
     total += evaluateTask(&Tracker::computeTrackletsHybrid, "Hybrid Tracklet finding", logger, iteration);
     logger(fmt::format("\t- Number of tracklets: {}", mTraits->getTFNumberOfTracklets()));
@@ -115,7 +116,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,
     }
     float trackletsPerCluster = mTraits->getTFNumberOfClusters() > 0 ? float(mTraits->getTFNumberOfTracklets()) / mTraits->getTFNumberOfClusters() : 0.f;
     if (trackletsPerCluster > mTrkParams[iteration].TrackletsPerClusterLimit) {
-      error(fmt::format("Too many tracklets per cluster ({}), check the detector status and/or the selections.", trackletsPerCluster));
+      error(fmt::format("Too many tracklets per cluster ({}), check the detector status and/or the selections. Current limit is {}", trackletsPerCluster, mTrkParams[iteration].TrackletsPerClusterLimit));
       break;
     }
 
@@ -127,7 +128,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,
     }
     float cellsPerCluster = mTraits->getTFNumberOfClusters() > 0 ? float(mTraits->getTFNumberOfCells()) / mTraits->getTFNumberOfClusters() : 0.f;
     if (cellsPerCluster > mTrkParams[iteration].CellsPerClusterLimit) {
-      error(fmt::format("Too many cells per cluster ({}), check the detector status and/or the selections.", cellsPerCluster));
+      error(fmt::format("Too many cells per cluster ({}), check the detector status and/or the selections. Current limit is {}", cellsPerCluster, mTrkParams[iteration].CellsPerClusterLimit));
       break;
     }
     total += evaluateTask(&Tracker::findCellsNeighboursHybrid, "Hybrid Neighbour finding", logger, iteration);
diff --git a/Detectors/ITSMFT/ITS/workflow/src/RecoWorkflow.cxx b/Detectors/ITSMFT/ITS/workflow/src/RecoWorkflow.cxx
@@ -62,6 +62,7 @@ framework::WorkflowSpec getWorkflow(bool useMC,
         cfg.runITSTracking = true;
         cfg.itsTriggerType = useTrig;
         cfg.itsOverrBeamEst = overrideBeamPosition;
+        cfg.itsTrackingMode = trmode == "sync" ? 0 : (trmode == "async" ? 1 : 2);
 
         Inputs ggInputs;
         auto ggRequest = std::make_shared<o2::base::GRPGeomRequest>(false, true, false, true, true, o2::base::GRPGeomRequest::Aligned, ggInputs, true);
diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx
@@ -676,7 +676,7 @@ void* GPUReconstruction::AllocateUnmanagedMemory(size_t size, int type)
     char* retVal;
     GPUProcessor::computePointerWithAlignment(pool, retVal, size);
     if (pool > poolend) {
-      GPUError("Insufficient unmanaged memory: missing %lu", (size_t)((char*)pool - (char*)poolend));
+      GPUError("Insufficient unmanaged memory: missing %lu bytes", (size_t)((char*)pool - (char*)poolend));
       throw std::bad_alloc();
     }
     UpdateMaxMemoryUsed();
diff --git a/GPU/GPUTracking/Global/GPUChainITS.cxx b/GPU/GPUTracking/Global/GPUChainITS.cxx
@@ -90,7 +90,7 @@ o2::its::TimeFrame* GPUChainITS::GetITSTimeframe()
     mFrameworkAllocator.reset(new o2::its::GPUFrameworkExternalAllocator);
     mFrameworkAllocator->setReconstructionFramework(rec());
     mITSTimeFrame->setExternalAllocator(mFrameworkAllocator.get());
-    LOGP(debug, "GPUChainITS is giving me ps: {} prop: {}", (void*)processorsShadow(), (void*)processorsShadow()->calibObjects.o2Propagator);
+    LOGP(info, "GPUChainITS is giving me ps: {} prop: {} allocator: {}", (void*)processorsShadow(), (void*)processorsShadow()->calibObjects.o2Propagator, (void*)mFrameworkAllocator.get());
     mITSTimeFrame->setDevicePropagator(processorsShadow()->calibObjects.o2Propagator);
   }
 #endif
diff --git a/GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h b/GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h
@@ -126,6 +126,7 @@ class GPURecoWorkflowSpec : public o2::framework::Task
     int lumiScaleType = 0; // 0=off, 1=CTP, 2=TPC scalers
     bool outputErrorQA = false;
     bool runITSTracking = false;
+    int itsTrackingMode = 0; // 0=sync, 1=async, 2=cosmics
     bool itsOverrBeamEst = false;
     bool tpcTriggerHandling = false;
   };
@@ -232,7 +233,6 @@ class GPURecoWorkflowSpec : public o2::framework::Task
   bool mPropagatorInstanceCreated = false;
   bool mITSRunVertexer = false;
   bool mITSCosmicsProcessing = false;
-  std::string mITSMode = "sync";
 };
 
 } // end namespace gpu
diff --git a/GPU/Workflow/src/GPUWorkflowITS.cxx b/GPU/Workflow/src/GPUWorkflowITS.cxx
@@ -184,6 +184,7 @@ int GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc)
   pattIt = patterns.begin();
   std::vector<int> savedROF;
   auto logger = [&](std::string s) { LOG(info) << s; };
+  auto fatalLogger = [&](std::string s) { LOG(fatal) << s; };
   auto errorLogger = [&](std::string s) { LOG(error) << s; };
 
   o2::its::FastMultEst multEst; // mult estimator
@@ -252,7 +253,11 @@ int GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc)
 
     mITSTimeFrame->setMultiplicityCutMask(processingMask);
     // Run CA tracker
-    mITSTracker->clustersToTracksHybrid(logger, errorLogger);
+    if (!mSpecConfig.itsTrackingMode) {
+      mITSTracker->clustersToTracksHybrid(logger, errorLogger);
+    } else {
+      mITSTracker->clustersToTracksHybrid(logger, fatalLogger);
+    }
     size_t totTracks{mITSTimeFrame->getNumberOfTracks()}, totClusIDs{mITSTimeFrame->getNumberOfUsedClusters()};
     allTracks.reserve(totTracks);
     allClusIdx.reserve(totClusIDs);
@@ -303,7 +308,6 @@ int GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc)
 
 void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
 {
-  std::transform(mITSMode.begin(), mITSMode.end(), mITSMode.begin(), [](unsigned char c) { return std::tolower(c); });
   o2::its::VertexerTraits* vtxTraits = nullptr;
   o2::its::TrackerTraits* trkTraits = nullptr;
   mGPUReco->GetITSTraits(trkTraits, vtxTraits, mITSTimeFrame);
@@ -315,25 +319,27 @@ void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
   mITSCosmicsProcessing = false;
   std::vector<o2::its::TrackingParameters> trackParams;
 
-  if (mITSMode == "async") {
+  if (mSpecConfig.itsTrackingMode == 1) {
     trackParams.resize(3);
     for (auto& param : trackParams) {
       param.ZBins = 64;
       param.PhiBins = 32;
+      param.CellsPerClusterLimit = 1.e3f;
+      param.TrackletsPerClusterLimit = 1.e3f;
     }
     trackParams[1].TrackletMinPt = 0.2f;
     trackParams[1].CellDeltaTanLambdaSigma *= 2.;
     trackParams[2].TrackletMinPt = 0.1f;
     trackParams[2].CellDeltaTanLambdaSigma *= 4.;
     trackParams[2].MinTrackLength = 4;
     LOG(info) << "Initializing tracker in async. phase reconstruction with " << trackParams.size() << " passes";
-  } else if (mITSMode == "sync") {
+  } else if (mSpecConfig.itsTrackingMode == 0) {
     trackParams.resize(1);
     trackParams[0].ZBins = 64;
     trackParams[0].PhiBins = 32;
     trackParams[0].MinTrackLength = 4;
     LOG(info) << "Initializing tracker in sync. phase reconstruction with " << trackParams.size() << " passes";
-  } else if (mITSMode == "cosmics") {
+  } else if (mSpecConfig.itsTrackingMode == 2) {
     mITSCosmicsProcessing = true;
     mITSRunVertexer = false;
     trackParams.resize(1);
@@ -347,8 +353,6 @@ void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
     trackParams[0].TrackletsPerClusterLimit = 100.;
     trackParams[0].CellsPerClusterLimit = 100.;
     LOG(info) << "Initializing tracker in reconstruction for cosmics with " << trackParams.size() << " passes";
-  } else {
-    throw std::runtime_error(fmt::format("Unsupported ITS tracking mode {:s} ", mITSMode));
   }
 
   for (auto& params : trackParams) {

Original file line number	Diff line number	Diff line change
`@@ -342,7 +342,7 @@ void TimeFrameGPU<nLayers>::allocMemAsync(void** ptr, size_t size, Stream* strPt`
`342`	`342`	`if (extAllocator) {`
`343`	`343`	`*ptr = mAllocator->allocate(size);`
`344`	`344`	`} else {`
`345`		`- LOGP(debug, "Calling default CUDA allocator");`
	`345`	`+ LOGP(info, "Calling default CUDA allocator");`
`346`	`346`	`checkGPUError(cudaMallocAsync(reinterpret_cast<void**>(ptr), size, strPtr->get()));`
`347`	`347`	`}`
`348`	`348`	`}`
`@@ -373,17 +373,10 @@ void TimeFrameGPU<nLayers>::registerHostMemory(const int maxLayers)`
`373`	`373`	`template <int nLayers>`
`374`	`374`	`void TimeFrameGPU<nLayers>::unregisterHostMemory(const int maxLayers)`
`375`	`375`	`{`
`376`		`- if (!mHostRegistered) {`
`377`		`- return;`
`378`		`- }`
`379`		`- for (auto iLayer{0}; iLayer < maxLayers; ++iLayer) {`
`380`		`- checkGPUError(cudaHostUnregister(mClusters[iLayer].data()));`
`381`		`- checkGPUError(cudaHostUnregister(mNClustersPerROF[iLayer].data()));`
`382`		`- checkGPUError(cudaHostUnregister(mIndexTables[iLayer].data()));`
	`376`	`+ for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {`
	`377`	`+ checkGPUError(cudaHostUnregister(mTrackingFrameInfo[iLayer].data()));`
`383`	`378`	`}`
`384`		`- checkGPUError(cudaHostUnregister(mHostNTracklets.data()));`
`385`		`- checkGPUError(cudaHostUnregister(mHostNCells.data()));`
`386`		`- mHostRegistered = false;`
	`379`	`+ checkGPUError(cudaHostUnregister(mTrackingFrameInfoDevice.data()));`
`387`	`380`	`}`
`388`	`381`
`389`	`382`	`template <int nLayers>`
`@@ -522,17 +515,21 @@ void TimeFrameGPU<nLayers>::loadClustersDevice()`
`522`	`515`	`}`
`523`	`516`
`524`	`517`	`template <int nLayers>`
`525`		`-void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice()`
	`518`	`+void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice(const int iteration)`
`526`	`519`	`{`
`527`	`520`	`for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {`
`528`	`521`	`LOGP(debug, "gpu-transfer: loading {} tfinfo on layer {}, for {} MB.", mTrackingFrameInfo[iLayer].size(), iLayer, mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo) / MB);`
`529`	`522`	`allocMemAsync(reinterpret_cast<void*>(&mTrackingFrameInfoDevice[iLayer]), mTrackingFrameInfo[iLayer].size() sizeof(TrackingFrameInfo), nullptr, getExtAllocator());`
`530`	`523`	`// Register and move data`
`531`		`- checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));`
	`524`	`+ if (!iteration) {`
	`525`	`+ checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));`
	`526`	`+ }`
`532`	`527`	`checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDevice[iLayer], mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`533`	`528`	`}`
`534`	`529`	`allocMemAsync(reinterpret_cast<void*>(&mTrackingFrameInfoDeviceArray), nLayers sizeof(TrackingFrameInfo*), nullptr, getExtAllocator());`
`535`		`- checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));`
	`530`	`+ if (!iteration) {`
	`531`	`+ checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));`
	`532`	`+ }`
`536`	`533`	`checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDeviceArray, mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`537`	`534`	`}`
`538`	`535`
`@@ -570,7 +567,7 @@ template <int nLayers>`
`570`	`567`	`void TimeFrameGPU<nLayers>::loadRoadsDevice()`
`571`	`568`	`{`
`572`	`569`	`LOGP(debug, "gpu-transfer: loading {} roads, for {} MB.", mRoads.size(), mRoads.size() * sizeof(Road<nLayers - 2>) / MB);`
`573`		`- allocMemAsync(reinterpret_cast<void*>(&mRoadsDevice), mRoads.size() sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), false);`
	`570`	`+ allocMemAsync(reinterpret_cast<void*>(&mRoadsDevice), mRoads.size() sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), getExtAllocator());`
`574`	`571`	`checkGPUError(cudaHostRegister(mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaHostRegisterPortable));`
`575`	`572`	`checkGPUError(cudaMemcpyAsync(mRoadsDevice, mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`576`	`573`	`}`
`@@ -579,7 +576,7 @@ template <int nLayers>`
`579`	`576`	`void TimeFrameGPU<nLayers>::loadTrackSeedsDevice(std::vector<CellSeed>& seeds)`
`580`	`577`	`{`
`581`	`578`	`LOGP(debug, "gpu-transfer: loading {} track seeds, for {} MB.", seeds.size(), seeds.size() * sizeof(CellSeed) / MB);`
`582`		`- allocMemAsync(reinterpret_cast<void*>(&mTrackSeedsDevice), seeds.size() sizeof(CellSeed), &(mGpuStreams[0]), false);`
	`579`	`+ allocMemAsync(reinterpret_cast<void*>(&mTrackSeedsDevice), seeds.size() sizeof(CellSeed), &(mGpuStreams[0]), getExtAllocator());`
`583`	`580`	`checkGPUError(cudaHostRegister(seeds.data(), seeds.size() * sizeof(CellSeed), cudaHostRegisterPortable));`
`584`	`581`	`checkGPUError(cudaMemcpyAsync(mTrackSeedsDevice, seeds.data(), seeds.size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`585`	`582`	`}`
`@@ -589,8 +586,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice(const std::vector<CellSeed>&`
`589`	`586`	`{`
`590`	`587`	`mTrackITSExt.clear();`
`591`	`588`	`mTrackITSExt.resize(seeds.size());`
`592`		`- LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);`
`593`		`- allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), seeds.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);`
	`589`	`+ LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);`
	`590`	`+ allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), seeds.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());`
`594`	`591`	`checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));`
`595`	`592`	`checkGPUError(cudaHostRegister(mTrackITSExt.data(), seeds.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));`
`596`	`593`	`}`
`@@ -600,8 +597,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice()`
`600`	`597`	`{`
`601`	`598`	`mTrackITSExt.clear();`
`602`	`599`	`mTrackITSExt.resize(mRoads.size());`
`603`		`- LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);`
`604`		`- allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), mRoads.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);`
	`600`	`+ LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);`
	`601`	`+ allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), mRoads.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());`
`605`	`602`	`checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, mRoads.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));`
`606`	`603`	`checkGPUError(cudaHostRegister(mTrackITSExt.data(), mRoads.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));`
`607`	`604`	`}`
Original file line number	Diff line number	Diff line change
`@@ -870,9 +870,7 @@ template <int nLayers>`
`870`	`870`	`void TrackerTraitsGPU<nLayers>::initialiseTimeFrame(const int iteration)`
`871`	`871`	`{`
`872`	`872`	`mTimeFrameGPU->initialiseHybrid(iteration, mTrkParams[iteration], nLayers);`
`873`		`- mTimeFrameGPU->loadClustersDevice();`
`874`		`- mTimeFrameGPU->loadUnsortedClustersDevice();`
`875`		`- mTimeFrameGPU->loadTrackingFrameInfoDevice();`
	`873`	`+ mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration);`
`876`	`874`	`}`
`877`	`875`
`878`	`876`	`template <int nLayers>`
`@@ -1165,14 +1163,12 @@ template <int nLayers>`
`1165`	`1163`	`void TrackerTraitsGPU<nLayers>::computeTrackletsHybrid(const int iteration)`
`1166`	`1164`	`{`
`1167`	`1165`	`TrackerTraits::computeLayerTracklets(iteration);`
`1168`		`- mTimeFrameGPU->loadTrackletsDevice();`
`1169`	`1166`	`}`
`1170`	`1167`
`1171`	`1168`	`template <int nLayers>`
`1172`	`1169`	`void TrackerTraitsGPU<nLayers>::computeCellsHybrid(const int iteration)`
`1173`	`1170`	`{`
`1174`	`1171`	`TrackerTraits::computeLayerCells(iteration);`
`1175`		`- mTimeFrameGPU->loadCellsDevice();`
`1176`	`1172`	`};`
`1177`	`1173`
`1178`	`1174`	`template <int nLayers>`
`@@ -1203,6 +1199,10 @@ void TrackerTraitsGPU<nLayers>::findRoadsHybrid(const int iteration)`
`1203`	`1199`	`}`
`1204`	`1200`	`trackSeeds.insert(trackSeeds.end(), updatedCellSeed.begin(), updatedCellSeed.end());`
`1205`	`1201`	`}`
	`1202`	`+ if (!trackSeeds.size()) {`
	`1203`	`+ LOGP(info, "No track seeds found, skipping track finding");`
	`1204`	`+ continue;`
	`1205`	`+ }`
`1206`	`1206`	`mTimeFrameGPU->createTrackITSExtDevice(trackSeeds);`
`1207`	`1207`	`mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);`
`1208`	`1208`
`@@ -1265,12 +1265,15 @@ void TrackerTraitsGPU<nLayers>::findRoadsHybrid(const int iteration)`
`1265`	`1265`	`mTimeFrame->getTracks(std::min(rofs[0], rofs[1])).emplace_back(track);`
`1266`	`1266`	`}`
`1267`	`1267`	`}`
	`1268`	`+ if (iteration == 2) {`
	`1269`	`+ mTimeFrameGPU->unregisterHostMemory(0); // FIXME this needs to work also with sync`
	`1270`	`+ }`
`1268`	`1271`	`};`
`1269`	`1272`
`1270`	`1273`	`template <int nLayers>`
`1271`	`1274`	`void TrackerTraitsGPU<nLayers>::findTracksHybrid(const int iteration)`
`1272`	`1275`	`{`
`1273`		`- // LOGP(info, "propagator device pointer: {}", (void*)mTimeFrameGPU->getDevicePropagator());`
	`1276`	`+ LOGP(info, "propagator device pointer: {}", (void*)mTimeFrameGPU->getDevicePropagator());`
`1274`	`1277`	`mTimeFrameGPU->createTrackITSExtDevice();`
`1275`	`1278`	`gpu::fitTracksKernel<<<20, 512>>>(mTimeFrameGPU->getDeviceArrayClusters(), // Cluster** foundClusters,`
`1276`	`1279`	`mTimeFrameGPU->getDeviceArrayUnsortedClusters(), // Cluster** foundUnsortedClusters,`
`@@ -1287,7 +1290,6 @@ void TrackerTraitsGPU<nLayers>::findTracksHybrid(const int iteration)`
`1287`	`1290`	`mTrkParams[0].MaxChi2NDF, // float maxChi2NDF,`
`1288`	`1291`	`mTimeFrameGPU->getDevicePropagator()); // const o2::base::Propagator* propagator`
`1289`	`1292`	`mTimeFrameGPU->downloadTrackITSExtDevice();`
`1290`		`- discardResult(cudaDeviceSynchronize());`
`1291`	`1293`	`auto& tracks = mTimeFrameGPU->getTrackITSExt();`
`1292`	`1294`	`std::sort(tracks.begin(), tracks.end(),`
`1293`	`1295`	`[](TrackITSExt& track1, TrackITSExt& track2) { return track1.isBetter(track2, 1.e6f); });`
Original file line number	Diff line number	Diff line change
`@@ -106,6 +106,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,`
`106`	`106`	`double total{0.};`
`107`	`107`	`mTraits->UpdateTrackingParameters(mTrkParams);`
`108`	`108`	`for (int iteration = 0; iteration < (int)mTrkParams.size(); ++iteration) {`
	`109`	`+ LOGP(info, "Iteration {}", iteration);`
`109`	`110`	`total += evaluateTask(&Tracker::initialiseTimeFrameHybrid, "Hybrid Timeframe initialisation", logger, iteration);`
`110`	`111`	`total += evaluateTask(&Tracker::computeTrackletsHybrid, "Hybrid Tracklet finding", logger, iteration);`
`111`	`112`	`logger(fmt::format("\t- Number of tracklets: {}", mTraits->getTFNumberOfTracklets()));`
`@@ -115,7 +116,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,`
`115`	`116`	`}`
`116`	`117`	`float trackletsPerCluster = mTraits->getTFNumberOfClusters() > 0 ? float(mTraits->getTFNumberOfTracklets()) / mTraits->getTFNumberOfClusters() : 0.f;`
`117`	`118`	`if (trackletsPerCluster > mTrkParams[iteration].TrackletsPerClusterLimit) {`
`118`		`- error(fmt::format("Too many tracklets per cluster ({}), check the detector status and/or the selections.", trackletsPerCluster));`
	`119`	`+ error(fmt::format("Too many tracklets per cluster ({}), check the detector status and/or the selections. Current limit is {}", trackletsPerCluster, mTrkParams[iteration].TrackletsPerClusterLimit));`
`119`	`120`	`break;`
`120`	`121`	`}`
`121`	`122`
`@@ -127,7 +128,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,`
`127`	`128`	`}`
`128`	`129`	`float cellsPerCluster = mTraits->getTFNumberOfClusters() > 0 ? float(mTraits->getTFNumberOfCells()) / mTraits->getTFNumberOfClusters() : 0.f;`
`129`	`130`	`if (cellsPerCluster > mTrkParams[iteration].CellsPerClusterLimit) {`
`130`		`- error(fmt::format("Too many cells per cluster ({}), check the detector status and/or the selections.", cellsPerCluster));`
	`131`	`+ error(fmt::format("Too many cells per cluster ({}), check the detector status and/or the selections. Current limit is {}", cellsPerCluster, mTrkParams[iteration].CellsPerClusterLimit));`
`131`	`132`	`break;`
`132`	`133`	`}`
`133`	`134`	`total += evaluateTask(&Tracker::findCellsNeighboursHybrid, "Hybrid Neighbour finding", logger, iteration);`
Original file line number	Diff line number	Diff line change
`@@ -676,7 +676,7 @@ void* GPUReconstruction::AllocateUnmanagedMemory(size_t size, int type)`
`676`	`676`	`char* retVal;`
`677`	`677`	`GPUProcessor::computePointerWithAlignment(pool, retVal, size);`
`678`	`678`	`if (pool > poolend) {`
`679`		`- GPUError("Insufficient unmanaged memory: missing %lu", (size_t)((char)pool - (char)poolend));`
	`679`	`+ GPUError("Insufficient unmanaged memory: missing %lu bytes", (size_t)((char)pool - (char)poolend));`
`680`	`680`	`throw std::bad_alloc();`
`681`	`681`	`}`
`682`	`682`	`UpdateMaxMemoryUsed();`
Original file line number	Diff line number	Diff line change
`@@ -90,7 +90,7 @@ o2::its::TimeFrame* GPUChainITS::GetITSTimeframe()`
`90`	`90`	`mFrameworkAllocator.reset(new o2::its::GPUFrameworkExternalAllocator);`
`91`	`91`	`mFrameworkAllocator->setReconstructionFramework(rec());`
`92`	`92`	`mITSTimeFrame->setExternalAllocator(mFrameworkAllocator.get());`
`93`		`- LOGP(debug, "GPUChainITS is giving me ps: {} prop: {}", (void)processorsShadow(), (void)processorsShadow()->calibObjects.o2Propagator);`
	`93`	`+ LOGP(info, "GPUChainITS is giving me ps: {} prop: {} allocator: {}", (void)processorsShadow(), (void)processorsShadow()->calibObjects.o2Propagator, (void*)mFrameworkAllocator.get());`
`94`	`94`	`mITSTimeFrame->setDevicePropagator(processorsShadow()->calibObjects.o2Propagator);`
`95`	`95`	`}`
`96`	`96`	`#endif`