New seeding

cmake/traccc-compiler-options-cuda.cmake

@@ -31,6 +31,8 @@ traccc_add_flag( CMAKE_CUDA_FLAGS "--expt-relaxed-constexpr" )
 # build.
 traccc_add_flag( CMAKE_CUDA_FLAGS_DEBUG "-G" )
 
+traccc_add_flag( CMAKE_CUDA_FLAGS_RELWITHDEBINFO "-lineinfo" )
+
 # More rigorous tests for the Debug builds.
 if( ( "${CUDAToolkit_VERSION}" VERSION_GREATER_EQUAL "10.2" ) AND
     ( "${CMAKE_CUDA_COMPILER_ID}" MATCHES "NVIDIA" ) )

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,

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_seed.hpp>
+#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 <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<internal_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,27 @@
+/** 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/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 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/alt_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<alt_seed_collection_types::buffer(
+                          const spacepoint_collection_types::const_view&)> {
+    public:
+    seed_finding2(const traccc::memory_resource& mr);
+
+    output_type operator()(
+        const spacepoint_collection_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/internal_sp.hpp

@@ -0,0 +1,143 @@
+/** 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>
+#include <traccc/edm/internal_spacepoint.hpp>
+#include <traccc/edm/spacepoint.hpp>
+#include <vecmem/memory/unique_ptr.hpp>
+
+namespace traccc::cuda {
+struct internal_sp_owning_t {
+    internal_sp_owning_t(vecmem::unique_alloc_ptr<float[]> &&_xs,
+                         vecmem::unique_alloc_ptr<float[]> &&_ys,
+                         vecmem::unique_alloc_ptr<float[]> &&_zs,
+                         vecmem::unique_alloc_ptr<float[]> &&_phis,
+                         vecmem::unique_alloc_ptr<float[]> &&_radii,
+                         vecmem::unique_alloc_ptr<unsigned int[]> &&_links)
+        : xs(std::forward<vecmem::unique_alloc_ptr<float[]>>(_xs)),
+          ys(std::forward<vecmem::unique_alloc_ptr<float[]>>(_ys)),
+          zs(std::forward<vecmem::unique_alloc_ptr<float[]>>(_zs)),
+          phis(std::forward<vecmem::unique_alloc_ptr<float[]>>(_phis)),
+          radii(std::forward<vecmem::unique_alloc_ptr<float[]>>(_radii)),
+          links(
+              std::forward<vecmem::unique_alloc_ptr<unsigned int[]>>(_links)) {}
+
+    internal_sp_owning_t() {}
+
+    vecmem::unique_alloc_ptr<float[]> xs, ys, zs, phis, radii;
+    vecmem::unique_alloc_ptr<unsigned int[]> links;
+
+    // internal_sp_owning_t(
+    //     vecmem::unique_alloc_ptr<internal_spacepoint<spacepoint>[]> && _sps
+    // ) :
+    // sps(std::forward<vecmem::unique_alloc_ptr<internal_spacepoint<spacepoint>[]>>(_sps))
+    // {}
+
+    // internal_sp_owning_t() {}
+
+    // vecmem::unique_alloc_ptr<internal_spacepoint<spacepoint>[]> sps;
+};
+
+struct internal_sp_t {
+    internal_sp_t(const internal_sp_owning_t &o)
+        : xs(o.xs.get()),
+          ys(o.ys.get()),
+          zs(o.zs.get()),
+          phis(o.phis.get()),
+          radii(o.radii.get()),
+          links(o.links.get()) {}
+
+    // internal_sp_t(const internal_sp_owning_t & o) : sps(o.sps.get()) {}
+
+    template <typename I>
+    __device__ __forceinline__ float &x(const I i) {
+        return xs[i];
+        // return sps[i].m_x;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float x(const I i) const {
+        return xs[i];
+        // return sps[i].m_x;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float &y(const I i) {
+        return ys[i];
+        // return sps[i].m_y;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float y(const I i) const {
+        return ys[i];
+        // return sps[i].m_y;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float &z(const I i) {
+        return zs[i];
+        // return sps[i].m_z;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float z(const I i) const {
+        return zs[i];
+        // return sps[i].m_z;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float &phi(const I i) {
+        return phis[i];
+        // return sps[i].m_phi;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float phi(const I i) const {
+        return phis[i];
+        // return sps[i].m_phi;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float &radius(const I i) {
+        return radii[i];
+        // return sps[i].m_r;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ float radius(const I i) const {
+        return radii[i];
+        // return sps[i].m_r;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ unsigned int &link(const I i) {
+        return links[i];
+        // return sps[i].m_link_alt;
+    }
+
+    template <typename I>
+    __device__ __forceinline__ unsigned int link(const I i) const {
+        return links[i];
+        // return sps[i].m_link_alt;
+    }
+
+    template <typename T, typename I>
+    __device__ __forceinline__ internal_spacepoint<T> get(const I i) const {
+        return internal_spacepoint<T>(x(i), y(i), z(i), radius(i), phi(i),
+                                      link(i));
+        // return sps[i];
+    }
+
+    private:
+    float *xs, *ys, *zs, *phis, *radii;
+    unsigned int *links;
+    // internal_spacepoint<spacepoint> * sps;
+};
+}  // namespace traccc::cuda