Merge pull request #415 from guilhermeAlmeida1/fixCrashOnNoSeeds

Fix: Handle empty inputs gracefully

Author: krasznaa

device/cuda/src/clusterization/clusterization_algorithm.cu

@@ -91,6 +91,11 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()(
     const cell_collection_types::view::size_type num_cells =
         m_copy.get_size(cells);
 
+    if (num_cells == 0) {
+        return {output_type::first_type{0, m_mr.main},
+                output_type::second_type{0, m_mr.main}};
+    }
+
     // Create result object for the CCL kernel with size overestimation
     alt_measurement_collection_types::buffer measurements_buffer(num_cells,
                                                                  m_mr.main);

device/cuda/src/seeding/seed_finding.cu

@@ -164,18 +164,21 @@ seed_finding::output_type seed_finding::operator()(
     vecmem::data::vector_buffer sp_grid_prefix_sum_buff =
         make_prefix_sum_buff(grid_sizes, m_copy, m_mr, m_stream);
 
+    const auto num_spacepoints = m_copy.get_size(sp_grid_prefix_sum_buff);
+    if (num_spacepoints == 0) {
+        return {0, m_mr.main};
+    }
+
     // Set up the doublet counter buffer.
     device::doublet_counter_collection_types::buffer doublet_counter_buffer = {
-        m_copy.get_size(sp_grid_prefix_sum_buff), m_mr.main,
-        vecmem::data::buffer_type::resizable};
+        num_spacepoints, m_mr.main, vecmem::data::buffer_type::resizable};
     m_copy.setup(doublet_counter_buffer);
 
     // Calculate the number of threads and thread blocks to run the doublet
     // counting kernel for.
     const unsigned int nDoubletCountThreads = WARP_SIZE * 2;
     const unsigned int nDoubletCountBlocks =
-        (m_copy.get_size(sp_grid_prefix_sum_buff) + nDoubletCountThreads - 1) /
-        nDoubletCountThreads;
+        (num_spacepoints + nDoubletCountThreads - 1) / nDoubletCountThreads;
 
     // Counter for the total number of doublets and triplets
     vecmem::unique_alloc_ptr<device::seeding_global_counter>
@@ -205,6 +208,11 @@ seed_finding::output_type seed_finding::operator()(
                                      cudaMemcpyDeviceToHost, stream));
     m_stream.synchronize();
 
+    if (globalCounter_host->m_nMidBot == 0 ||
+        globalCounter_host->m_nMidTop == 0) {
+        return {0, m_mr.main};
+    }
+
     // Set up the doublet counter buffers.
     device::device_doublet_collection_types::buffer doublet_buffer_mb = {
         globalCounter_host->m_nMidBot, m_mr.main};
@@ -275,6 +283,10 @@ seed_finding::output_type seed_finding::operator()(
                                      cudaMemcpyDeviceToHost, stream));
     m_stream.synchronize();
 
+    if (globalCounter_host->m_nTriplets == 0) {
+        return {0, m_mr.main};
+    }
+
     // Set up the triplet buffer.
     device::device_triplet_collection_types::buffer triplet_buffer = {
         globalCounter_host->m_nTriplets, m_mr.main};

device/cuda/src/seeding/spacepoint_binning.cu

@@ -60,7 +60,11 @@ sp_grid_buffer spacepoint_binning::operator()(
     cudaStream_t stream = details::get_stream(m_stream);
 
     // Get the spacepoint sizes from the view
-    auto sp_size = m_copy.get_size(spacepoints_view);
+    const auto sp_size = m_copy.get_size(spacepoints_view);
+
+    if (sp_size == 0) {
+        return {m_axes.first, m_axes.second, {}, m_mr.main, m_mr.host};
+    }
 
     // Set up the container that will be filled with the required capacities for
     // the spacepoint grid.

device/cuda/src/utils/make_prefix_sum_buff.cu

@@ -38,6 +38,10 @@ vecmem::data::vector_buffer<device::prefix_sum_element_t> make_prefix_sum_buff(
         sizes_sum_view = make_sum_result.view;
     const unsigned int totalSize = make_sum_result.totalSize;
 
+    if (totalSize == 0) {
+        return {0, mr.main};
+    }
+
     // Create buffer and view objects for prefix sum vector
     vecmem::data::vector_buffer<device::prefix_sum_element_t> prefix_sum_buff(
         totalSize, mr.main);
@@ -65,6 +69,10 @@ vecmem::data::vector_buffer<device::prefix_sum_element_t> make_prefix_sum_buff(
         sizes_sum_view = make_sum_result.view;
     const unsigned int totalSize = make_sum_result.totalSize;
 
+    if (totalSize == 0) {
+        return {0, mr.main};
+    }
+
     // Create buffer and view objects for prefix sum vector
     vecmem::data::vector_buffer<device::prefix_sum_element_t> prefix_sum_buff(
         totalSize, mr.main);

device/sycl/src/clusterization/clusterization_algorithm.sycl

@@ -65,6 +65,11 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()(
     const cell_collection_types::view::size_type num_cells =
         m_copy.get_size(cells);
 
+    if (num_cells == 0) {
+        return {output_type::first_type{0, m_mr.main},
+                output_type::second_type{0, m_mr.main}};
+    }
+
     // Create result object for the CCL kernel with size overestimation
     alt_measurement_collection_types::buffer measurements_buffer(num_cells,
                                                                  m_mr.main);

device/sycl/src/seeding/seed_finding.sycl

@@ -82,10 +82,14 @@ seed_finding::output_type seed_finding::operator()(
     vecmem::data::vector_view<device::prefix_sum_element_t>
         sp_grid_prefix_sum_view = sp_grid_prefix_sum_buff;
 
+    const auto num_spacepoints = m_copy.get_size(sp_grid_prefix_sum_view);
+    if (num_spacepoints == 0) {
+        return {0, m_mr.main};
+    }
+
     // Set up the doublet counter buffer.
     device::doublet_counter_collection_types::buffer doublet_counter_buffer = {
-        m_copy.get_size(sp_grid_prefix_sum_view), m_mr.main,
-        vecmem::data::buffer_type::resizable};
+        num_spacepoints, m_mr.main, vecmem::data::buffer_type::resizable};
     m_copy.setup(doublet_counter_buffer)->wait();
 
     // Counter for the total number of doublets and triplets
@@ -102,7 +106,7 @@ seed_finding::output_type seed_finding::operator()(
     // Calculate the range to run the doublet counting for.
     static constexpr unsigned int doubletCountLocalSize = 32 * 2;
     auto doubletCountRange = traccc::sycl::calculate1DimNdRange(
-        m_copy.get_size(sp_grid_prefix_sum_view), doubletCountLocalSize);
+        num_spacepoints, doubletCountLocalSize);
 
     // Count the number of doublets that we need to produce.
     device::doublet_counter_collection_types::view doublet_counter_view =
@@ -135,6 +139,11 @@ seed_finding::output_type seed_finding::operator()(
                 sizeof(device::seeding_global_counter))
         .wait_and_throw();
 
+    if (globalCounter_host->m_nMidBot == 0 ||
+        globalCounter_host->m_nMidTop == 0) {
+        return {0, m_mr.main};
+    }
+
     // Set up the doublet buffers.
     device::device_doublet_collection_types::buffer doublet_buffer_mb = {
         globalCounter_host->m_nMidBot, m_mr.main};
@@ -232,6 +241,10 @@ seed_finding::output_type seed_finding::operator()(
                 sizeof(device::seeding_global_counter))
         .wait_and_throw();
 
+    if (globalCounter_host->m_nTriplets == 0) {
+        return {0, m_mr.main};
+    }
+
     // Set up the triplet buffer and its view
     device::device_triplet_collection_types::buffer triplet_buffer = {
         globalCounter_host->m_nTriplets, m_mr.main};

device/sycl/src/seeding/spacepoint_binning.sycl

@@ -44,7 +44,11 @@ sp_grid_buffer spacepoint_binning::operator()(
     const spacepoint_collection_types::const_view& spacepoints_view) const {
 
     // Get the spacepoint sizes from the view
-    auto sp_size = m_copy.get_size(spacepoints_view);
+    const auto sp_size = m_copy.get_size(spacepoints_view);
+
+    if (sp_size == 0) {
+        return {m_axes.first, m_axes.second, {}, m_mr.main, m_mr.host};
+    }
 
     // Set up the container that will be filled with the required capacities for
     // the spacepoint grid.

device/sycl/src/utils/make_prefix_sum_buff.sycl

@@ -26,6 +26,10 @@ vecmem::data::vector_buffer<device::prefix_sum_element_t> make_prefix_sum_buff(
         sizes_sum_view = make_sum_result.view;
     const unsigned int totalSize = make_sum_result.totalSize;
 
+    if (totalSize == 0) {
+        return {0, mr.main};
+    }
+
     // Create buffer and view objects for prefix sum vector
     vecmem::data::vector_buffer<device::prefix_sum_element_t> prefix_sum_buff(
         totalSize, mr.main);

extern/vecmem/CMakeLists.txt

@@ -18,7 +18,7 @@ message( STATUS "Building VecMem as part of the TRACCC project" )
 
 # Declare where to get VecMem from.
 set( TRACCC_VECMEM_SOURCE
-   "URL;https://github.com/acts-project/vecmem/archive/refs/tags/v0.25.0.tar.gz;URL_MD5;7fc14592c693c9853d27ca21ec528555"
+   "URL;https://github.com/acts-project/vecmem/archive/refs/tags/v0.26.0.tar.gz;URL_MD5;3162c26ced6bb7ce25a5873002e8a2b4"
    CACHE STRING "Source for VecMem, when built as part of this project" )
 mark_as_advanced( TRACCC_VECMEM_SOURCE )
 FetchContent_Declare( VecMem ${TRACCC_VECMEM_SOURCE} )