Deduplicate max seeds per middle SP parameter

Currently, we have three seeding parameters which all control the
maximum number of seeds per middle spacepoint. This is unnecessary and
makes life much more complicated for everyone. Thus, this commit
deduplicates those parameters, leaving a single parameter in the seed
finder configuration.

Author: stephenswat

core/include/traccc/seeding/detail/seeding_config.hpp

@@ -62,7 +62,7 @@ struct seedfinder_config {
     float maxPtScattering = 10.f * unit<float>::GeV;
 
     // for how many seeds can one SpacePoint be the middle SpacePoint?
-    unsigned int maxSeedsPerSpM = 10;
+    unsigned int maxSeedsPerSpM = 5;
 
     float bFieldInZ = 1.99724f * unit<float>::T;
     // location of beam in x,y plane.
@@ -197,15 +197,9 @@ struct seedfilter_config {
     // minimum distance between compatible seeds to be considered for weight
     // boost
     float deltaRMin = 5.f * unit<float>::mm;
-    // in dense environments many seeds may be found per middle space point.
-    // only seeds with the highest weight will be kept if this limit is reached.
-    unsigned int maxSeedsPerSpM = 20;
     // how often do you want to increase the weight of a seed for finding a
     // compatible seed?
     size_t compatSeedLimit = 2;
-    // Tool to apply experiment specific cuts on collected middle space points
-
-    size_t max_triplets_per_spM = 5;
 
     // seed weight increase
     float good_spB_min_radius = 150.f * unit<float>::mm;

core/src/seeding/seed_filtering.cpp

@@ -16,10 +16,14 @@
 
 namespace traccc::host::details {
 
-seed_filtering::seed_filtering(const seedfilter_config& config,
+seed_filtering::seed_filtering(const seedfinder_config& finder_config,
+                               const seedfilter_config& filter_config,
                                vecmem::memory_resource& mr,
                                std::unique_ptr<const Logger> logger)
-    : messaging(std::move(logger)), m_filter_config(config), m_mr{mr} {}
+    : messaging(std::move(logger)),
+      m_finder_config(finder_config),
+      m_filter_config(filter_config),
+      m_mr{mr} {}
 
 void seed_filtering::operator()(
     const edm::spacepoint_collection::const_device& spacepoints,
@@ -89,7 +93,7 @@ void seed_filtering::operator()(
         // Consider only a maximum number of triplets for the final quality cut.
         const std::size_t itLength =
             std::min(triplets_passing_single_seed_cuts.size(),
-                     m_filter_config.max_triplets_per_spM);
+                     static_cast<std::size_t>(m_finder_config.maxSeedsPerSpM));
         for (std::size_t i = 1; i < itLength; ++i) {
             if (seed_selecting_helper::cut_per_middle_sp(
                     m_filter_config, spacepoints, sp_grid_data,
@@ -103,7 +107,7 @@ void seed_filtering::operator()(
     // Add the best remaining seeds to the output collection.
     for (std::size_t i = 0;
          const triplet& triplet : triplets_passing_final_cuts) {
-        if (i++ >= m_filter_config.maxSeedsPerSpM) {
+        if (i++ >= m_finder_config.maxSeedsPerSpM) {
             break;
         }
         seeds.push_back({sp_grid.bin(triplet.sp1.bin_idx)[triplet.sp1.sp_idx],

core/src/seeding/seed_filtering.hpp

@@ -29,7 +29,8 @@ class seed_filtering : public messaging {
     public:
     /// Constructor with the seed filter configuration
     seed_filtering(
-        const seedfilter_config& config, vecmem::memory_resource& mr,
+        const seedfinder_config& finding_config,
+        const seedfilter_config& filter_config, vecmem::memory_resource& mr,
         std::unique_ptr<const Logger> logger = getDummyLogger().clone());
 
     /// Callable operator for the seed filtering
@@ -46,6 +47,8 @@ class seed_filtering : public messaging {
                     edm::seed_collection::host& seeds) const;
 
     private:
+    /// Seed finder configuration
+    seedfinder_config m_finder_config;
     /// Seed filter configuration
     seedfilter_config m_filter_config;
     /// The memory resource to use

core/src/seeding/seed_finding.cpp

@@ -25,7 +25,7 @@ struct seed_finding::impl {
                            logger->cloneWithSuffix("MidTopAlg")),
           m_triplet_finding(finder_config, filter_config, mr,
                             logger->cloneWithSuffix("TripletAlg")),
-          m_seed_filtering(filter_config, mr,
+          m_seed_filtering(finder_config, filter_config, mr,
                            logger->cloneWithSuffix("FilterAlg")),
           m_mr{mr} {}
 

device/alpaka/src/seeding/seed_finding.cpp

@@ -160,7 +160,8 @@ struct UpdateTripletWeights {
 struct SelectSeeds {
     template <typename TAcc>
     ALPAKA_FN_ACC void operator()(
-        TAcc const& acc, seedfilter_config filter_config,
+        TAcc const& acc, seedfinder_config finder_config,
+        seedfilter_config filter_config,
         edm::spacepoint_collection::const_view spacepoints,
         traccc::details::spacepoint_grid_types::const_view sp_view,
         device::triplet_counter_spM_collection_types::const_view spM_tc,
@@ -176,13 +177,12 @@ struct SelectSeeds {
         // triplets within weight updating kernel
         triplet* const data = ::alpaka::getDynSharedMem<triplet>(acc);
 
-        // Each thread uses max_triplets_per_spM elements of the array
-        triplet* dataPos =
-            &data[localThreadIdx * filter_config.max_triplets_per_spM];
+        // Each thread uses maxSeedsPerSpM elements of the array
+        triplet* dataPos = &data[localThreadIdx * finder_config.maxSeedsPerSpM];
 
-        device::select_seeds(globalThreadIdx, filter_config, spacepoints,
-                             sp_view, spM_tc, midBot_tc, triplet_view, dataPos,
-                             seed_view);
+        device::select_seeds(globalThreadIdx, finder_config, filter_config,
+                             spacepoints, sp_view, spM_tc, midBot_tc,
+                             triplet_view, dataPos, seed_view);
     }
 };
 
@@ -379,8 +379,9 @@ edm::seed_collection::buffer seed_finding::operator()(
 
     // Create seeds out of selected triplets
     ::alpaka::exec<Acc>(
-        queue, workDiv, kernels::SelectSeeds{}, m_seedfilter_config,
-        spacepoints_view, g2_view, vecmem::get_data(triplet_counter_spM_buffer),
+        queue, workDiv, kernels::SelectSeeds{}, m_seedfinder_config,
+        m_seedfilter_config, spacepoints_view, g2_view,
+        vecmem::get_data(triplet_counter_spM_buffer),
         vecmem::get_data(triplet_counter_midBot_buffer),
         vecmem::get_data(triplet_buffer), vecmem::get_data(seed_buffer));
 
@@ -414,9 +415,9 @@ struct BlockSharedMemDynSizeBytes<traccc::alpaka::kernels::SelectSeeds, TAcc> {
     ALPAKA_FN_HOST_ACC static auto getBlockSharedMemDynSizeBytes(
         traccc::alpaka::kernels::SelectSeeds const& /* kernel */,
         TVec const& blockThreadExtent, TVec const& /* threadElemExtent */,
-        traccc::seedfilter_config filter_config, TArgs const&... /* args */
+        traccc::seedfinder_config finder_config, TArgs const&... /* args */
         ) -> std::size_t {
-        return static_cast<std::size_t>(filter_config.max_triplets_per_spM *
+        return static_cast<std::size_t>(finder_config.maxSeedsPerSpM *
                                         blockThreadExtent.prod()) *
                sizeof(traccc::triplet);
     }

device/common/include/traccc/seeding/device/impl/select_seeds.ipp

@@ -60,7 +60,8 @@ TRACCC_HOST_DEVICE void insertionSort(triplet* arr,
 // Select seeds kernel
 TRACCC_HOST_DEVICE
 inline void select_seeds(
-    const global_index_t globalIndex, const seedfilter_config& filter_config,
+    const global_index_t globalIndex, const seedfinder_config& finder_config,
+    const seedfilter_config& filter_config,
     const edm::spacepoint_collection::const_view& spacepoints_view,
     const traccc::details::spacepoint_grid_types::const_view& sp_view,
     const triplet_counter_spM_collection_types::const_view& spM_tc_view,
@@ -123,10 +124,10 @@ inline void select_seeds(
 
         // if the number of good triplets is larger than the threshold,
         // the triplet with the lowest weight is removed
-        if (n_triplets_per_spM >= filter_config.max_triplets_per_spM) {
+        if (n_triplets_per_spM >= finder_config.maxSeedsPerSpM) {
 
             const std::size_t min_index =
-                details::min_elem(data, 0, filter_config.max_triplets_per_spM,
+                details::min_elem(data, 0, finder_config.maxSeedsPerSpM,
                                   [](const triplet lhs, const triplet rhs) {
                                       return lhs.weight > rhs.weight;
                                   });
@@ -142,7 +143,7 @@ inline void select_seeds(
 
         // if the number of good triplets is below the threshold, add
         // the current triplet to the array
-        else if (n_triplets_per_spM < filter_config.max_triplets_per_spM) {
+        else if (n_triplets_per_spM < finder_config.maxSeedsPerSpM) {
             data[n_triplets_per_spM] = {spB_loc,         spM_loc,
                                         spT_loc,         aTriplet.curvature,
                                         aTriplet.weight, aTriplet.z_vertex};
@@ -165,7 +166,7 @@ inline void select_seeds(
         const sp_location& spT_loc = aTriplet.sp3;
 
         // if the number of seeds reaches the threshold, break
-        if (n_seeds_per_spM >= filter_config.maxSeedsPerSpM + 1) {
+        if (n_seeds_per_spM >= finder_config.maxSeedsPerSpM + 1) {
             break;
         }
 

device/common/include/traccc/seeding/device/select_seeds.hpp

@@ -35,7 +35,8 @@ namespace traccc::device {
 ///
 TRACCC_HOST_DEVICE
 inline void select_seeds(
-    global_index_t globalIndex, const seedfilter_config& filter_config,
+    global_index_t globalIndex, const seedfinder_config& finder_config,
+    const seedfilter_config& filter_config,
     const edm::spacepoint_collection::const_view& spacepoints_view,
     const traccc::details::spacepoint_grid_types::const_view& sp_grid_view,
     const triplet_counter_spM_collection_types::const_view& spM_tc_view,

device/cuda/src/seeding/seed_finding.cu

@@ -129,7 +129,7 @@ __global__ void update_triplet_weights(
 
 /// CUDA kernel for running @c traccc::device::select_seeds
 __global__ void select_seeds(
-    seedfilter_config filter_config,
+    seedfinder_config finder_config, seedfilter_config filter_config,
     edm::spacepoint_collection::const_view spacepoints,
     traccc::details::spacepoint_grid_types::const_view sp_view,
     device::triplet_counter_spM_collection_types::const_view spM_tc,
@@ -140,12 +140,12 @@ __global__ void select_seeds(
     // Array for temporary storage of triplets for comparing within seed
     // selecting kernel
     extern __shared__ triplet data2[];
-    // Each thread uses max_triplets_per_spM elements of the array
-    triplet* dataPos = &data2[threadIdx.x * filter_config.max_triplets_per_spM];
+    // Each thread uses maxSeedsPerSpM elements of the array
+    triplet* dataPos = &data2[threadIdx.x * finder_config.maxSeedsPerSpM];
 
-    device::select_seeds(details::global_index1(), filter_config, spacepoints,
-                         sp_view, spM_tc, midBot_tc, triplet_view, dataPos,
-                         seed_view);
+    device::select_seeds(details::global_index1(), finder_config, filter_config,
+                         spacepoints, sp_view, spM_tc, midBot_tc, triplet_view,
+                         dataPos, seed_view);
 }
 
 }  // namespace kernels
@@ -355,14 +355,14 @@ edm::seed_collection::buffer seed_finding::operator()(
         nSeedSelectingThreads;
 
     // Create seeds out of selected triplets
-    kernels::select_seeds<<<nSeedSelectingBlocks, nSeedSelectingThreads,
-                            sizeof(triplet) *
-                                m_seedfilter_config.max_triplets_per_spM *
-                                nSeedSelectingThreads,
-                            stream>>>(m_seedfilter_config, spacepoints_view,
-                                      g2_view, triplet_counter_spM_buffer,
-                                      triplet_counter_midBot_buffer,
-                                      triplet_buffer, seed_buffer);
+    kernels::
+        select_seeds<<<nSeedSelectingBlocks, nSeedSelectingThreads,
+                       sizeof(triplet) * m_seedfinder_config.maxSeedsPerSpM *
+                           nSeedSelectingThreads,
+                       stream>>>(
+            m_seedfinder_config, m_seedfilter_config, spacepoints_view, g2_view,
+            triplet_counter_spM_buffer, triplet_counter_midBot_buffer,
+            triplet_buffer, seed_buffer);
     TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
     return seed_buffer;

device/sycl/src/seeding/seed_finding.sycl

@@ -338,7 +338,7 @@ edm::seed_collection::buffer seed_finding::operator()(
     update_weights_kernel.wait_and_throw();
 
     // Check if device is capable of allocating sufficient local memory
-    assert(sizeof(triplet) * m_seedfilter_config.max_triplets_per_spM *
+    assert(sizeof(triplet) * m_seedfinder_config.maxSeedsPerSpM *
                seedSelectingLocalSize <
            details::get_queue(m_queue)
                .get_device()
@@ -350,25 +350,23 @@ edm::seed_collection::buffer seed_finding::operator()(
             // Array for temporary storage of triplets for comparing within
             // kernel
             vecmem::sycl::local_accessor<triplet> local_mem(
-                m_seedfilter_config.max_triplets_per_spM *
-                    seedSelectingLocalSize,
-                h);
+                m_seedfinder_config.maxSeedsPerSpM * seedSelectingLocalSize, h);
 
             h.parallel_for<kernels::select_seeds>(
                 seedSelectingRange,
-                [filter_config = m_seedfilter_config, spacepoints_view, g2_view,
+                [finder_config = m_seedfinder_config,
+                 filter_config = m_seedfilter_config, spacepoints_view, g2_view,
                  triplet_counter_spM_view, triplet_counter_midBot_view,
                  triplet_view, local_mem, seed_view](::sycl::nd_item<1> item) {
                     // Each thread uses compatSeedLimit elements of the array
-                    triplet* dataPos =
-                        &local_mem[item.get_local_id() *
-                                   filter_config.max_triplets_per_spM];
-
-                    device::select_seeds(details::global_index(item),
-                                         filter_config, spacepoints_view,
-                                         g2_view, triplet_counter_spM_view,
-                                         triplet_counter_midBot_view,
-                                         triplet_view, dataPos, seed_view);
+                    triplet* dataPos = &local_mem[item.get_local_id() *
+                                                  finder_config.maxSeedsPerSpM];
+
+                    device::select_seeds(
+                        details::global_index(item), finder_config,
+                        filter_config, spacepoints_view, g2_view,
+                        triplet_counter_spM_view, triplet_counter_midBot_view,
+                        triplet_view, dataPos, seed_view);
                 });
         })
         .wait_and_throw();

tests/cpu/compare_with_acts_seeding.cpp

@@ -164,7 +164,6 @@ TEST_P(CompareWithActsSeedingTests, Run) {
 
     // Start creating Seed filter object
     Acts::SeedFilterConfig sfconf;
-    sfconf.maxSeedsPerSpM = traccc::seedfilter_config().maxSeedsPerSpM;
     // there are a lot more variables here tbh
 
     // We also need some atlas-specific cut