Merge pull request #1089 from stephenswat/cleanup/max_seeds_per_spm

Deduplicate max seeds per middle SP parameter

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