Add a first version of SYCL track finding.

acts-project · Nov 15, 2024 · 7af0150 · 7af0150
1 parent c4df4c9
commit 7af0150
Show file tree

Hide file tree

Showing 3 changed files with 457 additions and 24 deletions.
diff --git a/core/include/traccc/edm/measurement.hpp b/core/include/traccc/edm/measurement.hpp
@@ -88,14 +88,14 @@ inline bool operator==(const measurement& lhs, const measurement& rhs) {
 /// Comparator based on detray barcode value
 struct measurement_sort_comp {
     TRACCC_HOST_DEVICE
-    bool operator()(const measurement& lhs, const measurement& rhs) {
+    bool operator()(const measurement& lhs, const measurement& rhs) const {
         return lhs.surface_link < rhs.surface_link;
     }
 };
 
 struct measurement_equal_comp {
     TRACCC_HOST_DEVICE
-    bool operator()(const measurement& lhs, const measurement& rhs) {
+    bool operator()(const measurement& lhs, const measurement& rhs) const {
         return lhs.surface_link == rhs.surface_link;
     }
 };

diff --git a/device/sycl/include/traccc/sycl/utils/thread_id.hpp b/device/sycl/include/traccc/sycl/utils/thread_id.hpp
@@ -14,32 +14,32 @@ namespace traccc::sycl {
 struct thread_id1 {
     TRACCC_DEVICE thread_id1(const ::sycl::nd_item<1>& item) : m_item(item) {}
 
-    auto inline TRACCC_DEVICE getLocalThreadId() const {
-        return m_item.get_local_linear_id();
+    unsigned int inline TRACCC_DEVICE getLocalThreadId() const {
+        return static_cast<unsigned int>(m_item.get_local_linear_id());
     }
 
-    auto inline TRACCC_DEVICE getLocalThreadIdX() const {
-        return m_item.get_local_linear_id();
+    unsigned int inline TRACCC_DEVICE getLocalThreadIdX() const {
+        return static_cast<unsigned int>(m_item.get_local_linear_id());
     }
 
-    auto inline TRACCC_DEVICE getGlobalThreadId() const {
-        return m_item.get_global_linear_id();
+    unsigned int inline TRACCC_DEVICE getGlobalThreadId() const {
+        return static_cast<unsigned int>(m_item.get_global_linear_id());
     }
 
-    auto inline TRACCC_DEVICE getGlobalThreadIdX() const {
-        return m_item.get_global_linear_id();
+    unsigned int inline TRACCC_DEVICE getGlobalThreadIdX() const {
+        return static_cast<unsigned int>(m_item.get_global_linear_id());
     }
 
-    auto inline TRACCC_DEVICE getBlockIdX() const {
-        return m_item.get_group_linear_id();
+    unsigned int inline TRACCC_DEVICE getBlockIdX() const {
+        return static_cast<unsigned int>(m_item.get_group_linear_id());
     }
 
-    auto inline TRACCC_DEVICE getBlockDimX() const {
-        return m_item.get_local_range(0);
+    unsigned int inline TRACCC_DEVICE getBlockDimX() const {
+        return static_cast<unsigned int>(m_item.get_local_range(0));
     }
 
-    auto inline TRACCC_DEVICE getGridDimX() const {
-        return m_item.get_global_range(0);
+    unsigned int inline TRACCC_DEVICE getGridDimX() const {
+        return static_cast<unsigned int>(m_item.get_global_range(0));
     }
 
     private: