GPU: Remove GPUTRDTracker DoTracking function, and do everything via GPUChainTracking

davidrohr · davidrohr · commit ee97a625c469 · 2022-01-30T19:16:48.000+01:00
diff --git a/Detectors/TRD/workflow/src/TRDGlobalTrackingSpec.cxx b/Detectors/TRD/workflow/src/TRDGlobalTrackingSpec.cxx
@@ -328,7 +328,7 @@ void TRDGlobalTracking::run(ProcessingContext& pc)
 
   // start the tracking
   //mTracker->DumpTracks();
-  mTracker->DoTracking(mChainTracking);
+  mChainTracking->DoTRDGPUTracking<GPUTRDTrackerKernels::o2Version>(mTracker);
   //mTracker->DumpTracks();
 
   // finished tracking, now collect the output
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx b/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx
@@ -78,15 +78,17 @@ int GPUChainTracking::DoTRDGPUTracking(GPUTRDTracker* externalInstance)
       throw std::runtime_error("Must not provide external instance that does not match template type");
     }
   }
+  Tracker->PrepareTracking(this);
 
   const auto& threadContext = GetThreadContext();
   SetupGPUProcessor(Tracker, false);
-  TrackerShadow->OverrideGPUGeometry(reinterpret_cast<GPUTRDGeometry*>(mFlatObjectsDevice.mCalibObjects.trdGeometry));
+  if (doGPU) {
+    TrackerShadow->OverrideGPUGeometry(reinterpret_cast<GPUTRDGeometry*>(mFlatObjectsDevice.mCalibObjects.trdGeometry));
+    WriteToConstantMemory(RecoStep::TRDTracking, (char*)&processors()->trdTrackerGPU - (char*)processors(), TrackerShadow, sizeof(*TrackerShadow), 0);
+  }
 
-  WriteToConstantMemory(RecoStep::TRDTracking, (char*)&processors()->trdTrackerGPU - (char*)processors(), TrackerShadow, sizeof(*TrackerShadow), 0);
   TransferMemoryResourcesToGPU(RecoStep::TRDTracking, Tracker, 0);
-
-  runKernel<GPUTRDTrackerKernels>(GetGridAuto(0), krnlRunRangeNone, krnlEventNone, nullptr);
+  runKernel<GPUTRDTrackerKernels, I>(GetGridAuto(0), krnlRunRangeNone, krnlEventNone, externalInstance);
   TransferMemoryResourcesToHost(RecoStep::TRDTracking, Tracker, 0);
   SynchronizeStream(0);
 
diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDTracker.cxx b/GPU/GPUTracking/TRDTracking/GPUTRDTracker.cxx
@@ -246,50 +246,6 @@ void GPUTRDTracker_t<TRDTRK, PROP>::PrepareTracking(GPUChainTracking* chainTrack
   if (mGenerateSpacePoints) {
     chainTracking->mIOPtrs.trdSpacePoints = mSpacePoints;
   }
-}
-
-template <class TRDTRK, class PROP>
-void GPUTRDTracker_t<TRDTRK, PROP>::DoTracking(GPUChainTracking* chainTracking)
-{
-  //--------------------------------------------------------------------
-  // Steering function for the tracking
-  //--------------------------------------------------------------------
-
-  PrepareTracking(chainTracking);
-
-  auto timeStart = std::chrono::high_resolution_clock::now();
-
-  if (mRec->GetRecoStepsGPU() & GPUDataTypes::RecoStep::TRDTracking) {
-    chainTracking->DoTRDGPUTracking<0>();
-  } else {
-#ifdef WITH_OPENMP
-#pragma omp parallel for num_threads(mRec->GetProcessingSettings().ompThreads)
-    for (int iTrk = 0; iTrk < mNTracks; ++iTrk) {
-      if (omp_get_num_threads() > mMaxThreads) {
-        GPUError("Number of parallel threads too high, aborting tracking");
-        // break statement not possible in OpenMP for loop
-        iTrk = mNTracks;
-        continue;
-      }
-      DoTrackingThread(iTrk, omp_get_thread_num());
-    }
-#else
-    for (int iTrk = 0; iTrk < mNTracks; ++iTrk) {
-      DoTrackingThread(iTrk);
-    }
-#endif
-  }
-
-  auto duration = std::chrono::high_resolution_clock::now() - timeStart;
-  (void)duration; // suppress warning about unused variable
-  /*
-  std::cout << "--->  -----> -------> ---------> ";
-  std::cout << "Time for event " << mNEvents << ": " << std::chrono::duration_cast<std::chrono::microseconds>(duration).count() << " us ";
-  std::cout << "nTracks: " << mNTracks;
-  std::cout << " nTracklets: " << GetConstantMem()->ioPtrs.nTRDTracklets;
-  std::cout << std::endl;
-  */
-  //DumpTracks();
   mNEvents++;
 }
 
diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h b/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h
@@ -66,7 +66,6 @@ class GPUTRDTracker_t : public GPUProcessor
   void* SetPointersTracks(void* base);
 
   void PrepareTracking(GPUChainTracking* chainTracking);
-  void DoTracking(GPUChainTracking* chainTracking);
   void SetNCandidates(int n);
   void PrintSettings() const;
   bool IsInitialized() const { return mIsInitialized; }
