Make ITS GPU running on PbPb in async

mconcas · mconcas · commit 8f2cdfcee71d · 2023-12-27T18:28:45.000+01:00
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
@@ -345,7 +345,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()));
   }
 }
@@ -376,17 +376,11 @@ 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()));
+  mAllocator->deallocate();
 }
 
 template <int nLayers>
@@ -525,17 +519,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()));
 }
 
@@ -573,7 +571,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()));
 }
@@ -582,7 +580,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()));
 }
@@ -592,8 +590,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));
 }
@@ -603,8 +601,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,6 +1265,9 @@ 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>
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/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/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);
@@ -319,6 +324,8 @@ void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
     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.;

Original file line number	Diff line number	Diff line change
`@@ -345,7 +345,7 @@ void TimeFrameGPU<nLayers>::allocMemAsync(void** ptr, size_t size, Stream* strPt`
`345`	`345`	`if (extAllocator) {`
`346`	`346`	`*ptr = mAllocator->allocate(size);`
`347`	`347`	`} else {`
`348`		`- LOGP(debug, "Calling default CUDA allocator");`
	`348`	`+ LOGP(info, "Calling default CUDA allocator");`
`349`	`349`	`checkGPUError(cudaMallocAsync(reinterpret_cast<void**>(ptr), size, strPtr->get()));`
`350`	`350`	`}`
`351`	`351`	`}`
`@@ -376,17 +376,11 @@ void TimeFrameGPU<nLayers>::registerHostMemory(const int maxLayers)`
`376`	`376`	`template <int nLayers>`
`377`	`377`	`void TimeFrameGPU<nLayers>::unregisterHostMemory(const int maxLayers)`
`378`	`378`	`{`
`379`		`- if (!mHostRegistered) {`
`380`		`- return;`
`381`		`- }`
`382`		`- for (auto iLayer{0}; iLayer < maxLayers; ++iLayer) {`
`383`		`- checkGPUError(cudaHostUnregister(mClusters[iLayer].data()));`
`384`		`- checkGPUError(cudaHostUnregister(mNClustersPerROF[iLayer].data()));`
`385`		`- checkGPUError(cudaHostUnregister(mIndexTables[iLayer].data()));`
	`379`	`+ for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {`
	`380`	`+ checkGPUError(cudaHostUnregister(mTrackingFrameInfo[iLayer].data()));`
`386`	`381`	`}`
`387`		`- checkGPUError(cudaHostUnregister(mHostNTracklets.data()));`
`388`		`- checkGPUError(cudaHostUnregister(mHostNCells.data()));`
`389`		`- mHostRegistered = false;`
	`382`	`+ checkGPUError(cudaHostUnregister(mTrackingFrameInfoDevice.data()));`
	`383`	`+ mAllocator->deallocate();`
`390`	`384`	`}`
`391`	`385`
`392`	`386`	`template <int nLayers>`
`@@ -525,17 +519,21 @@ void TimeFrameGPU<nLayers>::loadClustersDevice()`
`525`	`519`	`}`
`526`	`520`
`527`	`521`	`template <int nLayers>`
`528`		`-void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice()`
	`522`	`+void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice(const int iteration)`
`529`	`523`	`{`
`530`	`524`	`for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {`
`531`	`525`	`LOGP(debug, "gpu-transfer: loading {} tfinfo on layer {}, for {} MB.", mTrackingFrameInfo[iLayer].size(), iLayer, mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo) / MB);`
`532`	`526`	`allocMemAsync(reinterpret_cast<void*>(&mTrackingFrameInfoDevice[iLayer]), mTrackingFrameInfo[iLayer].size() sizeof(TrackingFrameInfo), nullptr, getExtAllocator());`
`533`	`527`	`// Register and move data`
`534`		`- checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));`
	`528`	`+ if (!iteration) {`
	`529`	`+ checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));`
	`530`	`+ }`
`535`	`531`	`checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDevice[iLayer], mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`536`	`532`	`}`
`537`	`533`	`allocMemAsync(reinterpret_cast<void*>(&mTrackingFrameInfoDeviceArray), nLayers sizeof(TrackingFrameInfo*), nullptr, getExtAllocator());`
`538`		`- checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));`
	`534`	`+ if (!iteration) {`
	`535`	`+ checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));`
	`536`	`+ }`
`539`	`537`	`checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDeviceArray, mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`540`	`538`	`}`
`541`	`539`
`@@ -573,7 +571,7 @@ template <int nLayers>`
`573`	`571`	`void TimeFrameGPU<nLayers>::loadRoadsDevice()`
`574`	`572`	`{`
`575`	`573`	`LOGP(debug, "gpu-transfer: loading {} roads, for {} MB.", mRoads.size(), mRoads.size() * sizeof(Road<nLayers - 2>) / MB);`
`576`		`- allocMemAsync(reinterpret_cast<void*>(&mRoadsDevice), mRoads.size() sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), false);`
	`574`	`+ allocMemAsync(reinterpret_cast<void*>(&mRoadsDevice), mRoads.size() sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), getExtAllocator());`
`577`	`575`	`checkGPUError(cudaHostRegister(mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaHostRegisterPortable));`
`578`	`576`	`checkGPUError(cudaMemcpyAsync(mRoadsDevice, mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`579`	`577`	`}`
`@@ -582,7 +580,7 @@ template <int nLayers>`
`582`	`580`	`void TimeFrameGPU<nLayers>::loadTrackSeedsDevice(std::vector<CellSeed>& seeds)`
`583`	`581`	`{`
`584`	`582`	`LOGP(debug, "gpu-transfer: loading {} track seeds, for {} MB.", seeds.size(), seeds.size() * sizeof(CellSeed) / MB);`
`585`		`- allocMemAsync(reinterpret_cast<void*>(&mTrackSeedsDevice), seeds.size() sizeof(CellSeed), &(mGpuStreams[0]), false);`
	`583`	`+ allocMemAsync(reinterpret_cast<void*>(&mTrackSeedsDevice), seeds.size() sizeof(CellSeed), &(mGpuStreams[0]), getExtAllocator());`
`586`	`584`	`checkGPUError(cudaHostRegister(seeds.data(), seeds.size() * sizeof(CellSeed), cudaHostRegisterPortable));`
`587`	`585`	`checkGPUError(cudaMemcpyAsync(mTrackSeedsDevice, seeds.data(), seeds.size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0].get()));`
`588`	`586`	`}`
`@@ -592,8 +590,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice(const std::vector<CellSeed>&`
`592`	`590`	`{`
`593`	`591`	`mTrackITSExt.clear();`
`594`	`592`	`mTrackITSExt.resize(seeds.size());`
`595`		`- LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);`
`596`		`- allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), seeds.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);`
	`593`	`+ LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);`
	`594`	`+ allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), seeds.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());`
`597`	`595`	`checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));`
`598`	`596`	`checkGPUError(cudaHostRegister(mTrackITSExt.data(), seeds.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));`
`599`	`597`	`}`
`@@ -603,8 +601,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice()`
`603`	`601`	`{`
`604`	`602`	`mTrackITSExt.clear();`
`605`	`603`	`mTrackITSExt.resize(mRoads.size());`
`606`		`- LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);`
`607`		`- allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), mRoads.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);`
	`604`	`+ LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);`
	`605`	`+ allocMemAsync(reinterpret_cast<void*>(&mTrackITSExtDevice), mRoads.size() sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());`
`608`	`606`	`checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, mRoads.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));`
`609`	`607`	`checkGPUError(cudaHostRegister(mTrackITSExt.data(), mRoads.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));`
`610`	`608`	`}`
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,6 +1265,9 @@ 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>`
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`