fix2

Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>

Author: f3sch

Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h

@@ -157,6 +157,7 @@ class TimeFrameGPU final : public TimeFrame<NLayers>
  private:
   void allocMemAsync(void**, size_t, Stream&, bool, int32_t = o2::gpu::GPUMemoryResource::MEMORY_GPU); // Abstract owned and unowned memory allocations on specific stream
   void allocMem(void**, size_t, bool, int32_t = o2::gpu::GPUMemoryResource::MEMORY_GPU);               // Abstract owned and unowned memory allocations on default stream
+  void releaseTrackingParametersDevice();
 
   // Host-available device buffer sizes
   std::array<int, MaxTransitions> mNTracklets{};

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu

@@ -80,6 +80,25 @@ void TimeFrameGPU<NLayers>::loadTrackingParametersDevice(const TrackingParameter
   GPUChkErrS(cudaMemcpy(mLayerxX0Device, trkParam.LayerxX0.data(), trkParam.LayerxX0.size() * sizeof(float), cudaMemcpyHostToDevice));
 }
 
+template <int NLayers>
+void TimeFrameGPU<NLayers>::releaseTrackingParametersDevice()
+{
+  if (!this->hasFrameworkAllocator()) {
+    if (mLayerRadiiDevice != nullptr) {
+      GPUChkErrS(cudaFree(mLayerRadiiDevice));
+    }
+    if (mMinPtsDevice != nullptr) {
+      GPUChkErrS(cudaFree(mMinPtsDevice));
+    }
+    if (mLayerxX0Device != nullptr) {
+      GPUChkErrS(cudaFree(mLayerxX0Device));
+    }
+  }
+  mLayerRadiiDevice = nullptr;
+  mMinPtsDevice = nullptr;
+  mLayerxX0Device = nullptr;
+}
+
 template <int NLayers>
 void TimeFrameGPU<NLayers>::createUnsortedClustersDeviceArray(const int maxLayers)
 {
@@ -679,6 +698,7 @@ template <int NLayers>
 void TimeFrameGPU<NLayers>::wipe()
 {
   unregisterHostMemory(0);
+  releaseTrackingParametersDevice();
   o2::its::TimeFrame<NLayers>::wipe();
 }
 

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx

@@ -387,8 +387,6 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
     auto& tracks = mTimeFrameGPU->getTrackITSExt();
     this->acceptTracks(iteration, tracks, firstClusters, sharedFirstClusters);
     mTimeFrameGPU->loadUsedClustersDevice();
-    // The next road-building level reads the updated used-cluster mask on the default stream.
-    mTimeFrameGPU->syncStreams(false);
   }
   this->markTracks(iteration, sharedFirstClusters);
   // wipe the artefact memory