Simplify the API of traccc::device::form_spacepoints.

There is no need for it to receive the size of the measurement
collection separately. It already has a view of the collection
with that information.

Author: krasznaa

device/alpaka/src/seeding/spacepoint_formation_algorithm.cpp

@@ -23,15 +23,13 @@ struct FormSpacepointsKernel {
     ALPAKA_FN_ACC void operator()(
         TAcc const& acc, typename detector_t::view_type det_view,
         measurement_collection_types::const_view measurements_view,
-        const unsigned int measurement_count,
         edm::spacepoint_collection::view spacepoints_view) const {
 
         auto const globalThreadIdx =
             ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Threads>(acc)[0u];
 
         device::form_spacepoints<detector_t>(
-            globalThreadIdx, det_view, measurements_view, measurement_count,
-            spacepoints_view);
+            globalThreadIdx, det_view, measurements_view, spacepoints_view);
     }
 };
 
@@ -58,6 +56,7 @@ spacepoint_formation_algorithm<detector_t>::operator()(
     edm::spacepoint_collection::buffer spacepoints(
         num_measurements, m_mr.main, vecmem::data::buffer_type::resizable);
     m_copy.get().setup(spacepoints)->ignore();
+    edm::spacepoint_collection::view spacepoints_view{spacepoints};
 
     // If there are no measurements, we can conclude here.
     if (num_measurements == 0) {
@@ -71,8 +70,7 @@ spacepoint_formation_algorithm<detector_t>::operator()(
 
     // Launch the spacepoint formation kernel.
     ::alpaka::exec<Acc>(queue, workDiv, FormSpacepointsKernel<detector_t>{},
-                        det_view, measurements_view, num_measurements,
-                        vecmem::get_data(spacepoints));
+                        det_view, measurements_view, spacepoints_view);
     ::alpaka::wait(queue);
 
     // Return the reconstructed spacepoints.

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

@@ -23,14 +23,12 @@ namespace traccc::device {
 /// @param[in] globalIndex          The index for the current thread
 /// @param[in] det_view             A view type object of detector
 /// @param[in] measurements_view    Collection of measurements
-/// @param[in] measurement_count    Number of measurements
 /// @param[out] spacepoints_view    Collection of spacepoints
 ///
 template <typename detector_t>
 TRACCC_HOST_DEVICE inline void form_spacepoints(
     global_index_t globalIndex, typename detector_t::view_type det_view,
     const measurement_collection_types::const_view& measurements_view,
-    unsigned int measurement_count,
     edm::spacepoint_collection::view spacepoints_view);
 
 }  // namespace traccc::device

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

@@ -19,21 +19,21 @@ template <typename detector_t>
 TRACCC_HOST_DEVICE inline void form_spacepoints(
     const global_index_t globalIndex, typename detector_t::view_type det_view,
     const measurement_collection_types::const_view& measurements_view,
-    unsigned int measurement_count,
     edm::spacepoint_collection::view spacepoints_view) {
 
+    // Set up the input container(s).
+    const measurement_collection_types::const_device measurements(
+        measurements_view);
+
     // Check if anything needs to be done
-    if (globalIndex >= measurement_count) {
+    if (globalIndex >= measurements.size()) {
         return;
     }
 
     // Create the tracking geometry
     detector_t det(det_view);
 
-    // Get device copy of input/output parameters
-    const measurement_collection_types::const_device measurements(
-        measurements_view);
-    assert(measurements.size() == measurement_count);
+    // Set up the output container(s).
     edm::spacepoint_collection::device spacepoints(spacepoints_view);
 
     const auto& meas = measurements.at(globalIndex);

device/cuda/src/seeding/spacepoint_formation_algorithm.cu

@@ -22,12 +22,10 @@ template <typename detector_t>
 __global__ void __launch_bounds__(1024, 1)
     form_spacepoints(typename detector_t::view_type det_view,
                      measurement_collection_types::const_view measurements_view,
-                     const unsigned int measurement_count,
                      edm::spacepoint_collection::view spacepoints_view) {
 
     device::form_spacepoints<detector_t>(details::global_index1(), det_view,
-                                         measurements_view, measurement_count,
-                                         spacepoints_view);
+                                         measurements_view, spacepoints_view);
 }
 
 }  // namespace kernels
@@ -66,7 +64,7 @@ spacepoint_formation_algorithm<detector_t>::operator()(
 
     // Launch the spacepoint formation kernel.
     kernels::form_spacepoints<detector_t><<<nBlocks, blockSize, 0, stream>>>(
-        det_view, measurements_view, num_measurements, spacepoints);
+        det_view, measurements_view, spacepoints);
     TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
     // Return the reconstructed spacepoints.

device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp

@@ -65,14 +65,13 @@ edm::spacepoint_collection::buffer silicon_pixel_spacepoint_formation(
     // Run the spacepoint formation on the device.
     queue
         .submit([&](::sycl::handler& h) {
-            h.parallel_for(
-                countRange, [det_view, measurements_view, n_measurements,
-                             spacepoints_view = vecmem::get_data(result)](
-                                ::sycl::nd_item<1> item) {
-                    device::form_spacepoints<detector_t>(
-                        details::global_index(item), det_view,
-                        measurements_view, n_measurements, spacepoints_view);
-                });
+            h.parallel_for(countRange, [det_view, measurements_view,
+                                        spacepoints_view = vecmem::get_data(
+                                            result)](::sycl::nd_item<1> item) {
+                device::form_spacepoints<detector_t>(
+                    details::global_index(item), det_view, measurements_view,
+                    spacepoints_view);
+            });
         })
         .wait_and_throw();