test

core/include/traccc/seeding/doublet_finding_helper.hpp

@@ -34,6 +34,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
@@ -48,6 +53,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 TRACCC_HOST_DEVICE
 doublet_finding_helper::isCompatible(const internal_spacepoint<spacepoint>& sp1,

device/cuda/CMakeLists.txt

@@ -39,6 +39,11 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "src/seeding/spacepoint_binning.cu"
   "include/traccc/cuda/seeding/experimental/spacepoint_formation.hpp"
   "src/seeding/experimental/spacepoint_formation.cu"
+  # Alternate seeding.
+  "src/seeding2/seed_finding2.cu"
+  "src/seeding2/kernels/seed_finding_kernel.cu"
+  "src/seeding2/kernels/kd_tree_kernel.cu"
+  "src/seeding2/kernels/write_output_kernel.cu"
   # Clusterization
   "include/traccc/cuda/clusterization/clusterization_algorithm.hpp"
   "src/clusterization/clusterization_algorithm.cu"

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,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_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 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, kd_tree_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

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_sp.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 <array>
+#include <cstdint>
+#include <traccc/edm/internal_spacepoint.hpp>
+#include <traccc/edm/spacepoint.hpp>
+#include <traccc/utils/array_wrapper.hpp>
+#include <vecmem/memory/unique_ptr.hpp>
+
+namespace traccc::cuda {
+template <template <typename> typename F>
+struct sp_t {
+    using tuple_t = std::tuple<float, float, float, float, float, unsigned int>;
+
+    F<float> x;
+    F<float> y;
+    F<float> z;
+    F<float> phi;
+    F<float> radius;
+    F<unsigned int> link;
+};
+
+using internal_sp = array_wrapper<soa, sp_t>;
+
+using internal_sp_owning_t = internal_sp::owner;
+
+using internal_sp_t = internal_sp::handle;
+}  // namespace traccc::cuda

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

@@ -0,0 +1,36 @@
+/** 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>
+#include <traccc/utils/array_wrapper.hpp>
+#include <vecmem/memory/unique_ptr.hpp>
+
+namespace traccc::cuda {
+enum class nodetype_e { LEAF, INTERNAL, NON_EXTANT };
+
+enum class pivot_e { Phi, R, Z };
+
+template <template <typename> typename F>
+struct node_t {
+    using tuple_t =
+        std::tuple<nodetype_e, range3d, uint32_t, uint32_t, pivot_e, float>;
+
+    F<nodetype_e> type;
+    F<range3d> range;
+    F<uint32_t> begin, end;
+    F<pivot_e> dim;
+    F<float> mid;
+};
+
+using kd_tree = array_wrapper<soa, node_t>;
+
+using kd_tree_owning_t = kd_tree::owner;
+
+using kd_tree_t = kd_tree::handle;
+}  // namespace traccc::cuda

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

@@ -0,0 +1,73 @@
+/** 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__ __forceinline__ 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__ __forceinline__ 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__ __forceinline__ static range3d Union(const range3d& a,
+                                                             const range3d& b) {
+        return {std::min(a.phi_min, b.phi_min), std::max(a.phi_max, b.phi_max),
+                std::min(a.r_min, b.r_min),     std::max(a.r_max, b.r_max),
+                std::min(a.z_min, b.z_min),     std::max(a.z_max, b.z_max)};
+    }
+
+    __host__ __device__ __forceinline__ 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__ __forceinline__ bool dominates(const range3d& o) const {
+        return phi_min <= o.phi_min && o.phi_max <= phi_max &&
+               r_min <= o.r_min && o.r_max <= r_max && z_min <= o.z_min &&
+               o.z_max <= z_max;
+    }
+
+    __host__ __device__ __forceinline__ 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

device/cuda/include/traccc/cuda/utils/device_traits.hpp

@@ -0,0 +1,56 @@
+/** 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
+
+#if defined(__CUDACC__) && defined(__CUDA_ARCH__)
+// Set values of maximum resident threads per SM
+#if __CUDA_ARCH__ < 500
+#pragma message \
+    "Very old CUDA architecture, setting maximum resident threads per SM to 1024."
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 1024u
+#elif __CUDA_ARCH__ <= 720
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 2048u
+#elif __CUDA_ARCH__ <= 750
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 1024u
+#elif __CUDA_ARCH__ <= 800
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 2048u
+#elif __CUDA_ARCH__ <= 890
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 1536u
+#elif __CUDA_ARCH__ <= 900
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 2048u
+#pragma message \
+    "Unknown CUDA architecture, setting maximum resident threads per SM to 1024."
+#define CUDA_MAX_RESIDENT_THREADS_PER_SM 1024u
+#endif
+
+// Set values of maximum resident blocks per SM
+#if __CUDA_ARCH__ < 500
+#pragma message \
+    "Very old CUDA architecture, setting maximum resident blocks per SM to 16."
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 16u
+#elif __CUDA_ARCH__ <= 720
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 32u
+#elif __CUDA_ARCH__ <= 750
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 16u
+#elif __CUDA_ARCH__ <= 800
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 32u
+#elif __CUDA_ARCH__ <= 870
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 16u
+#elif __CUDA_ARCH__ <= 890
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 24u
+#elif __CUDA_ARCH__ <= 900
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 32u
+#else
+#pragma message \
+    "Very old CUDA architecture, setting maximum resident blocks per SM to 16."
+#define CUDA_MAX_RESIDENT_BLOCKS_PER_SM 16u
+#endif
+
+// Set values of maximum threadss per block
+#define CUDA_MAX_THREADS_PER_BLOCK 1024u
+#endif