Fix bugs in the greedy resolver (#1083)

* Fix several bugs in the greedy resolver which includes wrong size allocation, deadlock and unsafe usage of a local varaiable

* Change the test name to register in CI run

* Add CUDA keywords to the test suites

* Rename the variables and use local memory for n_accepted_prev

* Use at instead of [] operator

---------

Co-authored-by: Stephen Nicholas Swatman <[email protected]>

Author: beomki-yeo

device/cuda/src/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.cu

@@ -233,7 +233,7 @@ greedy_ambiguity_resolution_algorithm::operator()(
 
     // Unique measurement ids
     vecmem::data::vector_buffer<measurement_id_type>
-        meas_id_to_unique_id_buffer{max_meas.measurement_id, m_mr.main};
+        meas_id_to_unique_id_buffer{max_meas.measurement_id + 1, m_mr.main};
 
     // Make meas_id to meas vector
     {

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

@@ -36,8 +36,9 @@ __device__ void count_tracks(int tid, int* sh_n_meas, int n_tracks,
                 offset = sh_n_meas[count];
                 add = stride * 2;
             }
-            __syncthreads();
         }
+
+        __syncthreads();
     }
 
     if (tid == 0) {

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

@@ -38,6 +38,9 @@ __global__ void remove_tracks(device::remove_tracks_payload payload) {
 
     auto threadIndex = threadIdx.x;
 
+    bool is_valid_thread = false;
+    bool is_duplicate = true;
+
     shared_tids[threadIndex] = std::numeric_limits<unsigned int>::max();
 
     vecmem::device_vector<const unsigned int> sorted_ids(
@@ -62,112 +65,128 @@ __global__ void remove_tracks(device::remove_tracks_payload payload) {
         payload.meas_to_remove_view);
     vecmem::device_vector<unsigned int> threads(payload.threads_view);
 
-    auto n_accepted_prev = (*payload.n_accepted);
+    const unsigned n_accepted_prev = *(payload.n_accepted);
+
+    __syncthreads();
+
     if (threadIndex == 0) {
         (*payload.n_accepted) -= *(payload.n_removable_tracks);
     }
 
     if (threadIndex < *(payload.n_meas_to_remove)) {
         sh_meas_ids[threadIndex] = meas_to_remove[threadIndex];
         sh_threads[threadIndex] = threads[threadIndex];
-    } else {
-        return;
+        is_valid_thread = true;
     }
 
-    const auto id = sh_meas_ids[threadIndex];
+    __syncthreads();
+
+    if (is_valid_thread) {
+
+        const auto id = sh_meas_ids[threadIndex];
+        is_duplicate = false;
 
-    bool is_duplicate = false;
-    for (unsigned int i = 0; i < threadIndex; ++i) {
-        if (sh_meas_ids[i] == id) {
-            is_duplicate = true;
-            break;
+        for (unsigned int i = 0; i < threadIndex; ++i) {
+            if (sh_meas_ids[i] == id) {
+                is_duplicate = true;
+                break;
+            }
         }
     }
-    if (is_duplicate) {
-        return;
-    }
 
-    const auto unique_meas_idx = meas_id_to_unique_id.at(id);
+    bool active = false;
 
-    // If there is only one track associated with measurement, the
-    // number of shared measurement can be reduced by one
-    const auto& tracks = tracks_per_measurement[unique_meas_idx];
-    auto track_status = track_status_per_measurement[unique_meas_idx];
+    if (!is_duplicate && is_valid_thread) {
 
-    auto trk_id = sorted_ids[n_accepted_prev - 1 - sh_threads[threadIndex]];
+        const auto id = sh_meas_ids[threadIndex];
+        const auto unique_meas_idx = meas_id_to_unique_id.at(id);
 
-    unsigned int worst_idx =
-        thrust::find(thrust::seq, tracks.begin(), tracks.end(), trk_id) -
-        tracks.begin();
+        // If there is only one track associated with measurement, the
+        // number of shared measurement can be reduced by one
+        const auto& tracks = tracks_per_measurement[unique_meas_idx];
+        auto track_status = track_status_per_measurement[unique_meas_idx];
 
-    track_status[worst_idx] = 0;
+        auto trk_id =
+            sorted_ids.at(n_accepted_prev - 1 - sh_threads[threadIndex]);
 
-    int n_sharing_tracks = 1;
-    for (unsigned int i = threadIndex + 1; i < *(payload.n_meas_to_remove);
-         ++i) {
+        unsigned int worst_idx =
+            thrust::find(thrust::seq, tracks.begin(), tracks.end(), trk_id) -
+            tracks.begin();
 
-        if (sh_meas_ids[i] == id && sh_threads[i] != sh_threads[i - 1]) {
-            n_sharing_tracks++;
+        track_status[worst_idx] = 0;
 
-            trk_id = sorted_ids[n_accepted_prev - 1 - sh_threads[i]];
+        int n_sharing_tracks = 1;
+        for (unsigned int i = threadIndex + 1; i < *(payload.n_meas_to_remove);
+             ++i) {
 
-            worst_idx = thrust::find(thrust::seq, tracks.begin(), tracks.end(),
-                                     trk_id) -
-                        tracks.begin();
+            if (sh_meas_ids[i] == id && sh_threads[i] != sh_threads[i - 1]) {
+                n_sharing_tracks++;
 
-            track_status[worst_idx] = 0;
+                trk_id = sorted_ids[n_accepted_prev - 1 - sh_threads[i]];
 
-        } else if (sh_meas_ids[i] != id) {
-            break;
-        }
-    }
+                worst_idx = thrust::find(thrust::seq, tracks.begin(),
+                                         tracks.end(), trk_id) -
+                            tracks.begin();
 
-    vecmem::device_atomic_ref<unsigned int> n_accepted_per_meas(
-        n_accepted_tracks_per_measurement.at(
-            static_cast<unsigned int>(unique_meas_idx)));
-    const unsigned int N_A = n_accepted_per_meas.fetch_sub(n_sharing_tracks);
+                track_status[worst_idx] = 0;
 
-    if (N_A != 1 + n_sharing_tracks) {
-        return;
-    }
+            } else if (sh_meas_ids[i] != id) {
+                break;
+            }
+        }
 
-    const unsigned int alive_idx =
-        thrust::find(thrust::seq, track_status.begin(), track_status.end(), 1) -
-        track_status.begin();
+        vecmem::device_atomic_ref<unsigned int> n_accepted_per_meas(
+            n_accepted_tracks_per_measurement.at(
+                static_cast<unsigned int>(unique_meas_idx)));
+        const unsigned int N_A =
+            n_accepted_per_meas.fetch_sub(n_sharing_tracks);
 
-    shared_tids[threadIndex] = static_cast<unsigned int>(tracks[alive_idx]);
+        if (N_A == 1 + n_sharing_tracks) {
+            active = true;
+            const unsigned int alive_idx =
+                thrust::find(thrust::seq, track_status.begin(),
+                             track_status.end(), 1) -
+                track_status.begin();
 
-    auto tid = shared_tids[threadIndex];
+            shared_tids[threadIndex] =
+                static_cast<unsigned int>(tracks[alive_idx]);
 
-    const auto m_count = static_cast<unsigned int>(thrust::count(
-        thrust::seq, meas_ids[tid].begin(), meas_ids[tid].end(), id));
+            auto tid = shared_tids[threadIndex];
 
-    const unsigned int N_S =
-        vecmem::device_atomic_ref<unsigned int>(n_shared.at(tid))
-            .fetch_sub(m_count);
+            const auto m_count = static_cast<unsigned int>(thrust::count(
+                thrust::seq, meas_ids[tid].begin(), meas_ids[tid].end(), id));
+
+            const unsigned int N_S =
+                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) {
-            already_pushed = true;
-            break;
+    if (active) {
+        auto tid = shared_tids[threadIndex];
+        bool already_pushed = false;
+        for (unsigned int i = 0; i < threadIndex; ++i) {
+            if (shared_tids[i] == tid) {
+                already_pushed = true;
+                break;
+            }
         }
-    }
-    if (!already_pushed) {
+        if (!already_pushed) {
 
-        // Write updated track IDs
-        vecmem::device_atomic_ref<unsigned int> num_updated_tracks(
-            *(payload.n_updated_tracks));
+            // Write updated track IDs
+            vecmem::device_atomic_ref<unsigned int> num_updated_tracks(
+                *(payload.n_updated_tracks));
 
-        const unsigned int pos = num_updated_tracks.fetch_add(1);
+            const unsigned int pos = num_updated_tracks.fetch_add(1);
 
-        updated_tracks[pos] = tid;
-        is_updated[tid] = 1;
+            updated_tracks[pos] = tid;
+            is_updated[tid] = 1;
 
-        rel_shared.at(tid) = static_cast<traccc::scalar>(n_shared.at(tid)) /
-                             static_cast<traccc::scalar>(n_meas.at(tid));
+            rel_shared.at(tid) = static_cast<traccc::scalar>(n_shared.at(tid)) /
+                                 static_cast<traccc::scalar>(n_meas.at(tid));
+        }
     }
 }