From f3c0e1370ab6f85ee7dc490d553449fa5826661a Mon Sep 17 00:00:00 2001
From: Stephen Nicholas Swatman <stephen.nicholas.swatman@cern.ch>
Date: Tue, 25 Jun 2024 16:37:03 +0200
Subject: [PATCH] UpdatE

---
 core/include/traccc/edm/spacepoint.hpp        |   2 +
 device/cuda/CMakeLists.txt                    |   1 -
 .../cuda/seeding2/kernels/kd_tree_kernel.hpp  |   7 +-
 .../seeding2/kernels/seed_finding_kernel.hpp  |  10 +-
 .../seeding2/kernels/write_output_kernel.hpp  |  26 ---
 .../traccc/cuda/seeding2/seed_finding.hpp     |  23 ++-
 .../cuda/seeding2/types/internal_sp.hpp       |  35 ----
 .../traccc/cuda/utils/device_traits.hpp       |  56 ------
 .../src/seeding2/kernels/kd_tree_kernel.cu    | 164 ++++++++----------
 .../seeding2/kernels/seed_finding_kernel.cu   | 139 +++++++--------
 .../seeding2/kernels/write_output_kernel.cu   |  82 ---------
 device/cuda/src/seeding2/seed_finding2.cu     |  79 ++-------
 12 files changed, 182 insertions(+), 442 deletions(-)
 delete mode 100644 device/cuda/include/traccc/cuda/seeding2/kernels/write_output_kernel.hpp
 delete mode 100644 device/cuda/include/traccc/cuda/seeding2/types/internal_sp.hpp
 delete mode 100644 device/cuda/include/traccc/cuda/utils/device_traits.hpp
 delete mode 100644 device/cuda/src/seeding2/kernels/write_output_kernel.cu

diff --git a/core/include/traccc/edm/spacepoint.hpp b/core/include/traccc/edm/spacepoint.hpp
index 4814629248..add6bc1f08 100644
--- a/core/include/traccc/edm/spacepoint.hpp
+++ b/core/include/traccc/edm/spacepoint.hpp
@@ -42,6 +42,8 @@ struct spacepoint {
     scalar radius() const {
         return std::sqrt(global[0] * global[0] + global[1] * global[1]);
     }
+    TRACCC_HOST_DEVICE
+    scalar phi() const { return std::atan2(y(), x()); }
 };
 
 /// Comparison / ordering operator for spacepoints
diff --git a/device/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt
index 0e7067222f..784495698a 100644
--- a/device/cuda/CMakeLists.txt
+++ b/device/cuda/CMakeLists.txt
@@ -43,7 +43,6 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "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"
diff --git a/device/cuda/include/traccc/cuda/seeding2/kernels/kd_tree_kernel.hpp b/device/cuda/include/traccc/cuda/seeding2/kernels/kd_tree_kernel.hpp
index 4ce1325af7..63afb2f3ac 100644
--- a/device/cuda/include/traccc/cuda/seeding2/kernels/kd_tree_kernel.hpp
+++ b/device/cuda/include/traccc/cuda/seeding2/kernels/kd_tree_kernel.hpp
@@ -7,10 +7,10 @@
 
 #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 <vecmem/utils/copy.hpp>
 #include <vector>
 
 namespace traccc::cuda {
@@ -19,6 +19,7 @@ namespace traccc::cuda {
  *
  * @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);
+std::tuple<kd_tree_owning_t, uint32_t, vecmem::data::vector_buffer<std::size_t>>
+create_kd_tree(vecmem::memory_resource&, vecmem::copy& copy,
+               const spacepoint_collection_types::const_view&);
 }  // namespace traccc::cuda
diff --git a/device/cuda/include/traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp b/device/cuda/include/traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp
index 464b4dd2af..a23cc37d56 100644
--- a/device/cuda/include/traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp
+++ b/device/cuda/include/traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp
@@ -7,12 +7,12 @@
 
 #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/seed.hpp>
 #include <traccc/edm/spacepoint.hpp>
 #include <traccc/seeding/detail/seeding_config.hpp>
+#include <vecmem/utils/copy.hpp>
 
 namespace traccc::cuda {
 /**
@@ -21,7 +21,7 @@ namespace traccc::cuda {
  * @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);
+seed_collection_types::buffer run_seeding(
+    seedfinder_config, seedfilter_config, vecmem::memory_resource &,
+    vecmem::copy &, const spacepoint_collection_types::const_view &, kd_tree_t);
 }  // namespace traccc::cuda
diff --git a/device/cuda/include/traccc/cuda/seeding2/kernels/write_output_kernel.hpp b/device/cuda/include/traccc/cuda/seeding2/kernels/write_output_kernel.hpp
deleted file mode 100644
index 0bddf482d0..0000000000
--- a/device/cuda/include/traccc/cuda/seeding2/kernels/write_output_kernel.hpp
+++ /dev/null
@@ -1,26 +0,0 @@
-/** 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
diff --git a/device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp b/device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp
index a350a73730..119a18188d 100644
--- a/device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp
+++ b/device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp
@@ -7,29 +7,40 @@
 
 #pragma once
 
-#include <traccc/edm/alt_seed.hpp>
+#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>
+#include <vecmem/utils/copy.hpp>
+
+#include "traccc/cuda/utils/stream.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(
+class seed_finding2 : public algorithm<seed_collection_types::buffer(
                           const spacepoint_collection_types::const_view&)> {
     public:
-    seed_finding2(const traccc::memory_resource& mr);
+    seed_finding2(const seedfinder_config& config,
+                  const seedfilter_config& filter_config,
+                  const traccc::memory_resource& mr, vecmem::copy& copy,
+                  stream& str);
 
     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;
+    traccc::memory_resource m_mr;
+    vecmem::copy& m_copy;
+    stream& m_stream;
+
+    seedfinder_config m_seedfinder_config;
+    seedfilter_config m_seedfilter_config;
+
+    int m_warp_size;
 };
 }  // namespace traccc::cuda
diff --git a/device/cuda/include/traccc/cuda/seeding2/types/internal_sp.hpp b/device/cuda/include/traccc/cuda/seeding2/types/internal_sp.hpp
deleted file mode 100644
index 94b375ff4d..0000000000
--- a/device/cuda/include/traccc/cuda/seeding2/types/internal_sp.hpp
+++ /dev/null
@@ -1,35 +0,0 @@
-/** 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
diff --git a/device/cuda/include/traccc/cuda/utils/device_traits.hpp b/device/cuda/include/traccc/cuda/utils/device_traits.hpp
deleted file mode 100644
index c1cc59eb1c..0000000000
--- a/device/cuda/include/traccc/cuda/utils/device_traits.hpp
+++ /dev/null
@@ -1,56 +0,0 @@
-/** 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
diff --git a/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu b/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu
index ff5d54c301..e41f07cbbc 100644
--- a/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu
+++ b/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu
@@ -14,8 +14,6 @@
 #include <traccc/cuda/seeding2/seed_finding.hpp>
 #include <traccc/cuda/seeding2/types/kd_tree.hpp>
 #include <traccc/cuda/seeding2/types/range3d.hpp>
-#include <traccc/cuda/utils/definitions.hpp>
-#include <traccc/cuda/utils/device_traits.hpp>
 #include <traccc/cuda/utils/sort.hpp>
 #include <traccc/edm/seed.hpp>
 #include <traccc/edm/spacepoint.hpp>
@@ -25,6 +23,8 @@
 #include <vecmem/memory/unique_ptr.hpp>
 #include <vector>
 
+#include "../../utils/cuda_error_handling.hpp"
+
 namespace traccc::cuda {
 /**
  * @brief Bit mask for halo point indices.
@@ -62,23 +62,24 @@ constexpr uint32_t AUXILIARY_INDEX_MASK = (1u << 31);
  *
  * @return The ϕ value of the spacepoint with the given index.
  */
-__device__ float get_coordinate(const internal_sp_t spacepoints, uint32_t i,
-                                pivot_e dim) {
+__device__ float get_coordinate(
+    const spacepoint_collection_types::const_device& spacepoints, uint32_t i,
+    pivot_e dim) {
     bool is_real = !(i & AUXILIARY_INDEX_MASK);
 
     if (is_real) {
         if (dim == pivot_e::Phi) {
-            return spacepoints[i].phi;
+            return spacepoints[i].phi();
         } else if (dim == pivot_e::R) {
-            return spacepoints[i].radius;
+            return spacepoints[i].radius();
         } else {
-            return spacepoints[i].z;
+            return spacepoints[i].z();
         }
     } else {
         uint32_t real_i = i & ~AUXILIARY_INDEX_MASK;
 
         if (dim == pivot_e::Phi) {
-            float phi = spacepoints[real_i].phi;
+            float phi = spacepoints[real_i].phi();
 
             if (phi >= 0) {
                 return phi - 2.f * static_cast<float>(M_PI);
@@ -86,9 +87,9 @@ __device__ float get_coordinate(const internal_sp_t spacepoints, uint32_t i,
                 return phi + 2.f * static_cast<float>(M_PI);
             }
         } else if (dim == pivot_e::R) {
-            return spacepoints[real_i].radius;
+            return spacepoints[real_i].radius();
         } else {
-            return spacepoints[real_i].z;
+            return spacepoints[real_i].z();
         }
     }
 }
@@ -105,10 +106,11 @@ __device__ float get_coordinate(const internal_sp_t spacepoints, uint32_t i,
  * @param[in] sps The total number of spacepoints.
  * @param[out] extra_indices The output halo point count.
  */
-__global__ void initialize_index_vector(internal_sp_t spacepoints,
-                                        uint32_t* __restrict__ indices,
-                                        uint32_t sps,
-                                        uint32_t* __restrict__ extra_indices) {
+__global__ void initialize_index_vector(
+    const spacepoint_collection_types::const_view& _spacepoints,
+    uint32_t* __restrict__ indices, uint32_t sps,
+    uint32_t* __restrict__ extra_indices) {
+    const spacepoint_collection_types::const_device spacepoints(_spacepoints);
     cooperative_groups::grid_group grid = cooperative_groups::this_grid();
 
     if (grid.thread_rank() < sps) {
@@ -196,9 +198,9 @@ __device__ pivot_e next_pivot(pivot_e p) {
  * @return The approximate median value of the given indices in the given
  * pivot axis.
  */
-__device__ float get_approximate_median(internal_sp_t spacepoints,
-                                        const uint32_t* __restrict__ indices,
-                                        uint32_t num_indices, pivot_e pivot) {
+__device__ float get_approximate_median(
+    const spacepoint_collection_types::const_device& spacepoints,
+    const uint32_t* __restrict__ indices, uint32_t num_indices, pivot_e pivot) {
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
 
@@ -255,10 +257,10 @@ __device__ float get_approximate_median(internal_sp_t spacepoints,
  * @param[in] num_indices The number of indices for this function call.
  * @param[in] pivot The requested pivot axis.
  */
-__device__ uint32_t sort_and_get_median(internal_sp_t spacepoints,
-                                        const uint32_t* __restrict__ indices,
-                                        uint32_t* __restrict__ index_buffer,
-                                        uint32_t num_indices, pivot_e pivot) {
+__device__ uint32_t sort_and_get_median(
+    const spacepoint_collection_types::const_device& spacepoints,
+    const uint32_t* __restrict__ indices, uint32_t* __restrict__ index_buffer,
+    uint32_t num_indices, pivot_e pivot) {
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
 
@@ -299,10 +301,11 @@ __device__ uint32_t sort_and_get_median(internal_sp_t spacepoints,
  * @param[in] nid The node index.
  * @param[in] pivot The requested pivot axis.
  */
-__device__ bool create_internal_node(kd_tree_t tree, internal_sp_t spacepoints,
-                                     const uint32_t* __restrict__ indices,
-                                     uint32_t* __restrict__ index_buffer,
-                                     int nid, pivot_e pivot) {
+__device__ bool create_internal_node(
+    kd_tree_t tree,
+    const spacepoint_collection_types::const_device& spacepoints,
+    const uint32_t* __restrict__ indices, uint32_t* __restrict__ index_buffer,
+    int nid, pivot_e pivot) {
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
 
@@ -468,16 +471,16 @@ __device__ bool create_internal_node(kd_tree_t tree, internal_sp_t spacepoints,
  * @param[inout] _index_buffer Temporary buffer space for indices.
  * @param[in] iteration The iteration count.
  */
-__global__ void __launch_bounds__(512)
-    construct_kd_tree_big_step(kd_tree_t tree, uint32_t num_nodes,
-                               internal_sp_t spacepoints,
-                               const uint32_t* __restrict__ _indices_old,
-                               uint32_t* __restrict__ _indices_new,
-                               uint32_t iteration,
-                               uint32_t* __restrict__ work_list_current,
-                               uint32_t* __restrict__ work_list_size_current,
-                               uint32_t* __restrict__ work_list_next,
-                               uint32_t* __restrict__ work_list_size_next) {
+__global__ void construct_kd_tree_big_step(
+    kd_tree_t tree, uint32_t num_nodes,
+    const spacepoint_collection_types::const_view& _spacepoints,
+    const uint32_t* __restrict__ _indices_old,
+    uint32_t* __restrict__ _indices_new, uint32_t iteration,
+    uint32_t* __restrict__ work_list_current,
+    uint32_t* __restrict__ work_list_size_current,
+    uint32_t* __restrict__ work_list_next,
+    uint32_t* __restrict__ work_list_size_next) {
+    const spacepoint_collection_types::const_device spacepoints(_spacepoints);
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
 
@@ -556,16 +559,16 @@ __global__ void __launch_bounds__(512)
     }
 }
 
-__global__ void __launch_bounds__(96)
-    construct_kd_tree_small_step(kd_tree_t tree, uint32_t num_nodes,
-                                 internal_sp_t spacepoints,
-                                 const uint32_t* __restrict__ _indices_old,
-                                 uint32_t* __restrict__ _indices_new,
-                                 uint32_t iteration,
-                                 uint32_t* __restrict__ work_list_current,
-                                 uint32_t* __restrict__ work_list_size_current,
-                                 uint32_t* __restrict__ work_list_next,
-                                 uint32_t* __restrict__ work_list_size_next) {
+__global__ void construct_kd_tree_small_step(
+    kd_tree_t tree, uint32_t num_nodes,
+    const spacepoint_collection_types::const_view& _spacepoints,
+    const uint32_t* __restrict__ _indices_old,
+    uint32_t* __restrict__ _indices_new, uint32_t iteration,
+    uint32_t* __restrict__ work_list_current,
+    uint32_t* __restrict__ work_list_size_current,
+    uint32_t* __restrict__ work_list_next,
+    uint32_t* __restrict__ work_list_size_next) {
+    const spacepoint_collection_types::const_device spacepoints(_spacepoints);
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
 
@@ -654,28 +657,11 @@ uint32_t round_to_power_2(uint32_t i) {
     return power;
 }
 
-__global__ void bake_spacepoints_kernel(uint32_t n, internal_sp_t oldsps,
-                                        internal_sp_t newsps,
-                                        uint32_t* __restrict__ indices) {
-    cooperative_groups::grid_group grid = cooperative_groups::this_grid();
-
-    uint32_t i = grid.thread_rank();
+std::tuple<kd_tree_owning_t, uint32_t, vecmem::data::vector_buffer<std::size_t>>
+create_kd_tree(vecmem::memory_resource& mr, vecmem::copy& copy,
+               const spacepoint_collection_types::const_view& spacepoints) {
+    std::size_t n_sp = copy.get_size(spacepoints);
 
-    if (i < n) {
-        uint32_t ni = indices[i] & ~AUXILIARY_INDEX_MASK;
-
-        newsps[i].x = oldsps[ni].x;
-        newsps[i].y = oldsps[ni].y;
-        newsps[i].z = oldsps[ni].z;
-        newsps[i].phi = oldsps[ni].phi;
-        newsps[i].radius = oldsps[ni].radius;
-        newsps[i].link = oldsps[ni].link;
-    }
-}
-
-std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
-    vecmem::memory_resource& mr, internal_sp_owning_t&& spacepoints,
-    uint32_t n_sp) {
     /*
      * Allocate space for the indices. Since each spacepoint can produce at
      * most two indices, we allocate enough space to support this theoretical
@@ -691,7 +677,8 @@ std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
     vecmem::unique_alloc_ptr<uint32_t> extra_indices =
         vecmem::make_unique_alloc<uint32_t>(mr);
 
-    CUDA_ERROR_CHECK(cudaMemset(extra_indices.get(), 0, sizeof(uint32_t)));
+    TRACCC_CUDA_ERROR_CHECK(
+        cudaMemset(extra_indices.get(), 0, sizeof(uint32_t)));
 
     /*
      * Launch the index initialization kernel, which turns n spacepoints into
@@ -703,16 +690,17 @@ std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
                               threads_per_block1>>>(spacepoints, indices.get(),
                                                     n_sp, extra_indices.get());
 
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
     /*
      * Retrieve the number of halo points from the device.
      */
     uint32_t extra_indices_h;
 
-    CUDA_ERROR_CHECK(cudaMemcpy(&extra_indices_h, extra_indices.get(),
-                                sizeof(uint32_t), cudaMemcpyDeviceToHost));
+    TRACCC_CUDA_ERROR_CHECK(cudaMemcpy(&extra_indices_h, extra_indices.get(),
+                                       sizeof(uint32_t),
+                                       cudaMemcpyDeviceToHost));
 
     /*
      * The total number of indices is equal to the number of spacepoints plus
@@ -743,8 +731,10 @@ std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
     vecmem::unique_alloc_ptr<uint32_t> work_list_size_2 =
         vecmem::make_unique_alloc<uint32_t>(mr);
 
-    CUDA_ERROR_CHECK(cudaMemset(work_list_size_1.get(), 0, sizeof(uint32_t)));
-    CUDA_ERROR_CHECK(cudaMemset(work_list_size_2.get(), 0, sizeof(uint32_t)));
+    TRACCC_CUDA_ERROR_CHECK(
+        cudaMemset(work_list_size_1.get(), 0, sizeof(uint32_t)));
+    TRACCC_CUDA_ERROR_CHECK(
+        cudaMemset(work_list_size_2.get(), 0, sizeof(uint32_t)));
 
     /*
      * Run the initialization kernel for the tree itself, which sets the
@@ -757,8 +747,8 @@ std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
                                                num_indices, work_list_1.get(),
                                                work_list_size_1.get());
 
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
     /*
      * Allocate some temporary space for the index buffer, one half for "lower"
@@ -801,33 +791,25 @@ std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
                 work_list_size_2.get());
         }
 
-        CUDA_ERROR_CHECK(cudaGetLastError());
-        CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+        TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+        TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
         iteration++;
 
-        CUDA_ERROR_CHECK(cudaMemcpy(&remaining, work_list_size_2.get(),
-                                    sizeof(uint32_t), cudaMemcpyDeviceToHost));
+        TRACCC_CUDA_ERROR_CHECK(cudaMemcpy(&remaining, work_list_size_2.get(),
+                                           sizeof(uint32_t),
+                                           cudaMemcpyDeviceToHost));
 
         std::swap(work_list_1, work_list_2);
         std::swap(work_list_size_1, work_list_size_2);
         std::swap(indices, index_buffer);
 
-        CUDA_ERROR_CHECK(
+        TRACCC_CUDA_ERROR_CHECK(
             cudaMemset(work_list_size_2.get(), 0, sizeof(uint32_t)));
     }
 
-    internal_sp_owning_t new_sps(mr, num_indices);
-
-    bake_spacepoints_kernel<<<
-        num_indices / 1024u + (num_indices % 1024u == 0u ? 0u : 1u), 1024u>>>(
-        num_indices, internal_sp_t(spacepoints), internal_sp_t(new_sps),
-        indices.get());
-
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+    vecmem::data::vector_buffer<std::size_t> out;
 
-    return {kd_tree_owning_t{std::move(tree_owner)}, num_nodes,
-            std::move(new_sps)};
+    return {kd_tree_owning_t{std::move(tree_owner)}, num_nodes, std::move(out)};
 }
 }  // namespace traccc::cuda
diff --git a/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu b/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu
index 9e1204e7fd..601cef0b41 100644
--- a/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu
+++ b/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu
@@ -9,24 +9,25 @@
 
 #include <iostream>
 #include <traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp>
-#include <traccc/cuda/seeding2/types/internal_sp.hpp>
 #include <traccc/cuda/seeding2/types/kd_tree.hpp>
 #include <traccc/cuda/seeding2/types/range3d.hpp>
-#include <traccc/cuda/utils/definitions.hpp>
-#include <traccc/cuda/utils/device_traits.hpp>
 #include <traccc/cuda/utils/sort.hpp>
 #include <traccc/cuda/utils/sync.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/seeding/detail/lin_circle.hpp>
+#include <traccc/seeding/detail/spacepoint_type.hpp>
 #include <traccc/seeding/doublet_finding_helper.hpp>
 #include <traccc/seeding/triplet_finding_helper.hpp>
 
+#include "../../utils/cuda_error_handling.hpp"
+
 #define MAX_LOWER_SP_PER_MIDDLE 100u
 #define MAX_UPPER_SP_PER_MIDDLE 100u
 #define KD_TREE_TRAVERSAL_STACK_SIZE 128u
 #define WARPS_PER_BLOCK 8u
+#define WARP_SIZE 32u
 
 /**
  * @brief Maximum difference in φ between two spacepoints adjacent in the same
@@ -65,10 +66,10 @@ namespace traccc::cuda {
  */
 __device__ void retrieve_from_tree(
     const seedfinder_config finder_conf, const seedfilter_config filter_conf,
-    const internal_sp_t spacepoints, const kd_tree_t tree,
-    const internal_spacepoint<spacepoint> mi, const range3d range,
-    uint32_t mi_idx, bool upper, uint32_t* __restrict__ output_arr,
-    uint32_t* __restrict__ output_cnt) {
+    const spacepoint_collection_types::const_device& spacepoints,
+    const kd_tree_t tree, const internal_spacepoint<spacepoint> mi,
+    const range3d range, uint32_t mi_idx, bool upper,
+    uint32_t* __restrict__ output_arr, uint32_t* __restrict__ output_cnt) {
     __shared__ uint32_t stack[WARPS_PER_BLOCK][KD_TREE_TRAVERSAL_STACK_SIZE];
     __shared__ uint32_t leaves[WARPS_PER_BLOCK][WARP_SIZE];
 
@@ -137,16 +138,15 @@ __device__ void retrieve_from_tree(
                  * are more specific than those encoded in the search
                  * range.
                  */
-                float phi = spacepoints[point_id].phi,
-                      radius = spacepoints[point_id].radius,
-                      z = spacepoints[point_id].z;
+                float phi = spacepoints[point_id].phi(),
+                      radius = spacepoints[point_id].radius(),
+                      z = spacepoints[point_id].z();
 
                 if (range.contains(phi, radius, z) && point_id != mi_idx &&
                     doublet_finding_helper::isCompatible(
                         upper, mi,
-                        internal_spacepoint<spacepoint>(
-                            spacepoints[point_id].x, spacepoints[point_id].y, z,
-                            radius, phi, spacepoints[point_id].link),
+                        internal_spacepoint<spacepoint>(spacepoints[point_id],
+                                                        point_id, {0., 0.}),
                         finder_conf)) {
                     /*
                      * Reserve a spot in the output for this point, and
@@ -195,24 +195,21 @@ __device__ void retrieve_from_tree(
  */
 __device__ internal_seed make_seed(
     const seedfinder_config finder_conf, const seedfilter_config filter_conf,
-    const internal_sp_t spacepoints, const kd_tree_t tree, uint32_t lower,
-    uint32_t middle, internal_spacepoint<spacepoint> mi, uint32_t upper) {
-    const internal_spacepoint<spacepoint> lo(
-        spacepoints[lower].x, spacepoints[lower].y, spacepoints[lower].z,
-        spacepoints[lower].radius, spacepoints[lower].phi,
-        spacepoints[lower].link);
-    const internal_spacepoint<spacepoint> hi(
-        spacepoints[upper].x, spacepoints[upper].y, spacepoints[upper].z,
-        spacepoints[upper].radius, spacepoints[upper].phi,
-        spacepoints[upper].link);
+    const spacepoint_collection_types::const_device& spacepoints,
+    const kd_tree_t tree, uint32_t lower, uint32_t middle,
+    internal_spacepoint<spacepoint> mi, uint32_t upper) {
+    const internal_spacepoint<spacepoint> lo(spacepoints[lower], lower,
+                                             {0., 0.});
+    const internal_spacepoint<spacepoint> hi(spacepoints[upper], upper,
+                                             {0., 0.});
 
     /*
      * Find the lin-circles for the bottom and top pair.
      */
     lin_circle lm = doublet_finding_helper::transform_coordinates<
-        details::spacepoint_type::bottom>(mi, lo);
+        ::traccc::details::spacepoint_type::bottom>(mi, lo);
     lin_circle mh = doublet_finding_helper::transform_coordinates<
-        details::spacepoint_type::top>(mi, hi);
+        ::traccc::details::spacepoint_type::top>(mi, hi);
 
     scalar iSinTheta2 = 1 + lm.cotTheta() * lm.cotTheta();
     scalar scatteringInRegion2 = finder_conf.maxScatteringAngle2 * iSinTheta2;
@@ -352,13 +349,11 @@ __device__ range3d get_upper_range3d(const seedfinder_config finder_conf,
  * @param[inout] internal_seeds The array to write internal seeds to.
  * @param[in] increasing_z True iff we are looking for increasing-z seeds.
  */
-__device__ void run_helper(const seedfinder_config finder_conf,
-                           const seedfilter_config filter_conf,
-                           const internal_sp_t spacepoints,
-                           const kd_tree_t tree,
-                           const internal_spacepoint<spacepoint> sp,
-                           std::size_t sp_idx, internal_seed* internal_seeds,
-                           bool increasing_z) {
+__device__ void run_helper(
+    const seedfinder_config finder_conf, const seedfilter_config filter_conf,
+    const spacepoint_collection_types::const_device& spacepoints,
+    const kd_tree_t tree, const internal_spacepoint<spacepoint> sp,
+    std::size_t sp_idx, internal_seed* internal_seeds, bool increasing_z) {
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
     cooperative_groups::thread_block_tile<WARP_SIZE> warp =
@@ -528,11 +523,13 @@ __device__ void run_helper(const seedfinder_config finder_conf,
  * one-block-per-middle-SP fashion.
  *
  */
-__global__ void __launch_bounds__(WARP_SIZE* WARPS_PER_BLOCK)
-    seed_finding_kernel(const seedfinder_config finder_conf,
-                        const seedfilter_config filter_conf,
-                        const internal_sp_t spacepoints, const kd_tree_t tree,
-                        alt_seed* output_seeds, uint32_t* output_seed_size) {
+__global__ void seed_finding_kernel(
+    const seedfinder_config finder_conf, const seedfilter_config filter_conf,
+    const spacepoint_collection_types::const_view& _spacepoints,
+    const kd_tree_t tree, seed_collection_types::view _seeds) {
+    const spacepoint_collection_types::const_device spacepoints(_spacepoints);
+    seed_collection_types::device seeds(_seeds);
+
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
     cooperative_groups::thread_block_tile<WARP_SIZE> warp =
@@ -548,19 +545,17 @@ __global__ void __launch_bounds__(WARP_SIZE* WARPS_PER_BLOCK)
      * ignore some of the middle spacepoint candidates.
      */
     if (spacepoint_id >= spacepoints.size() ||
-        spacepoints[spacepoint_id].phi < finder_conf.phiMin ||
-        spacepoints[spacepoint_id].phi > finder_conf.phiMax ||
-        spacepoints[spacepoint_id].radius < finder_conf.rMin ||
-        spacepoints[spacepoint_id].radius > finder_conf.rMax ||
-        spacepoints[spacepoint_id].z < finder_conf.zMin ||
-        spacepoints[spacepoint_id].z > finder_conf.zMax) {
+        spacepoints[spacepoint_id].phi() < finder_conf.phiMin ||
+        spacepoints[spacepoint_id].phi() > finder_conf.phiMax ||
+        spacepoints[spacepoint_id].radius() < finder_conf.rMin ||
+        spacepoints[spacepoint_id].radius() > finder_conf.rMax ||
+        spacepoints[spacepoint_id].z() < finder_conf.zMin ||
+        spacepoints[spacepoint_id].z() > finder_conf.zMax) {
         return;
     }
 
-    const internal_spacepoint<spacepoint> sp(
-        spacepoints[spacepoint_id].x, spacepoints[spacepoint_id].y,
-        spacepoints[spacepoint_id].z, spacepoints[spacepoint_id].radius,
-        spacepoints[spacepoint_id].phi, spacepoints[spacepoint_id].link);
+    const internal_spacepoint<spacepoint> sp(spacepoints[spacepoint_id],
+                                             spacepoint_id, {0., 0.});
 
     /*
      * The internal seeds are stored in this shared array.
@@ -614,7 +609,7 @@ __global__ void __launch_bounds__(WARP_SIZE* WARPS_PER_BLOCK)
          * Next, reserve space atomically in the output array for our new
          * seeds.
          */
-        uint32_t idx = atomicAdd(output_seed_size, num_valid);
+        uint32_t idx = seeds.bulk_append_implicit_unsafe(num_valid);
 
         /*
          * Finally, actually write the output to global memory. Again, this is
@@ -624,15 +619,15 @@ __global__ void __launch_bounds__(WARP_SIZE* WARPS_PER_BLOCK)
          * than invalid ones (if we have any).
          */
         for (uint32_t i = 0; i < num_valid; ++i) {
-            alt_seed& s = output_seeds[idx + i];
+            seed& s = seeds[idx + i];
             uint32_t j = WARP_SIZE + finder_conf.maxSeedsPerSpM - (i + 1);
 
             /*
              * Write the data back to the output.
              */
-            s.spB_link = spacepoints[internal_seeds[j].spacepoints[0]].link;
-            s.spM_link = spacepoints[spacepoint_id].link;
-            s.spT_link = spacepoints[internal_seeds[j].spacepoints[1]].link;
+            s.spB_link = internal_seeds[j].spacepoints[0];
+            s.spM_link = spacepoint_id;
+            s.spT_link = internal_seeds[j].spacepoints[1];
 
             s.weight = internal_seeds[i].weight;
             s.z_vertex = 0.f;
@@ -654,23 +649,29 @@ __global__ void __launch_bounds__(WARP_SIZE* WARPS_PER_BLOCK)
  * @return A unique pointer to an array of internal seeds, and the size of that
  * array.
  */
-std::pair<vecmem::unique_alloc_ptr<alt_seed[]>, uint32_t> run_seeding(
+seed_collection_types::buffer run_seeding(
     seedfinder_config finder_conf, seedfilter_config filter_conf,
-    vecmem::memory_resource& mr, const internal_sp_t spacepoints,
+    vecmem::memory_resource& mr, vecmem::copy& copy,
+    const spacepoint_collection_types::const_view& spacepoints,
     const kd_tree_t tree) {
+    std::size_t n_sp = copy.get_size(spacepoints);
+
     /*
      * Allocate space for output of seeds on the device.
      */
-    vecmem::unique_alloc_ptr<alt_seed[]> seeds_device =
-        vecmem::make_unique_alloc<alt_seed[]>(
-            mr, finder_conf.maxSeedsPerSpM * spacepoints.size());
+    std::size_t output_size = finder_conf.maxSeedsPerSpM * n_sp;
+
+    seed_collection_types::buffer seed_buffer(
+        output_size, mr, vecmem::data::buffer_type::resizable);
+    copy.setup(seed_buffer)->wait();
 
     /*
      * Allocate space for seed count on the device.
      */
     vecmem::unique_alloc_ptr<uint32_t> seed_count_device =
         vecmem::make_unique_alloc<uint32_t>(mr);
-    CUDA_ERROR_CHECK(cudaMemset(seed_count_device.get(), 0, sizeof(uint32_t)));
+    TRACCC_CUDA_ERROR_CHECK(
+        cudaMemset(seed_count_device.get(), 0, sizeof(uint32_t)));
 
     /*
      * Calculate the total amount of shared memory on top of that which is
@@ -680,27 +681,19 @@ std::pair<vecmem::unique_alloc_ptr<alt_seed[]>, uint32_t> run_seeding(
     std::size_t total_shared_memory =
         (WARPS_PER_BLOCK * (WARP_SIZE + finder_conf.maxSeedsPerSpM)) *
         sizeof(internal_seed);
-    std::size_t grid_size = (spacepoints.size() / WARPS_PER_BLOCK) +
-                            (spacepoints.size() % WARPS_PER_BLOCK == 0 ? 0 : 1);
+    std::size_t grid_size =
+        (n_sp / WARPS_PER_BLOCK) + (n_sp % WARPS_PER_BLOCK == 0 ? 0 : 1);
     std::size_t block_size = WARPS_PER_BLOCK * WARP_SIZE;
 
     /*
      * Call the kernel and all that jazz.
      */
     seed_finding_kernel<<<grid_size, block_size, total_shared_memory>>>(
-        finder_conf, filter_conf, spacepoints, tree, seeds_device.get(),
-        seed_count_device.get());
-
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+        finder_conf, filter_conf, spacepoints, tree, seed_buffer);
 
-    /*
-     * Transfer the seed count back to the host and then hand it to the user.
-     */
-    uint32_t seed_count_host;
-    CUDA_ERROR_CHECK(cudaMemcpy(&seed_count_host, seed_count_device.get(),
-                                sizeof(uint32_t), cudaMemcpyDeviceToHost));
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
-    return {std::move(seeds_device), seed_count_host};
+    return seed_buffer;
 }
 }  // namespace traccc::cuda
diff --git a/device/cuda/src/seeding2/kernels/write_output_kernel.cu b/device/cuda/src/seeding2/kernels/write_output_kernel.cu
deleted file mode 100644
index d9bb4a73f2..0000000000
--- a/device/cuda/src/seeding2/kernels/write_output_kernel.cu
+++ /dev/null
@@ -1,82 +0,0 @@
-/** 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
- */
-
-#include <cooperative_groups.h>
-
-#include <cstdint>
-#include <traccc/cuda/seeding2/kernels/write_output_kernel.hpp>
-#include <traccc/cuda/seeding2/types/internal_sp.hpp>
-#include <traccc/cuda/utils/definitions.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>
-#include <vecmem/containers/device_vector.hpp>
-#include <vecmem/containers/vector.hpp>
-#include <vecmem/utils/cuda/copy.hpp>
-
-namespace traccc::cuda {
-/**
- * @brief Main output writing kernel.
- *
- * @param[in] n_seeds The number of seeds to write.
- * @param[in] spacepoints Spacepoint array into which seeds are indexed.
- * @param[in] in_seeds Input internal seeds to read from.
- * @param[out] out_seeds Output seeds to write to.
- */
-__global__ void write_output_kernel(uint32_t n_seeds,
-                                    const internal_sp_t spacepoints,
-                                    const alt_seed* const in_seeds,
-                                    alt_seed_collection_types::view out_seeds) {
-    cooperative_groups::grid_group grid = cooperative_groups::this_grid();
-
-    /*
-     * Create a writable device vector for the seeds.
-     */
-    vecmem::device_vector<alt_seed> dvec(out_seeds);
-
-    int i = grid.thread_rank();
-
-    if (i < n_seeds) {
-        /*
-         * Get the output seed for the current thread.
-         */
-        dvec.at(i) = in_seeds[i];
-    }
-}
-
-alt_seed_collection_types::buffer write_output(
-    const traccc::memory_resource& mr, uint32_t n_seeds,
-    const internal_sp_t spacepoints, const alt_seed* const seeds) {
-    static constexpr std::size_t threads_per_block = 256;
-
-    /*
-     * Create a "vector buffer" that we can write our output to.
-     */
-    alt_seed_collection_types::buffer out(n_seeds, mr.main);
-
-    /*
-     * Fetch the current CUDA device number.
-     */
-    int device;
-    CUDA_ERROR_CHECK(cudaGetDevice(&device));
-
-    /*
-     * Execute the main output kernel.
-     */
-    write_output_kernel<<<(n_seeds / threads_per_block) +
-                              (n_seeds % threads_per_block == 0 ? 0 : 1),
-                          threads_per_block>>>(n_seeds, spacepoints, seeds,
-                                               out);
-
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
-
-    return out;
-}
-}  // namespace traccc::cuda
diff --git a/device/cuda/src/seeding2/seed_finding2.cu b/device/cuda/src/seeding2/seed_finding2.cu
index 0837873889..222660c2a6 100644
--- a/device/cuda/src/seeding2/seed_finding2.cu
+++ b/device/cuda/src/seeding2/seed_finding2.cu
@@ -9,10 +9,8 @@
 #include <memory>
 #include <traccc/cuda/seeding2/kernels/kd_tree_kernel.hpp>
 #include <traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp>
-#include <traccc/cuda/seeding2/kernels/write_output_kernel.hpp>
 #include <traccc/cuda/seeding2/seed_finding.hpp>
 #include <traccc/cuda/seeding2/types/kd_tree.hpp>
-#include <traccc/cuda/utils/definitions.hpp>
 #include <traccc/edm/seed.hpp>
 #include <traccc/edm/spacepoint.hpp>
 #include <traccc/seeding/detail/lin_circle.hpp>
@@ -22,75 +20,28 @@
 #include <vecmem/utils/cuda/copy.hpp>
 #include <vector>
 
-namespace {
-/// Helper function that would produce a default seed-finder configuration
-__host__ traccc::seedfinder_config default_seedfinder2_config() {
-    traccc::seedfinder_config config;
-    traccc::seedfinder_config config_copy = config.toInternalUnits();
-    config.highland = 13.6f * std::sqrt(config_copy.radLengthPerSeed) *
-                      (1.f + 0.038f * std::log(config_copy.radLengthPerSeed));
-    float maxScatteringAngle = config.highland / config_copy.minPt;
-    config.maxScatteringAngle2 = maxScatteringAngle * maxScatteringAngle;
-    // helix radius in homogeneous magnetic field. Units are Kilotesla, MeV
-    // and millimeter
-    config.pTPerHelixRadius = 300.f * config_copy.bFieldInZ;
-    config.minHelixDiameter2 =
-        std::pow(config_copy.minPt * 2.f / config.pTPerHelixRadius, 2.f);
-    config.pT2perRadius =
-        std::pow(config.highland / config.pTPerHelixRadius, 2.f);
-    return config;
-}
-}  // namespace
+#include "../utils/cuda_error_handling.hpp"
+#include "../utils/utils.hpp"
 
 namespace traccc::cuda {
-__global__ void construct_internal_sp_kernel(internal_sp_t out,
-                                             const spacepoint* in,
-                                             std::size_t n_sp) {
-    int tid = blockIdx.x * blockDim.x + threadIdx.x;
-
-    if (tid < n_sp) {
-        out[tid].x = in[tid].x();
-        out[tid].y = in[tid].y();
-        out[tid].z = in[tid].z();
-        out[tid].phi = atan2f(in[tid].y(), in[tid].x());
-        out[tid].radius = in[tid].radius();
-        out[tid].link = tid;
-    }
-}
-
-seed_finding2::seed_finding2(const traccc::memory_resource& mr)
-    : m_output_mr(mr), m_finder_conf(default_seedfinder2_config()) {}
+seed_finding2::seed_finding2(const seedfinder_config& config,
+                             const seedfilter_config& filter_config,
+                             const traccc::memory_resource& mr,
+                             vecmem::copy& copy, stream& str)
+    : m_seedfinder_config(config),
+      m_seedfilter_config(filter_config),
+      m_mr(mr),
+      m_copy(copy),
+      m_stream(str),
+      m_warp_size(details::get_warp_size(str.device())) {}
 
 seed_finding2::output_type seed_finding2::operator()(
     const spacepoint_collection_types::const_view& sps) const {
-    vecmem::cuda::copy copy;
-    vecmem::cuda::device_memory_resource mr;
-    std::size_t n_sp = copy.get_size(sps);
-
-    internal_sp_owning_t internal_sp_device(mr, n_sp);
-
-    static constexpr std::size_t copy_kernel_threads_per_block = 256;
-    construct_internal_sp_kernel<<<
-        (n_sp / copy_kernel_threads_per_block) +
-            (n_sp % copy_kernel_threads_per_block == 0 ? 0 : 1),
-        copy_kernel_threads_per_block>>>(internal_sp_t(internal_sp_device),
-                                         sps.ptr(), n_sp);
-
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
     auto [kd_tree_device, kd_tree_size, internal_sp_device_new] =
-        create_kd_tree(mr, std::move(internal_sp_device), n_sp);
-
-    vecmem::unique_alloc_ptr<alt_seed[]> seeds_device;
-    uint32_t seed_count;
-
-    std::tie(seeds_device, seed_count) = run_seeding(
-        m_finder_conf, m_filter_conf, mr, internal_sp_t(internal_sp_device_new),
-        kd_tree_t(kd_tree_device));
+        create_kd_tree(m_mr.main, m_copy, sps);
 
-    return write_output(m_output_mr, seed_count,
-                        internal_sp_t(internal_sp_device_new),
-                        seeds_device.get());
+    return run_seeding(m_seedfinder_config, m_seedfilter_config, m_mr.main,
+                       m_copy, sps, kd_tree_t(kd_tree_device));
 }
 }  // namespace traccc::cuda