Merge branch 'main' of https://github.com/acts-project/traccc into da…

…ta-plot

.gitignore

@@ -30,3 +30,6 @@
 *.exe
 *.out
 *.app
+
+# VSCode build output.
+out/

.gitlab-ci.yml

@@ -1,3 +1,9 @@
+# TRACCC library, part of the ACTS project (R&D line)
+#
+# (c) 2022-2023 CERN for the benefit of the ACTS project
+#
+# Mozilla Public License Version 2.0
+
 stages:
   - build
   - test
@@ -8,23 +14,19 @@ build_cuda:
   stage: build
   image: ghcr.io/acts-project/ubuntu2004_cuda:v30
   artifacts:
-    paths: 
+    paths:
       - build
   script:
       - git clone $CLONE_URL src
       - git -C src checkout $HEAD_SHA
-      - >
-        cmake -S src -B build -GNinja
-        -DCMAKE_BUILD_TYPE=Release
-        -DTRACCC_BUILD_TESTING=ON
-        -DTRACCC_BUILD_CUDA=ON
+      - cmake --preset cuda -DCMAKE_BUILD_TYPE=Release -S src -B build -G Ninja
       - cmake --build build --parallel 2
 
 test_cuda:
   stage: test
   tags: [docker-gpu-nvidia]
   image: ghcr.io/acts-project/ubuntu2004_cuda:v30
-  needs: 
+  needs:
     - build_cuda
   script:
     - git clone $CLONE_URL src
@@ -35,4 +37,3 @@ test_cuda:
     - cd build
     - nvidia-smi
     - ctest --output-on-failure -E "^SeedingValidation/CompareWithActsSeedingTests.*"
-

.vscode/settings.json

@@ -0,0 +1,7 @@
+{
+   "files.associations": {
+      "*.hip": "cpp",
+      "*.sycl": "cpp",
+      "*.ipp": "cpp"
+   }
+}

CMakeLists.txt

@@ -14,6 +14,7 @@ set( CMAKE_CXX_EXTENSIONS FALSE CACHE BOOL "Disable (host) C++ extensions" )
 set( CMAKE_CUDA_STANDARD 17 CACHE STRING "The (CUDA) C++ standard to use" )
 set( CMAKE_CUDA_EXTENSIONS FALSE CACHE BOOL "Disable (CUDA) C++ extensions" )
 set( CMAKE_SYCL_STANDARD 17 CACHE STRING "The (SYCL) C++ standard to use" )
+set( CMAKE_HIP_STANDARD 17 CACHE STRING "The (HIP) C++ standard to use" )
 
 # Flag controlling whether warnings should make the build fail.
 option( TRACCC_FAIL_ON_WARNINGS

CMakePresets.json

@@ -0,0 +1,83 @@
+{
+   "version" : 3,
+   "configurePresets": [
+      {
+         "name" : "base",
+         "displayName" : "Base Configuration",
+         "warnings": {
+            "deprecated": true
+         },
+         "cacheVariables": {
+            "CMAKE_BUILD_TYPE"        : "RelWithDebInfo",
+            "TRACCC_BUILD_TESTING"    : "TRUE",
+            "TRACCC_BUILD_EXAMPLES"   : "TRUE",
+            "TRACCC_FAIL_ON_WARNINGS" : "TRUE",
+            "TRACCC_USE_ROOT"         : "FALSE"
+         }
+      },
+      {
+         "name" : "cuda",
+         "displayName" : "CUDA Code Development",
+         "inherits": ["base"],
+         "cacheVariables": {
+            "TRACCC_BUILD_CUDA" : "TRUE",
+            "VECMEM_BUILD_CUDA_LIBRARY" : "TRUE"
+         }
+      },
+      {
+         "name" : "sycl",
+         "displayName" : "SYCL Code Development",
+         "inherits": ["base"],
+         "cacheVariables": {
+            "TRACCC_BUILD_SYCL" : "TRUE",
+            "VECMEM_BUILD_SYCL_LIBRARY" : "TRUE"
+         }
+      },
+      {
+         "name" : "alpaka",
+         "displayName" : "Alpaka Code Development",
+         "inherits": ["cuda"],
+         "cacheVariables": {
+            "TRACCC_BUILD_ALPAKA" : "TRUE"
+         }
+      },
+      {
+         "name" : "kokkos",
+         "displayName" : "Kokkos Code Development",
+         "inherits": ["cuda"],
+         "cacheVariables": {
+            "TRACCC_BUILD_KOKKOS" : "TRUE"
+         }
+      },
+      {
+         "name" : "root",
+         "displayName" : "ROOT Enabled Code Development",
+         "inherits": ["base"],
+         "cacheVariables": {
+            "TRACCC_USE_ROOT" : "TRUE",
+            "TRACCC_USE_SYSTEM_TBB" : "TRUE"
+         }
+      },
+      {
+         "name" : "array",
+         "displayName" : "ARRAY Backended Code Development",
+         "inherits": ["base"],
+         "cacheVariables": {
+            "TRACCC_ALGEBRA_PLUGINS" : "ARRAY"
+         }
+      },
+      {
+         "name" : "eigen",
+         "displayName" : "EIGEN Backended Code Development",
+         "inherits": ["base"],
+         "cacheVariables": {
+            "TRACCC_ALGEBRA_PLUGINS" : "EIGEN"
+         }
+      },
+      {
+         "name" : "default",
+         "displayName": "Default Build Configuration",
+         "inherits": ["root", "array"]
+      }
+   ]
+}

core/CMakeLists.txt

@@ -87,7 +87,10 @@ traccc_add_library( traccc_core core TYPE SHARED
   "include/traccc/seeding/seed_finding.hpp"
   "src/seeding/seed_finding.cpp"
   "include/traccc/seeding/spacepoint_binning.hpp"
-  "src/seeding/spacepoint_binning.cpp" )
+  "src/seeding/spacepoint_binning.cpp"
+  # Ambiguity resolution
+  "include/traccc/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp"
+  "src/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.cpp" )
 target_link_libraries( traccc_core
   PUBLIC Eigen3::Eigen vecmem::core detray::core traccc::Thrust
          traccc::algebra )

core/include/traccc/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp

@@ -0,0 +1,155 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// System include
+#include <algorithm>
+#include <initializer_list>
+#include <iostream>
+#include <limits>
+#include <map>
+#include <set>
+#include <unordered_map>
+#include <vector>
+
+// VecMem include(s).
+#include <vecmem/containers/vector.hpp>
+
+// Project include(s).
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/track_candidate.hpp"
+#include "traccc/edm/track_state.hpp"
+#include "traccc/utils/algorithm.hpp"
+
+// Greedy ambiguity resolution adapted from ACTS code
+
+namespace traccc {
+
+/// Evicts tracks that seem to be duplicates or fakes. This algorithm takes a
+/// greedy approach in the sense that it will remove the track which looks "most
+/// duplicate/fake" first and continues the same process with the rest. That
+/// process continues until the final state conditions are met.
+///
+/// The implementation works as follows:
+///  1) Calculate shared hits per track.
+///  2) If the maximum shared hits criteria is met, we are done.
+///     This is the configurable amount of shared hits we are ok with
+///     in our experiment.
+///  3) Else, remove the track with the highest relative shared hits (i.e.
+///     shared hits / hits).
+///  4) Back to square 1.
+class greedy_ambiguity_resolution_algorithm
+    : public algorithm<track_state_container_types::host(
+          const typename track_state_container_types::host&)> {
+
+    public:
+    struct config_t {
+
+        config_t(){};
+
+        /// Maximum amount of shared hits per track. One (1) means "no shared
+        /// hit allowed".
+        std::uint32_t maximum_shared_hits = 1;
+
+        /// Maximum number of iterations.
+        std::uint32_t maximum_iterations = 1000000;
+
+        /// Minimum number of measurement to form a track.
+        std::size_t n_measurements_min = 3;
+
+        // True if obvious errors should be checked after the completion
+        // of the algorithm.
+        bool check_obvious_errs = true;
+
+        bool verbose_info = true;
+        bool verbose_error = true;
+        bool verbose_flood = false;
+    };
+
+    struct state_t {
+        std::size_t number_of_tracks{};
+
+        /// For this whole comment section, track_index refers to the index of a
+        /// track in the initial input container.
+        ///
+        /// There is no (track_id) in this algorithm, only (track_index).
+
+        /// Associates each track_index with the track's chi2 value
+        std::vector<traccc::scalar> track_chi2;
+
+        /// Associates each track_index to the track's (measurement_id)s list
+        std::vector<std::vector<std::size_t>> measurements_per_track;
+
+        /// Associates each measurement_id to a set of (track_index)es sharing
+        /// it
+        std::unordered_map<std::size_t, std::set<std::size_t>>
+            tracks_per_measurement;
+
+        /// Associates each track_index to its number of shared measurements
+        /// (among other tracks)
+        std::vector<std::size_t> shared_measurements_per_track;
+
+        /// Keeps the selected tracks indexes that have not (yet) been removed
+        /// by the algorithm
+        std::set<std::size_t> selected_tracks;
+    };
+
+    /// Constructor for the greedy ambiguity resolution algorithm
+    ///
+    /// @param cfg  Configuration object
+    // greedy_ambiguity_resolution_algorithm(const config_type& cfg) :
+    // _config(cfg) {}
+    greedy_ambiguity_resolution_algorithm(const config_t cfg = {})
+        : _config{cfg} {}
+
+    /// Run the algorithm
+    ///
+    /// @param track_states the container of the fitted track parameters
+    /// @return the container without ambiguous tracks
+    track_state_container_types::host operator()(
+        const typename track_state_container_types::host& track_states)
+        const override;
+
+    private:
+    /// Computes the initial state for the input data. This function accumulates
+    /// information that will later be used to accelerate the ambiguity
+    /// resolution.
+    ///
+    /// @param track_states The input track container (output of the fitting
+    /// algorithm).
+    /// @param state An empty state object which is expected to be default
+    /// constructed.
+    void compute_initial_state(
+        const typename track_state_container_types::host& track_states,
+        state_t& state) const;
+
+    /// Updates the state iteratively by evicting one track after the other
+    /// until the final state conditions are met.
+    ///
+    /// @param state A state object that was previously filled by the
+    /// initialization.
+    void resolve(state_t& state) const;
+
+    /// Check for obvious errors returned by the algorithm:
+    /// - Returned tracks should be independent of each other: they should share
+    ///   a maximum of (_config.maximum_shared_hits - 1) hits per track.
+    /// - Each removed track should share at least (_config.maximum_shared_hits)
+    ///   with another initial track.
+    ///
+    /// @param initial_track_states The input track container, as given to
+    /// compute_initial_state.
+    /// @param final_state The state object after the resolve method has been
+    /// called.
+    bool check_obvious_errors(
+        const typename track_state_container_types::host& initial_track_states,
+        state_t& final_state) const;
+
+    config_t _config;
+};
+
+}  // namespace traccc

core/include/traccc/definitions/qualifiers.hpp

@@ -8,25 +8,25 @@
 
 #pragma once
 
-#if defined(__CUDACC__)
+#if defined(__CUDACC__) || defined(__HIP__)
 #define TRACCC_DEVICE __device__
 #else
 #define TRACCC_DEVICE
 #endif
 
-#if defined(__CUDACC__)
+#if defined(__CUDACC__) || defined(__HIP__)
 #define TRACCC_HOST __host__
 #else
 #define TRACCC_HOST
 #endif
 
-#if defined(__CUDACC__)
+#if defined(__CUDACC__) || defined(__HIP__)
 #define TRACCC_HOST_DEVICE __host__ __device__
 #else
 #define TRACCC_HOST_DEVICE
 #endif
 
-#if defined(__CUDACC__)
+#if defined(__CUDACC__) || defined(__HIP__)
 #define TRACCC_ALIGN(x) __align__(x)
 #else
 #define TRACCC_ALIGN(x) alignas(x)

core/include/traccc/edm/measurement.hpp

@@ -39,6 +39,9 @@ struct measurement {
     /// Geometry ID
     detray::geometry::barcode surface_link;
 
+    // Unique measurement ID
+    std::size_t measurement_id = 0;
+
     /// Link to Module vector index
     using link_type = cell_module_collection_types::view::size_type;
     link_type module_link = 0;

core/include/traccc/finding/candidate_link.hpp

@@ -30,6 +30,9 @@ struct candidate_link {
 
     // Measurement index
     unsigned int meas_idx;
+
+    // Index to the initial seed
+    unsigned int seed_idx;
 };
 
-}  // namespace traccc
+}  // namespace traccc