New seeding

device/cuda/CMakeLists.txt

@@ -30,6 +30,11 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   # Clusterization
   "include/traccc/cuda/clusterization/clusterization_algorithm.hpp"
   "src/clusterization/clusterization_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> )

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/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::pair<vecmem::unique_alloc_ptr<kd_tree_t[]>, uint32_t> create_kd_tree(
+    vecmem::memory_resource&, const internal_spacepoint<spacepoint>* const,
+    uint32_t);
+}  // namespace traccc::cuda

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

@@ -0,0 +1,38 @@
+/** 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_seed.hpp>
+#include <traccc/cuda/seeding2/types/kd_tree.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_spacepoint<spacepoint>* 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<internal_seed[]>, uint32_t> run_seeding(
+    seedfinder_config, seedfilter_config, vecmem::memory_resource&,
+    const internal_spacepoint<spacepoint>* const, uint32_t,
+    vecmem::unique_alloc_ptr<kd_tree_t[]>&&, uint32_t);
+}  // namespace traccc::cuda

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

@@ -0,0 +1,25 @@
+/** 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_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 GPU-unfriendly EDM.
+ *
+ * @return A vector buffer containing the output seeds.
+ */
+vecmem::data::vector_buffer<seed> write_output(
+    const traccc::memory_resource &, uint32_t,
+    const internal_spacepoint<spacepoint> *const, const internal_seed *const);
+}  // namespace traccc::cuda

device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp

@@ -0,0 +1,35 @@
+/** 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/edm/seed.hpp>
+#include <traccc/edm/spacepoint.hpp>
+#include <traccc/seeding/detail/seeding_config.hpp>
+#include <traccc/utils/algorithm.hpp>
+#include <traccc/utils/memory_resource.hpp>
+#include <vecmem/memory/cuda/device_memory_resource.hpp>
+
+namespace traccc::cuda {
+/**
+ * @brief Alternative seed finding algorithm, using orthogonal range search
+ * implemented through a k-d tree.
+ */
+class seed_finding2 : public algorithm<vecmem::data::vector_buffer<seed>(
+                          const spacepoint_container_types::const_view&)> {
+    public:
+    seed_finding2(const traccc::memory_resource& mr);
+
+    vecmem::data::vector_buffer<seed> operator()(
+        const spacepoint_container_types::const_view& sps) const override;
+
+    private:
+    traccc::memory_resource m_output_mr;
+    seedfinder_config m_finder_conf;
+    seedfilter_config m_filter_conf;
+};
+}  // namespace traccc::cuda

device/cuda/include/traccc/cuda/seeding2/types/internal_seed.hpp

@@ -0,0 +1,29 @@
+/** 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 <array>
+#include <cstdint>
+
+namespace traccc::cuda {
+/**
+ * @brief Internal structure representing a single seed.
+ */
+struct internal_seed {
+    std::array<uint32_t, 3> spacepoints;
+    float weight;
+
+    __host__ __device__ static internal_seed Invalid() {
+        internal_seed r;
+
+        r.weight = std::numeric_limits<float>::lowest();
+
+        return r;
+    }
+};
+}  // namespace traccc::cuda

device/cuda/include/traccc/cuda/seeding2/types/kd_tree.hpp

@@ -0,0 +1,54 @@
+/** 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/range3d.hpp>
+
+namespace traccc::cuda {
+/**
+ * @brief Data structure representing a node in a k-d tree.
+ *
+ * @tparam _value_t Type of values in the tree.
+ * @tparam _nodes_i Number of values in each leaf node.
+ */
+template <typename _value_t, std::size_t _nodes_i>
+struct kd_tree_node {
+    struct element_wrapper_t {
+        _value_t value;
+        float phi, r, z;
+    };
+
+    enum class nodetype_e { LEAF, INTERNAL, NON_EXTANT, INCOMPLETE };
+
+    enum class pivot_e { Phi, R, Z };
+
+    static constexpr std::size_t MAX_NODES_PER_LEAF = _nodes_i;
+
+    nodetype_e type;
+
+    range3d range;
+
+    union {
+        struct {
+            element_wrapper_t points[MAX_NODES_PER_LEAF];
+            uint32_t point_count;
+        } leaf;
+        struct {
+            pivot_e dim;
+            float mid;
+        } internal;
+        struct {
+        } non_extant;
+        struct {
+            uint32_t begin, end;
+        } incomplete;
+    };
+};
+
+using kd_tree_t = kd_tree_node<uint32_t, 8>;
+}  // namespace traccc::cuda

device/cuda/include/traccc/cuda/seeding2/types/range3d.hpp

@@ -0,0 +1,58 @@
+/** 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 <array>
+#include <cstdint>
+
+namespace traccc::cuda {
+/**
+ * @brief Range in three dimensions, defined by minima and maxima in the phi,
+ * r, and z dimensions.
+ */
+struct range3d {
+    float phi_min, phi_max, r_min, r_max, z_min, z_max;
+
+    __host__ __device__ static range3d Infinite() {
+        range3d r;
+
+        r.phi_min = std::numeric_limits<float>::lowest();
+        r.phi_max = std::numeric_limits<float>::max();
+        r.r_min = 0.f;
+        r.r_max = std::numeric_limits<float>::max();
+        r.z_min = std::numeric_limits<float>::lowest();
+        r.z_max = std::numeric_limits<float>::max();
+
+        return r;
+    }
+
+    __host__ __device__ static range3d Degenerate() {
+        range3d r;
+
+        r.phi_min = std::numeric_limits<float>::max();
+        r.phi_max = std::numeric_limits<float>::lowest();
+        r.r_min = std::numeric_limits<float>::max();
+        r.r_max = 0.f;
+        r.z_min = std::numeric_limits<float>::max();
+        r.z_max = std::numeric_limits<float>::lowest();
+
+        return r;
+    }
+
+    __host__ __device__ bool intersects(const range3d& o) const {
+        return phi_min <= o.phi_max && o.phi_min < phi_max &&
+               r_min <= o.r_max && o.r_min < r_max && z_min <= o.z_max &&
+               o.z_min < z_max;
+    }
+
+    __host__ __device__ bool contains(float phi, float r, float z) const {
+        return phi_min <= phi && phi <= phi_max && r_min <= r && r <= r_max &&
+               z_min <= z && z <= z_max;
+    }
+};
+}  // namespace traccc::cuda