New seeding

core/include/traccc/edm/internal_spacepoint.hpp

@@ -76,6 +76,11 @@ struct internal_spacepoint {
         m_phi = 0;
     }
 
+    TRACCC_HOST_DEVICE
+    internal_spacepoint(scalar x, scalar y, scalar z, scalar r, scalar phi,
+                        unsigned int link)
+        : m_x(x), m_y(y), m_z(z), m_r(r), m_phi(phi), m_link_alt(link) {}
+
     TRACCC_HOST_DEVICE
     static inline internal_spacepoint<spacepoint_t> invalid_value() {
 

core/include/traccc/seeding/doublet_finding_helper.hpp

@@ -32,6 +32,11 @@ struct doublet_finding_helper {
         const internal_spacepoint<spacepoint>& sp2,
         const seedfinder_config& config);
 
+    static inline TRACCC_HOST_DEVICE bool isCompatible(
+        bool top, const internal_spacepoint<spacepoint>& sp1,
+        const internal_spacepoint<spacepoint>& sp2,
+        const seedfinder_config& config);
+
     /// Do the conformal transformation on doublet's coordinate
     ///
     /// @param sp1 is middle spacepoint
@@ -46,6 +51,17 @@ struct doublet_finding_helper {
                           const internal_spacepoint<spacepoint>& sp2);
 };
 
+bool doublet_finding_helper::isCompatible(
+    bool top, const internal_spacepoint<spacepoint>& sp1,
+    const internal_spacepoint<spacepoint>& sp2,
+    const seedfinder_config& config) {
+    if (top) {
+        return isCompatible<details::spacepoint_type::top>(sp1, sp2, config);
+    } else {
+        return isCompatible<details::spacepoint_type::bottom>(sp1, sp2, config);
+    }
+}
+
 template <details::spacepoint_type otherSpType>
 bool doublet_finding_helper::isCompatible(
     const internal_spacepoint<spacepoint>& sp1,

core/include/traccc/utils/array_wrapper.hpp

@@ -0,0 +1,203 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2021 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+#include <traccc/utils/functor.hpp>
+#include <type_traits>
+#include <utility>
+#include <vecmem/memory/memory_resource.hpp>
+#include <vecmem/memory/unique_ptr.hpp>
+
+namespace traccc::cuda {
+
+template <typename... Ts>
+struct pod {};
+
+template <>
+struct pod<> {};
+
+template <typename T, typename... Ts>
+struct pod<T, Ts...> {
+    T v;
+    pod<Ts...> r;
+};
+
+template <std::size_t I, typename... Ts>
+constexpr auto& pod_get(pod<Ts...>& p) {
+    if constexpr (I == 0) {
+        return p.v;
+    } else {
+        return pod_get<I - 1>(p.r);
+    }
+}
+
+template <template <typename...> typename F,
+          template <template <typename> typename> typename T>
+struct array_wrapper {
+    struct owner {
+        owner(vecmem::memory_resource& mr, std::size_t n) : data(mr, n) {}
+
+        typename details::functor::reapply<
+            F, typename T<details::functor::identity>::tuple_t>::type::owner
+            data;
+    };
+
+    struct handle {
+        using handle_t = typename details::functor::reapply<
+            F, typename T<details::functor::identity>::tuple_t>::type::handle;
+
+        handle(const owner& o) : data(o.data) {}
+
+        std::size_t size() const { return data.size(); }
+
+        template <std::size_t I>
+        __host__ __device__ auto& get(std::size_t i) {
+            return data.get<I>(i);
+        }
+
+        template <std::size_t I>
+        __host__ __device__ auto get(std::size_t i) const {
+            return data.get<I>(i);
+        }
+
+        template <std::size_t I>
+        __host__ __device__ auto get_identity(std::size_t i) const {
+            return data.get<I>(i);
+        }
+
+        template <std::size_t I>
+        __host__ __device__ auto& get_reference(std::size_t i) {
+            return data.get<I>(i);
+        }
+
+        template <std::size_t... Ns>
+        constexpr __host__ __device__ T<details::functor::identity>
+        _construct_helper_identity(std::index_sequence<Ns...>,
+                                   std::size_t i) const {
+            return T<details::functor::identity>{get_identity<Ns>(i)...};
+        }
+
+        template <std::size_t... Ns>
+        constexpr __host__ __device__ T<details::functor::reference>
+        _construct_helper_reference(std::index_sequence<Ns...>, std::size_t i) {
+            return T<details::functor::reference>{get_reference<Ns>(i)...};
+        }
+
+        __host__ __device__ T<details::functor::identity> operator[](
+            std::size_t i) const {
+            return _construct_helper_identity(
+                std::make_index_sequence<std::tuple_size_v<
+                    typename T<details::functor::identity>::tuple_t>>(),
+                i);
+        }
+
+        __host__ __device__ T<details::functor::reference> operator[](
+            std::size_t i) {
+            return _construct_helper_reference(
+                std::make_index_sequence<std::tuple_size_v<
+                    typename T<details::functor::identity>::tuple_t>>(),
+                i);
+        }
+
+        handle_t data;
+    };
+};
+
+template <std::size_t... Ns, typename... Ts>
+std::tuple<Ts*...> _get_ptrs(
+    std::index_sequence<Ns...>,
+    const std::tuple<vecmem::unique_alloc_ptr<Ts[]>...>& o) {
+    return {std::get<Ns>(o).get()...};
+}
+
+template <typename... Ts>
+struct soa {
+    struct owner {
+        owner(vecmem::memory_resource& mr, std::size_t n)
+            : _size(n), _ptrs{vecmem::make_unique_alloc<Ts[]>(mr, n)...} {}
+
+        std::size_t size() const { return _size; }
+
+        const std::tuple<vecmem::unique_alloc_ptr<Ts[]>...>& pointers() const {
+            return _ptrs;
+        }
+
+        private:
+        std::size_t _size;
+        std::tuple<vecmem::unique_alloc_ptr<Ts[]>...> _ptrs;
+    };
+
+    struct handle {
+        private:
+        using tuple_t = std::tuple<Ts*...>;
+
+        public:
+        handle(const owner& o)
+            : _size(o.size()),
+              _ptrs(_get_ptrs(std::make_index_sequence<sizeof...(Ts)>(),
+                              o.pointers())) {}
+
+        __host__ __device__ std::size_t size() const { return _size; }
+
+        template <std::size_t I>
+        __host__ __device__ auto& get(std::size_t i) {
+            return std::get<I>(_ptrs)[i];
+        }
+
+        template <std::size_t I>
+        __host__ __device__ const auto& get(std::size_t i) const {
+            return std::get<I>(_ptrs)[i];
+        }
+
+        private:
+        std::size_t _size;
+        std::tuple<Ts*...> _ptrs;
+    };
+};
+
+template <typename... Ts>
+struct aos {
+    struct owner {
+        owner(vecmem::memory_resource& mr, std::size_t n)
+            : _size(n), _ptr{vecmem::make_unique_alloc<pod<Ts...>[]>(mr, n)} {}
+
+        std::size_t size() const { return _size; }
+
+        const vecmem::unique_alloc_ptr<pod<Ts...>[]>& pointer() const {
+            return _ptr;
+        }
+
+        private : std::size_t _size;
+        vecmem::unique_alloc_ptr<pod<Ts...>[]> _ptr;
+    };
+
+    struct handle {
+        private:
+        using tuple_t = pod<Ts...>;
+
+        public:
+        handle(const owner& o) : _size(o.size()), _ptr(o.pointer().get()) {}
+
+        __host__ __device__ std::size_t size() const { return _size; }
+
+        template <std::size_t I>
+        __host__ __device__ auto& get(std::size_t i) {
+            return pod_get<I>(_ptr[i]);
+        }
+
+        template <std::size_t I>
+        __host__ __device__ const auto& get(std::size_t i) const {
+            return pod_get<I>(_ptr[i]);
+        }
+
+        private:
+        std::size_t _size;
+        tuple_t* _ptr;
+    };
+};
+}  // namespace traccc::cuda

core/include/traccc/utils/functor.hpp

@@ -0,0 +1,28 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2021-2022 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+namespace traccc::details::functor {
+template <typename T>
+using identity = T;
+
+template <typename T>
+using reference = T&;
+
+template <typename T>
+using const_reference = const T&;
+
+template <template <typename...> typename F, typename T>
+struct reapply {};
+
+template <template <typename...> typename F1,
+          template <typename...> typename F2, typename... Ts>
+struct reapply<F1, F2<Ts...>> {
+    using type = F1<Ts...>;
+};
+}  // namespace traccc::details::functor

device/cuda/CMakeLists.txt

@@ -43,8 +43,14 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "src/clusterization/clusterization_algorithm.cu"
   # Fitting
   "include/traccc/cuda/fitting/fitting_algorithm.hpp"
-  "src/fitting/fitting_algorithm.cu")
-
+  "src/fitting/fitting_algorithm.cu"
+  # Alternate seeding.
+  "src/seeding2/seed_finding.cu"
+  "src/seeding2/kernels/seed_finding_kernel.cu"
+  "src/seeding2/kernels/kd_tree_kernel.cu"
+  "src/seeding2/kernels/write_output_kernel.cu"
+)
+
 target_compile_options( traccc_cuda
   PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr> )
 target_link_libraries( traccc_cuda

device/cuda/include/traccc/cuda/seeding2/kernels/kd_tree_kernel.hpp

@@ -0,0 +1,24 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+#include <traccc/cuda/seeding2/types/internal_sp.hpp>
+#include <traccc/cuda/seeding2/types/kd_tree.hpp>
+#include <traccc/edm/internal_spacepoint.hpp>
+#include <traccc/edm/spacepoint.hpp>
+#include <vector>
+
+namespace traccc::cuda {
+/**
+ * @brief Creates a k-d tree from a given set of spacepoints.
+ *
+ * @return A pair containing the k-d tree nodes as well as the number of nodes.
+ */
+std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
+    vecmem::memory_resource&, internal_sp_owning_t&&, uint32_t);
+}  // namespace traccc::cuda

device/cuda/include/traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp

@@ -0,0 +1,39 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+#include <traccc/cuda/seeding2/types/internal_sp.hpp>
+#include <traccc/cuda/seeding2/types/kd_tree.hpp>
+#include <traccc/edm/alt_seed.hpp>
+#include <traccc/edm/internal_spacepoint.hpp>
+#include <traccc/edm/spacepoint.hpp>
+#include <traccc/seeding/detail/seeding_config.hpp>
+
+namespace traccc::cuda {
+/**
+ * @brief Convenience data structure for all the data we need for seeding.
+ */
+struct seed_finding_data_t {
+    const seedfinder_config finder_conf;
+    const seedfilter_config filter_conf;
+    const internal_sp_t spacepoints;
+    const std::size_t n_spacepoints;
+    const kd_tree_t tree;
+    const uint32_t tree_nodes;
+};
+
+/**
+ * @brief Execute the seed finding kernel itself.
+ *
+ * @return A pair containing the list of internal seeds as well as the number
+ * of seeds.
+ */
+std::pair<vecmem::unique_alloc_ptr<alt_seed[]>, uint32_t> run_seeding(
+    seedfinder_config, seedfilter_config, vecmem::memory_resource&,
+    internal_sp_t, uint32_t, kd_tree_t, uint32_t);
+}  // namespace traccc::cuda

device/cuda/include/traccc/cuda/seeding2/kernels/write_output_kernel.hpp

@@ -0,0 +1,26 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+#include <traccc/cuda/seeding2/types/internal_sp.hpp>
+#include <traccc/edm/alt_seed.hpp>
+#include <traccc/edm/internal_spacepoint.hpp>
+#include <traccc/edm/seed.hpp>
+#include <traccc/edm/spacepoint.hpp>
+#include <traccc/utils/memory_resource.hpp>
+
+namespace traccc::cuda {
+/**
+ * @brief Kernel to write output data back into traccc's EDM.
+ *
+ * @return A vector buffer containing the output seeds.
+ */
+alt_seed_collection_types::buffer write_output(const traccc::memory_resource &,
+                                               uint32_t, const internal_sp_t,
+                                               const alt_seed *const);
+}  // namespace traccc::cuda