Synchronize the threads to have n_shared correctly before calculation of the relative number of shared measurements (#1014)

Author: beomki-yeo

device/cuda/src/ambiguity_resolution/kernels/remove_tracks.cu

@@ -140,8 +140,6 @@ __global__ void remove_tracks(device::remove_tracks_payload payload) {
 
     shared_tids[threadIndex] = static_cast<unsigned int>(tracks[alive_idx]);
 
-    __syncthreads();
-
     auto tid = shared_tids[threadIndex];
 
     const auto m_count = static_cast<unsigned int>(thrust::count(
@@ -151,6 +149,8 @@ __global__ void remove_tracks(device::remove_tracks_payload payload) {
         vecmem::device_atomic_ref<unsigned int>(n_shared.at(tid))
             .fetch_sub(m_count);
 
+    __syncthreads();
+
     bool already_pushed = false;
     for (unsigned int i = 0; i < threadIndex; ++i) {
         if (shared_tids[i] == tid) {

tests/cuda/test_ambiguity_resolution.cpp

@@ -942,6 +942,15 @@ INSTANTIATE_TEST_SUITE_P(
                                       std::array<std::size_t, 2u>{3u, 10u},
                                       100u, false)));
 
+INSTANTIATE_TEST_SUITE_P(
+    Long, GreedyResolutionCompareToCPU,
+    ::testing::Values(std::make_tuple(3u, 10000u,
+                                      std::array<std::size_t, 2u>{3u, 500u},
+                                      10000u, true),
+                      std::make_tuple(3u, 10000u,
+                                      std::array<std::size_t, 2u>{3u, 500u},
+                                      10000u, false)));
+
 INSTANTIATE_TEST_SUITE_P(
     Simple, GreedyResolutionCompareToCPU,
     ::testing::Values(