From 7f19454c456d1b669ce47480a1b43235b61ceab6 Mon Sep 17 00:00:00 2001
From: Felix Schlepper <felix.schlepper@cern.ch>
Date: Mon, 23 Jun 2025 21:34:03 +0200
Subject: [PATCH] ITS: GPU add needed synchronization

Otherwise we can end up with partial writes to the
foundSeedsTable.
---
 Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
index 2191880374548..8c6367c221583 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
@@ -1217,6 +1217,8 @@ void processNeighboursHandler(const int startLayer,
     maxChi2ClusterAttachment,
     propagator,
     matCorrType);
+  GPUChkErrS(cudaPeekAtLastError());
+  GPUChkErrS(cudaDeviceSynchronize());
 
   int level = startLevel;
   thrust::device_vector<int, gpu::TypedAllocator<int>> lastCellId(allocInt);
@@ -1276,6 +1278,8 @@ void processNeighboursHandler(const int startLayer,
       maxChi2ClusterAttachment,
       propagator,
       matCorrType);
+    GPUChkErrS(cudaPeekAtLastError());
+    GPUChkErrS(cudaDeviceSynchronize());
   }
   thrust::device_vector<CellSeed, gpu::TypedAllocator<CellSeed>> outSeeds(updatedCellSeed.size(), allocCellSeed);
   auto end = thrust::copy_if(updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), gpu::seed_selector(1.e3, maxChi2NDF * ((startLevel + 2) * 2 - 5)));