From c79b0972ceafce9a97c5e726cb04ba694b1b2437 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

---
 .../cuda/seeding2/kernels/kd_tree_kernel.hpp  |  1 -
 .../seeding2/kernels/seed_finding_kernel.hpp  |  7 +-
 .../seeding2/kernels/write_output_kernel.hpp  | 26 -------
 .../traccc/cuda/seeding2/seed_finding.hpp     | 22 ++++--
 .../cuda/seeding2/types/internal_sp.hpp       | 35 ---------
 .../src/seeding2/kernels/kd_tree_kernel.cu    | 30 ++++----
 .../seeding2/kernels/seed_finding_kernel.cu   | 39 +++++-----
 .../seeding2/kernels/write_output_kernel.cu   | 24 +++----
 device/cuda/src/seeding2/seed_finding2.cu     | 72 ++++---------------
 9 files changed, 79 insertions(+), 177 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

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..b21342135e 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,7 +7,6 @@
 
 #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>
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..2d04d73f7f 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,9 +7,8 @@
 
 #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/seed.hpp>
 #include <traccc/edm/internal_spacepoint.hpp>
 #include <traccc/edm/spacepoint.hpp>
 #include <traccc/seeding/detail/seeding_config.hpp>
@@ -21,7 +20,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(
+std::pair<vecmem::unique_alloc_ptr<seed[]>, uint32_t> run_seeding(
     seedfinder_config, seedfilter_config, vecmem::memory_resource&,
-    internal_sp_t, kd_tree_t);
+    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..1c7c8d4a85 100644
--- a/device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp
+++ b/device/cuda/include/traccc/cuda/seeding2/seed_finding.hpp
@@ -7,29 +7,39 @@
 
 #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 "traccc/cuda/utils/stream.hpp"
+#include <vecmem/utils/copy.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/src/seeding2/kernels/kd_tree_kernel.cu b/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu
index ff5d54c301..dc28a5f8b7 100644
--- a/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu
+++ b/device/cuda/src/seeding2/kernels/kd_tree_kernel.cu
@@ -9,12 +9,12 @@
 
 #include <iostream>
 #include <limits>
+#include "../../utils/cuda_error_handling.hpp"
 #include <memory>
 #include <traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp>
 #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>
@@ -691,7 +691,7 @@ 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,15 +703,15 @@ 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(),
+    TRACCC_CUDA_ERROR_CHECK(cudaMemcpy(&extra_indices_h, extra_indices.get(),
                                 sizeof(uint32_t), cudaMemcpyDeviceToHost));
 
     /*
@@ -743,8 +743,8 @@ 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 +757,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,19 +801,19 @@ 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(),
+        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)));
     }
 
@@ -824,8 +824,8 @@ std::tuple<kd_tree_owning_t, uint32_t, internal_sp_owning_t> create_kd_tree(
         num_indices, internal_sp_t(spacepoints), internal_sp_t(new_sps),
         indices.get());
 
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
     return {kd_tree_owning_t{std::move(tree_owner)}, num_nodes,
             std::move(new_sps)};
diff --git a/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu b/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu
index 9e1204e7fd..cc6302dd5f 100644
--- a/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu
+++ b/device/cuda/src/seeding2/kernels/seed_finding_kernel.cu
@@ -9,24 +9,24 @@
 
 #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/seed.hpp>
 #include <traccc/edm/internal_spacepoint.hpp>
 #include <traccc/edm/spacepoint.hpp>
 #include <traccc/seeding/detail/lin_circle.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
@@ -145,8 +145,7 @@ __device__ void retrieve_from_tree(
                     doublet_finding_helper::isCompatible(
                         upper, mi,
                         internal_spacepoint<spacepoint>(
-                            spacepoints[point_id].x, spacepoints[point_id].y, z,
-                            radius, phi, spacepoints[point_id].link),
+                            point_id, spacepoints[point_id], {0., 0.}),
                         finder_conf)) {
                     /*
                      * Reserve a spot in the output for this point, and
@@ -198,13 +197,13 @@ __device__ internal_seed make_seed(
     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);
+        spacepoints[lower].link, spacepoints[lower].x, spacepoints[lower].y, spacepoints[lower].z,
+        spacepoints[lower].radius, spacepoints[lower].phi
+        );
     const internal_spacepoint<spacepoint> hi(
-        spacepoints[upper].x, spacepoints[upper].y, spacepoints[upper].z,
-        spacepoints[upper].radius, spacepoints[upper].phi,
-        spacepoints[upper].link);
+        spacepoints[upper].link, spacepoints[upper].x, spacepoints[upper].y, spacepoints[upper].z,
+        spacepoints[upper].radius, spacepoints[upper].phi
+        );
 
     /*
      * Find the lin-circles for the bottom and top pair.
@@ -532,7 +531,7 @@ __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) {
+                        seed* output_seeds, uint32_t* output_seed_size) {
     cooperative_groups::thread_block block =
         cooperative_groups::this_thread_block();
     cooperative_groups::thread_block_tile<WARP_SIZE> warp =
@@ -624,7 +623,7 @@ __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 = output_seeds[idx + i];
             uint32_t j = WARP_SIZE + finder_conf.maxSeedsPerSpM - (i + 1);
 
             /*
@@ -654,15 +653,15 @@ __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(
+std::pair<vecmem::unique_alloc_ptr<seed[]>, uint32_t> run_seeding(
     seedfinder_config finder_conf, seedfilter_config filter_conf,
     vecmem::memory_resource& mr, const internal_sp_t spacepoints,
     const kd_tree_t tree) {
     /*
      * Allocate space for output of seeds on the device.
      */
-    vecmem::unique_alloc_ptr<alt_seed[]> seeds_device =
-        vecmem::make_unique_alloc<alt_seed[]>(
+    vecmem::unique_alloc_ptr<seed[]> seeds_device =
+        vecmem::make_unique_alloc<seed[]>(
             mr, finder_conf.maxSeedsPerSpM * spacepoints.size());
 
     /*
@@ -670,7 +669,7 @@ std::pair<vecmem::unique_alloc_ptr<alt_seed[]>, uint32_t> run_seeding(
      */
     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
@@ -691,14 +690,14 @@ std::pair<vecmem::unique_alloc_ptr<alt_seed[]>, uint32_t> run_seeding(
         finder_conf, filter_conf, spacepoints, tree, seeds_device.get(),
         seed_count_device.get());
 
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
     /*
      * 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(),
+    TRACCC_CUDA_ERROR_CHECK(cudaMemcpy(&seed_count_host, seed_count_device.get(),
                                 sizeof(uint32_t), cudaMemcpyDeviceToHost));
 
     return {std::move(seeds_device), seed_count_host};
diff --git a/device/cuda/src/seeding2/kernels/write_output_kernel.cu b/device/cuda/src/seeding2/kernels/write_output_kernel.cu
index d9bb4a73f2..fd961f0623 100644
--- a/device/cuda/src/seeding2/kernels/write_output_kernel.cu
+++ b/device/cuda/src/seeding2/kernels/write_output_kernel.cu
@@ -7,11 +7,9 @@
 
 #include <cooperative_groups.h>
 
+#include "../../utils/cuda_error_handling.hpp"
 #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/seed.hpp>
 #include <traccc/edm/internal_spacepoint.hpp>
 #include <traccc/edm/seed.hpp>
 #include <traccc/edm/spacepoint.hpp>
@@ -31,14 +29,14 @@ namespace traccc::cuda {
  */
 __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) {
+                                    const seed* const in_seeds,
+                                    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);
+    vecmem::device_vector<seed> dvec(out_seeds);
 
     int i = grid.thread_rank();
 
@@ -50,21 +48,21 @@ __global__ void write_output_kernel(uint32_t n_seeds,
     }
 }
 
-alt_seed_collection_types::buffer write_output(
+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) {
+    const internal_sp_t spacepoints, const 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);
+    seed_collection_types::buffer out(n_seeds, mr.main);
 
     /*
      * Fetch the current CUDA device number.
      */
     int device;
-    CUDA_ERROR_CHECK(cudaGetDevice(&device));
+    TRACCC_CUDA_ERROR_CHECK(cudaGetDevice(&device));
 
     /*
      * Execute the main output kernel.
@@ -74,8 +72,8 @@ alt_seed_collection_types::buffer write_output(
                           threads_per_block>>>(n_seeds, spacepoints, seeds,
                                                out);
 
-    CUDA_ERROR_CHECK(cudaGetLastError());
-    CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
     return out;
 }
diff --git a/device/cuda/src/seeding2/seed_finding2.cu b/device/cuda/src/seeding2/seed_finding2.cu
index 0837873889..8f426e8d37 100644
--- a/device/cuda/src/seeding2/seed_finding2.cu
+++ b/device/cuda/src/seeding2/seed_finding2.cu
@@ -7,12 +7,12 @@
 
 #include <limits>
 #include <memory>
+#include "../utils/utils.hpp"
+#include "../utils/cuda_error_handling.hpp"
 #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,74 +22,32 @@
 #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
-
 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;
+    vecmem::unique_alloc_ptr<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),
+        m_seedfinder_config, m_seedfilter_config, mr, internal_sp_t(internal_sp_device_new),
         kd_tree_t(kd_tree_device));
 
-    return write_output(m_output_mr, seed_count,
+    return write_output(m_mr, seed_count,
                         internal_sp_t(internal_sp_device_new),
                         seeds_device.get());
 }