Remove unused cell links from CCL kernel

This commit removes the `cell_links` vector from the CCL kernels. This
vector is never used for any meaningful output, so it's just unnecessary
memory accesses.

Author: stephenswat

device/alpaka/src/clusterization/clusterization_algorithm.cpp

@@ -42,8 +42,7 @@ struct ccl_kernel {
         uint32_t* backup_mutex_ptr,
         vecmem::data::vector_view<unsigned int> disjoint_set_view,
         vecmem::data::vector_view<unsigned int> cluster_size_view,
-        edm::measurement_collection::view measurements_view,
-        vecmem::data::vector_view<unsigned int> cell_links) const {
+        edm::measurement_collection::view measurements_view) const {
 
         details::thread_id1 thread_id(acc);
 
@@ -69,7 +68,7 @@ struct ccl_kernel {
                            f_view, gf_view, f_backup_view, gf_backup_view,
                            adjc_backup_view, adjv_backup_view, backup_mutex,
                            disjoint_set_view, cluster_size_view, barry_r,
-                           measurements_view, cell_links);
+                           measurements_view);
     }
 
 };  // struct ccl_kernel
@@ -120,7 +119,7 @@ void clusterization_algorithm::ccl_kernel(
         payload.config, payload.cells, payload.det_descr, payload.det_cond,
         payload.f_backup, payload.gf_backup, payload.adjc_backup,
         payload.adjv_backup, payload.backup_mutex, payload.disjoint_set,
-        payload.cluster_sizes, payload.measurements, payload.cell_links);
+        payload.cluster_sizes, payload.measurements);
 }
 
 void clusterization_algorithm::cluster_maker_kernel(

device/common/include/traccc/clusterization/device/aggregate_cluster.hpp

@@ -49,8 +49,7 @@ TRACCC_HOST_DEVICE inline void aggregate_cluster(
     const detector_conditions_description::const_device& det_cond,
     const vecmem::device_vector<index_t>& f, unsigned int start,
     unsigned int end, unsigned int cid,
-    edm::measurement_collection::device::proxy_type out,
-    vecmem::data::vector_view<unsigned int> cell_links, unsigned int link,
+    edm::measurement_collection::device::proxy_type out, unsigned int link,
     vecmem::device_vector<unsigned int>& disjoint_set,
     std::optional<std::reference_wrapper<unsigned int>> cluster_size);
 

device/common/include/traccc/clusterization/device/ccl_kernel.hpp

@@ -60,8 +60,6 @@ namespace traccc::device {
 ///     corresponding measurement.
 /// @param barrier  A generic object for block-wide synchronisation
 /// @param[out] measurements_view collection of measurements
-/// @param[out] cell_links    collection of links to measurements each cell is
-/// put into
 template <device::concepts::barrier barrier_t,
           device::concepts::thread_id1 thread_id_t>
 TRACCC_HOST_DEVICE inline void ccl_kernel(
@@ -80,8 +78,7 @@ TRACCC_HOST_DEVICE inline void ccl_kernel(
     vecmem::data::vector_view<unsigned int> disjoint_set_view,
     vecmem::data::vector_view<unsigned int> cluster_size_view,
     const barrier_t& barrier,
-    edm::measurement_collection::view measurements_view,
-    vecmem::data::vector_view<unsigned int> cell_links);
+    edm::measurement_collection::view measurements_view);
 
 }  // namespace traccc::device
 

device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp

@@ -131,8 +131,6 @@ class clusterization_algorithm
         const detector_conditions_description::const_view& det_cond;
         /// The measurement collection to fill
         edm::measurement_collection::view& measurements;
-        /// Buffer for linking cells to measurements
-        vecmem::data::vector_view<unsigned int>& cell_links;
         /// Buffer for backup of the first element links
         vecmem::data::vector_view<details::fallback_index_t>& f_backup;
         /// Buffer for backup of the group first element links

device/common/include/traccc/clusterization/device/impl/aggregate_cluster.ipp

@@ -23,10 +23,8 @@ TRACCC_HOST_DEVICE inline void aggregate_cluster(
     const vecmem::device_vector<index_t>& fll, const unsigned int start,
     const unsigned int end, const unsigned int cid,
     edm::measurement_collection::device::proxy_type out,
-    vecmem::data::vector_view<unsigned int> cell_links, const unsigned int link,
-    vecmem::device_vector<unsigned int>& disjoint_set,
+    const unsigned int link, vecmem::device_vector<unsigned int>& disjoint_set,
     std::optional<std::reference_wrapper<unsigned int>> cluster_size) {
-    vecmem::device_vector<unsigned int> cell_links_device(cell_links);
 
     /*
      * Now, we iterate over all other cells to check if they belong to our
@@ -110,7 +108,6 @@ TRACCC_HOST_DEVICE inline void aggregate_cluster(
         var[1] = (1.f - weight_factor) * var[1] +
                  weight_factor * (diff_old[1] * diff_new[1]);
 
-        cell_links_device.at(pos) = link;
         tmp_cluster_size++;
 
         if (disjoint_set.capacity()) {

device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp

@@ -147,8 +147,7 @@ TRACCC_HOST_DEVICE inline void ccl_core(
     const clustering_config& cfg, const thread_id_t& thread_id,
     std::size_t& partition_start, std::size_t& partition_end,
     vecmem::device_vector<index_t> f, vecmem::device_vector<index_t> gf,
-    vecmem::data::vector_view<unsigned int> cell_links, index_t* adjv,
-    unsigned char* adjc,
+    index_t* adjv, unsigned char* adjc,
     const edm::silicon_cell_collection::const_device& cells_device,
     const detector_design_description::const_device& det_desc,
     const detector_conditions_description::const_device& det_cond,
@@ -267,8 +266,7 @@ TRACCC_HOST_DEVICE inline void ccl_core(
                 cfg, cells_device, det_desc, det_cond, gf,
                 static_cast<unsigned int>(partition_start),
                 static_cast<unsigned int>(partition_end), cid,
-                measurements_device.at(meas_pos), cell_links, meas_pos,
-                disjoint_set,
+                measurements_device.at(meas_pos), meas_pos, disjoint_set,
                 (cluster_size.capacity()
                      ? std::optional<std::reference_wrapper<
                            unsigned int>>{cluster_size.at(meas_pos)}
@@ -295,8 +293,7 @@ TRACCC_HOST_DEVICE inline void ccl_kernel(
     vecmem::data::vector_view<unsigned int> disjoint_set_view,
     vecmem::data::vector_view<unsigned int> cluster_size_view,
     const barrier_t& barrier,
-    edm::measurement_collection::view measurements_view,
-    vecmem::data::vector_view<unsigned int> cell_links) {
+    edm::measurement_collection::view measurements_view) {
 
     // Construct device containers around the views.
     const edm::silicon_cell_collection::const_device cells_device(cells_view);
@@ -404,9 +401,8 @@ TRACCC_HOST_DEVICE inline void ccl_kernel(
             (thread_id.getLocalThreadIdX() * 4 * cfg.max_cells_per_thread *
              cfg.backup_size_multiplier);
         ccl_core(cfg, thread_id, partition_start, partition_end, f_backup,
-                 gf_backup, cell_links, adjv, adjc, cells_device, det_desc,
-                 det_cond, measurements_device, barrier, disjoint_set,
-                 cluster_size);
+                 gf_backup, adjv, adjc, cells_device, det_desc, det_cond,
+                 measurements_device, barrier, disjoint_set, cluster_size);
     } else {
         /*
          * Vector of indices of the adjacent cells.
@@ -422,9 +418,8 @@ TRACCC_HOST_DEVICE inline void ccl_kernel(
         unsigned char adjc[details::CELLS_PER_THREAD_STACK_LIMIT];
 
         ccl_core(cfg, thread_id, partition_start, partition_end, f_primary,
-                 gf_primary, cell_links, adjv, adjc, cells_device, det_desc,
-                 det_cond, measurements_device, barrier, disjoint_set,
-                 cluster_size);
+                 gf_primary, adjv, adjc, cells_device, det_desc, det_cond,
+                 measurements_device, barrier, disjoint_set, cluster_size);
     }
 
     barrier.blockBarrier();

device/common/src/clusterization/clusterization_algorithm.cpp

@@ -115,10 +115,6 @@ clusterization_algorithm::execute_impl(
         num_cells, mr().main, vecmem::data::buffer_type::resizable};
     copy().setup(measurements)->ignore();
 
-    // Create buffer for linking cells to their measurements.
-    vecmem::data::vector_buffer<unsigned int> cell_links(num_cells, mr().main);
-    copy().setup(cell_links)->ignore();
-
     // Ensure that the chosen maximum cell count is compatible with the maximum
     // stack size.
     assert(m_config.max_cells_per_thread <=
@@ -134,9 +130,8 @@ clusterization_algorithm::execute_impl(
 
     // Launch the CCL kernel.
     ccl_kernel({num_cells, m_config, cells, det_descr, det_cond, measurements,
-                cell_links, m_f_backup, m_gf_backup, m_adjc_backup,
-                m_adjv_backup, m_backup_mutex.get(), disjoint_set,
-                cluster_sizes});
+                m_f_backup, m_gf_backup, m_adjc_backup, m_adjv_backup,
+                m_backup_mutex.get(), disjoint_set, cluster_sizes});
 
     std::optional<edm::silicon_cluster_collection::buffer> cluster_data =
         std::nullopt;

device/cuda/src/clusterization/clusterization_algorithm.cu

@@ -49,9 +49,9 @@ void clusterization_algorithm::ccl_kernel(
                               sizeof(device::details::index_t),
                           details::get_stream(stream())>>>(
         payload.config, payload.cells, payload.det_descr, payload.det_cond,
-        payload.measurements, payload.cell_links, payload.f_backup,
-        payload.gf_backup, payload.adjc_backup, payload.adjv_backup,
-        payload.backup_mutex, payload.disjoint_set, payload.cluster_sizes);
+        payload.measurements, payload.f_backup, payload.gf_backup,
+        payload.adjc_backup, payload.adjv_backup, payload.backup_mutex,
+        payload.disjoint_set, payload.cluster_sizes);
     TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 }
 

device/cuda/src/clusterization/kernels/ccl_kernel.cu

@@ -24,7 +24,6 @@ __global__ void ccl_kernel(
     const detector_design_description::const_view det_desc_view,
     const detector_conditions_description::const_view det_cond_view,
     edm::measurement_collection::view measurements_view,
-    vecmem::data::vector_view<unsigned int> cell_links,
     vecmem::data::vector_view<device::details::fallback_index_t> f_backup_view,
     vecmem::data::vector_view<device::details::fallback_index_t> gf_backup_view,
     vecmem::data::vector_view<unsigned char> adjc_backup_view,
@@ -54,7 +53,6 @@ __global__ void ccl_kernel(
                        partition_start, partition_end, outi, f_view, gf_view,
                        f_backup_view, gf_backup_view, adjc_backup_view,
                        adjv_backup_view, backup_mutex, disjoint_set_view,
-                       cluster_size_view, barry_r, measurements_view,
-                       cell_links);
+                       cluster_size_view, barry_r, measurements_view);
 }
 }  // namespace traccc::cuda::kernels

device/cuda/src/clusterization/kernels/ccl_kernel.cuh

@@ -27,7 +27,6 @@ __global__ void ccl_kernel(
     const detector_design_description::const_view det_descr_view,
     const detector_conditions_description::const_view det_cond_view,
     edm::measurement_collection::view measurements_view,
-    vecmem::data::vector_view<unsigned int> cell_links,
     vecmem::data::vector_view<device::details::fallback_index_t> f_backup_view,
     vecmem::data::vector_view<device::details::fallback_index_t> gf_backup_view,
     vecmem::data::vector_view<unsigned char> adjc_backup_view,