diff --git a/CMakeLists.txt b/CMakeLists.txt index 92bdc8d82b..e0189d7023 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -421,6 +421,25 @@ if( TRACCC_SETUP_INDICATORS ) endif() endif() +# Set up stdexec. +option( TRACCC_SETUP_STDEXEC + "Set up the stdexec target(s) explicitly" TRUE ) +if (TRACCC_USE_SYSTEM_LIBS OR TRACCC_USE_SPACK_LIBS) + set(TRACCC_USE_SYSTEM_STDEXEC_DEFAULT ON) +else() + set(TRACCC_USE_SYSTEM_STDEXEC_DEFAULT OFF) +endif() +option( TRACCC_USE_SYSTEM_STDEXEC + "Pick up an existing installation of stdexec from the build environment" + ${TRACCC_USE_SYSTEM_STDEXEC_DEFAULT} ) +unset(TRACCC_USE_SYSTEM_STDEXEC_DEFAULT) +if( TRACCC_SETUP_STDEXEC ) + if( TRACCC_USE_SYSTEM_STDEXEC ) + find_package( stdexec REQUIRED ) + else() + add_subdirectory( extern/stdexec ) + endif() +endif() option( TRACCC_ENABLE_NVTX_PROFILING "Use instrument functions to enable fine grained profiling" FALSE ) diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 17b2ba4a73..0da0c047ab 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -65,6 +65,8 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/utils/logging.hpp" "include/traccc/utils/prob.hpp" "src/utils/logging.cpp" + # Execution. + "include/traccc/execution/task.hpp" # Clusterization algorithmic code. "include/traccc/clusterization/details/sparse_ccl.hpp" "include/traccc/clusterization/impl/sparse_ccl.ipp" @@ -138,7 +140,7 @@ traccc_add_library( traccc_core core TYPE SHARED "src/ambiguity_resolution/legacy/greedy_ambiguity_resolution_algorithm.cpp") target_link_libraries( traccc_core PUBLIC Eigen3::Eigen vecmem::core covfie::core detray::core detray::detectors - traccc::algebra Acts::Core ) + traccc::algebra Acts::Core STDEXEC::stdexec ) # Prevent Eigen from getting confused when building code for a # CUDA or HIP backend with SYCL. diff --git a/core/include/traccc/execution/task.hpp b/core/include/traccc/execution/task.hpp new file mode 100644 index 0000000000..71bf2269b6 --- /dev/null +++ b/core/include/traccc/execution/task.hpp @@ -0,0 +1,10 @@ +#pragma once + +// Stdexec include(s). +#include + +namespace traccc { + +template +using task = exec::task; +} diff --git a/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index 8866f4373c..4e96372dde 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -36,6 +36,7 @@ traccc_add_library( traccc_device_common device_common "include/traccc/clusterization/device/ccl_kernel.hpp" "include/traccc/clusterization/device/sorting_index_filler.hpp" "include/traccc/clusterization/device/clusterization_algorithm.hpp" + "include/traccc/clusterization/device/clusterization_kernel_payload.hpp" "src/clusterization/clusterization_algorithm.cpp" # Spacepoint binning function(s). "include/traccc/seeding/device/count_grid_capacities.hpp" @@ -60,10 +61,13 @@ traccc_add_library( traccc_device_common device_common "include/traccc/seeding/device/select_seeds.hpp" "include/traccc/seeding/device/impl/select_seeds.ipp" "include/traccc/seeding/device/triplet_seeding_algorithm.hpp" + "include/traccc/seeding/device/triplet_seeding_kernel_payloads.hpp" "src/seeding/triplet_seeding_algorithm.cpp" "include/traccc/seeding/device/silicon_pixel_spacepoint_formation_algorithm.hpp" + "include/traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp" "src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp" "include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp" + "include/traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp" "src/seeding/seed_parameter_estimation_algorithm.cpp" # Track parameters estimation function(s). "include/traccc/seeding/device/estimate_track_params.hpp" diff --git a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp index 89636c7842..a8c0d349f1 100644 --- a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp +++ b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp @@ -9,8 +9,10 @@ // Local include(s). #include "traccc/clusterization/device/ccl_kernel_definitions.hpp" +#include "traccc/clusterization/device/clusterization_kernel_payload.hpp" #include "traccc/clusterization/device/tags.hpp" #include "traccc/device/algorithm_base.hpp" +#include "traccc/execution/task.hpp" // Project include(s). #include "traccc/clusterization/clustering_config.hpp" @@ -41,16 +43,18 @@ namespace traccc::device { /// synchronisation statement is required before destroying the buffer. /// class clusterization_algorithm - : public algorithm::buffer( - const edm::silicon_cell_collection::const_view&, - const silicon_detector_description::const_view&)>, - public algorithm::buffer( - const edm::silicon_cell_collection::const_view&, - const silicon_detector_description::const_view&, - clustering_discard_disjoint_set&&)>, + : public algorithm< + task::buffer>( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>, + public algorithm< + task::buffer>( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&, + clustering_discard_disjoint_set&&)>, public algorithm< - std::pair::buffer, - edm::silicon_cluster_collection::buffer>( + task::buffer, + edm::silicon_cluster_collection::buffer>>( const edm::silicon_cell_collection::const_view&, const silicon_detector_description::const_view&, clustering_keep_disjoint_set&&)>, @@ -81,18 +85,18 @@ class clusterization_algorithm /// @return a measurement collection (buffer) /// /// @{ - edm::measurement_collection::buffer operator()( + task::buffer> operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr) const override; - edm::measurement_collection::buffer operator()( + task::buffer> operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, clustering_discard_disjoint_set&&) const override; - std::pair::buffer, - edm::silicon_cluster_collection::buffer> + task::buffer, + edm::silicon_cluster_collection::buffer>> operator()(const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, clustering_keep_disjoint_set&&) const override; @@ -111,34 +115,7 @@ class clusterization_algorithm const edm::silicon_cell_collection::const_view& cells) const = 0; /// Payload for the @c ccl_kernel function - struct ccl_kernel_payload { - /// Number of cells in the event - unsigned int n_cells; - /// The clustering configuration - const config_type& config; - /// All cells in an event - const edm::silicon_cell_collection::const_view& cells; - /// The detector description - const silicon_detector_description::const_view& det_descr; - /// The measurement collection to fill - edm::measurement_collection::view& measurements; - /// Buffer for linking cells to measurements - vecmem::data::vector_view& cell_links; - /// Buffer for backup of the first element links - vecmem::data::vector_view& f_backup; - /// Buffer for backup of the group first element links - vecmem::data::vector_view& gf_backup; - /// Buffer for backup of the adjacency matrix (counts) - vecmem::data::vector_view& adjc_backup; - /// Buffer for backup of the adjacency matrix (values) - vecmem::data::vector_view& adjv_backup; - /// Mutex for the backup structures - unsigned int* backup_mutex; - /// Buffer for the disjoint set data structure - vecmem::data::vector_view& disjoint_set; - /// Buffer for the sizes of the clusters - vecmem::data::vector_view& cluster_sizes; - }; + using ccl_kernel_payload = clusterization_ccl_kernel_payload; /// Main CCL kernel launcher /// @@ -159,13 +136,13 @@ class clusterization_algorithm /// @} - /// Possibly suspend execution until all asynchronous operations are done - virtual void await(vecmem::abstract_event& event) const = 0; + /// Suspend execution until all asynchronous operations are done + virtual task await(vecmem::abstract_event& event) const = 0; private: /// Main algorithmic implementation of the clusterization algorithm - std::pair::buffer, - std::optional> + task::buffer, + std::optional>> execute_impl(const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, bool keep_disjoint_set) const; diff --git a/device/common/include/traccc/clusterization/device/clusterization_kernel_payload.hpp b/device/common/include/traccc/clusterization/device/clusterization_kernel_payload.hpp new file mode 100644 index 0000000000..833bbda0f1 --- /dev/null +++ b/device/common/include/traccc/clusterization/device/clusterization_kernel_payload.hpp @@ -0,0 +1,53 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2022-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/clusterization/clustering_config.hpp" +#include "traccc/clusterization/device/ccl_kernel_definitions.hpp" +#include "traccc/edm/measurement_collection.hpp" +#include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/geometry/silicon_detector_description.hpp" + +// VecMem include(s). +#include + +namespace traccc::device { + +/// Payload for the @c traccc::device::clusterization_algorithm::ccl_kernel +/// function. +struct clusterization_ccl_kernel_payload { + /// Number of cells in the event + unsigned int n_cells; + /// The clustering configuration + const clustering_config& config; + /// All cells in an event + const edm::silicon_cell_collection::const_view& cells; + /// The detector description + const silicon_detector_description::const_view& det_descr; + /// The measurement collection to fill + edm::measurement_collection::view& measurements; + /// Buffer for linking cells to measurements + vecmem::data::vector_view& cell_links; + /// Buffer for backup of the first element links + vecmem::data::vector_view& f_backup; + /// Buffer for backup of the group first element links + vecmem::data::vector_view& gf_backup; + /// Buffer for backup of the adjacency matrix (counts) + vecmem::data::vector_view& adjc_backup; + /// Buffer for backup of the adjacency matrix (values) + vecmem::data::vector_view& adjv_backup; + /// Mutex for the backup structures + unsigned int* backup_mutex; + /// Buffer for the disjoint set data structure + vecmem::data::vector_view& disjoint_set; + /// Buffer for the sizes of the clusters + vecmem::data::vector_view& cluster_sizes; +}; + +} // namespace traccc::device diff --git a/device/common/include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp b/device/common/include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp index 52a25329c6..da878c3ef9 100644 --- a/device/common/include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp +++ b/device/common/include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp @@ -9,6 +9,8 @@ // Local include(s). #include "traccc/device/algorithm_base.hpp" +#include "traccc/execution/task.hpp" +#include "traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp" // Project include(s) #include "traccc/bfield/magnetic_field.hpp" @@ -29,7 +31,7 @@ namespace traccc::device { /// synchronisation statement is required before destroying this buffer. /// struct seed_parameter_estimation_algorithm - : public algorithm( const magnetic_field&, const edm::measurement_collection::const_view&, const edm::spacepoint_collection::const_view&, @@ -59,7 +61,8 @@ struct seed_parameter_estimation_algorithm /// @param measurements All measurements of the event /// @param spacepoints All spacepoints of the event /// @param seeds The reconstructed track seeds of the event - /// @return A vector of bound track parameters for the seeds + /// @return A task returning a vector of bound track parameters for the + /// seeds /// output_type operator()( const magnetic_field& bfield, @@ -73,35 +76,20 @@ struct seed_parameter_estimation_algorithm /// @{ /// Payload for the @c estimate_seed_params_kernel function - struct estimate_seed_params_kernel_payload { - /// The number of seeds - edm::seed_collection::const_view::size_type n_seeds; - /// The track parameter estimation configuration - const track_params_estimation_config& config; - /// The magnetic field object - const magnetic_field& bfield; - /// All measurements of the event - const edm::measurement_collection::const_view& - measurements; - /// All spacepoints of the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The reconstructed track seeds of the event - const edm::seed_collection::const_view& seeds; - /// The output buffer for the bound track parameters - bound_track_parameters_collection_types::view& params; - }; + using estimate_seed_params_kernel_payload = + struct traccc::device::estimate_seed_params_kernel_payload; /// Seed parameter estimation kernel launcher /// /// @param payload The payload for the kernel /// virtual void estimate_seed_params_kernel( - const struct estimate_seed_params_kernel_payload& payload) const = 0; + const estimate_seed_params_kernel_payload& payload) const = 0; /// @} - /// Possibly suspend execution until all asynchronous operations are done - virtual void await(vecmem::abstract_event& event) const = 0; + /// Suspend execution until all asynchronous operations are done + virtual task await(vecmem::abstract_event& event) const = 0; private: /// Internal data type diff --git a/device/common/include/traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp b/device/common/include/traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp new file mode 100644 index 0000000000..c7fe389374 --- /dev/null +++ b/device/common/include/traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp @@ -0,0 +1,40 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s) +#include "traccc/bfield/magnetic_field.hpp" +#include "traccc/edm/measurement_collection.hpp" +#include "traccc/edm/seed_collection.hpp" +#include "traccc/edm/spacepoint_collection.hpp" +#include "traccc/edm/track_parameters.hpp" +#include "traccc/seeding/detail/track_params_estimation_config.hpp" + +namespace traccc::device { + +/// Payload for the +/// @c seed_parameter_estimation_algorithm::estimate_seed_params_kernel function +struct estimate_seed_params_kernel_payload { + /// The number of seeds + edm::seed_collection::const_view::size_type n_seeds; + /// The track parameter estimation configuration + const track_params_estimation_config& config; + /// The magnetic field object + const magnetic_field& bfield; + /// All measurements of the event + const edm::measurement_collection::const_view& + measurements; + /// All spacepoints of the event + const edm::spacepoint_collection::const_view& spacepoints; + /// The reconstructed track seeds of the event + const edm::seed_collection::const_view& seeds; + /// The output buffer for the bound track parameters + bound_track_parameters_collection_types::view& params; +}; + +} // namespace traccc::device diff --git a/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_algorithm.hpp b/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_algorithm.hpp index cc8074eddc..524a43d1b2 100644 --- a/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_algorithm.hpp +++ b/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_algorithm.hpp @@ -9,6 +9,8 @@ // Local include(s). #include "traccc/device/algorithm_base.hpp" +#include "traccc/execution/task.hpp" +#include "traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp" // Project include(s). #include "traccc/edm/measurement_collection.hpp" @@ -26,7 +28,7 @@ namespace traccc::device { /// measurements made on every detector module, into 3D spacepoint coordinates. /// class silicon_pixel_spacepoint_formation_algorithm - : public algorithm( const detector_buffer&, const edm::measurement_collection::const_view&)>, public messaging, @@ -48,7 +50,8 @@ class silicon_pixel_spacepoint_formation_algorithm /// /// @param det Detector object /// @param measurements A collection of measurements - /// @return A spacepoint buffer, with one spacepoint for every + /// @return A task returning a spacepoint buffer, with one spacepoint for + /// every /// silicon pixel measurement /// output_type operator()( @@ -61,18 +64,8 @@ class silicon_pixel_spacepoint_formation_algorithm /// @{ /// Payload for the @c form_spacepoints_kernel function - struct form_spacepoints_kernel_payload { - /// The number of measurements in the event - edm::measurement_collection::const_view::size_type - n_measurements; - /// The detector object - const detector_buffer& detector; - /// The input measurements - const edm::measurement_collection::const_view& - measurements; - /// The output spacepoints - edm::spacepoint_collection::view& spacepoints; - }; + using form_spacepoints_kernel_payload = + silicon_pixel_spacepoint_formation_kernel_payload; /// Launch the spacepoint formation kernel /// @@ -83,8 +76,8 @@ class silicon_pixel_spacepoint_formation_algorithm /// @} - /// Possibly suspend execution until all asynchronous operations are done - virtual void await(vecmem::abstract_event& event) const = 0; + /// Suspend execution until all asynchronous operations are done + virtual task await(vecmem::abstract_event& event) const = 0; }; // class silicon_pixel_spacepoint_formation_algorithm diff --git a/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp b/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp new file mode 100644 index 0000000000..9076bff872 --- /dev/null +++ b/device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp @@ -0,0 +1,34 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/edm/measurement_collection.hpp" +#include "traccc/edm/spacepoint_collection.hpp" +#include "traccc/geometry/detector_buffer.hpp" + +namespace traccc::device { + +/// Payload for the +/// @c +/// traccc::device::silicon_pixel_spacepoint_formation_algorithm::form_spacepoints_kernel +/// function. +struct silicon_pixel_spacepoint_formation_kernel_payload { + /// The number of measurements in the event + edm::measurement_collection::const_view::size_type + n_measurements; + /// The detector object + const detector_buffer& detector; + /// The input measurements + const edm::measurement_collection::const_view& + measurements; + /// The output spacepoints + edm::spacepoint_collection::view& spacepoints; +}; + +} // namespace traccc::device diff --git a/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp b/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp index 47aaed83f7..59cb2c8261 100644 --- a/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp +++ b/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp @@ -14,12 +14,14 @@ #include "traccc/edm/device/device_triplet.hpp" #include "traccc/edm/device/doublet_counter.hpp" #include "traccc/edm/device/triplet_counter.hpp" +#include "traccc/execution/task.hpp" // Project include(s). #include "traccc/edm/seed_collection.hpp" #include "traccc/edm/spacepoint_collection.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" +#include "traccc/seeding/device/triplet_seeding_kernel_payloads.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -35,7 +37,7 @@ namespace traccc::device { /// synchronisation statement is required before destroying this buffer. /// class triplet_seeding_algorithm - : public algorithm( const edm::spacepoint_collection::const_view&)>, public messaging, public algorithm_base { @@ -63,7 +65,8 @@ class triplet_seeding_algorithm /// Operator executing the algorithm. /// /// @param spacepoints is a view of all spacepoints in the event - /// @return the buffer of track seeds reconstructed from the spacepoints + /// @return A task returning the buffer of track seeds reconstructed from + /// the spacepoints /// output_type operator()(const edm::spacepoint_collection::const_view& spacepoints) const override; @@ -73,22 +76,8 @@ class triplet_seeding_algorithm /// @{ /// Payload for the @c count_grid_capacities_kernel function - struct count_grid_capacities_kernel_payload { - /// The number of spacepoints in the event - edm::spacepoint_collection::const_view::size_type n_spacepoints; - /// The seed finding configuration - const seedfinder_config& config; - /// The phi axis of the spacepoint grid - const traccc::details::spacepoint_grid_types::host::axis_p0_type& - phi_axis; - /// The z axis of the spacepoint grid - const traccc::details::spacepoint_grid_types::host::axis_p1_type& - z_axis; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The buffer to write the grid capacities into - vecmem::data::vector_view& grid_capacities; - }; + using count_grid_capacities_kernel_payload = + triplet_seeding_count_grid_capacities_kernel_payload; /// Spacepoint grid capacity counting kernel launcher /// @@ -98,18 +87,8 @@ class triplet_seeding_algorithm const count_grid_capacities_kernel_payload& payload) const = 0; /// Payload for the @c populate_grid_kernel function - struct populate_grid_kernel_payload { - /// The number of spacepoints in the event - edm::spacepoint_collection::const_view::size_type n_spacepoints; - /// The seed finding configuration - const seedfinder_config& config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The spacepoint grid to populate - traccc::details::spacepoint_grid_types::view& grid; - /// A prefix sum describing the grid contents - vecmem::data::vector_view& grid_prefix_sum; - }; + using populate_grid_kernel_payload = + triplet_seeding_populate_grid_kernel_payload; /// Spacepoint grid population kernel launcher /// @@ -119,25 +98,8 @@ class triplet_seeding_algorithm const populate_grid_kernel_payload& payload) const = 0; /// Payload for the @c count_doublets_kernel function - struct count_doublets_kernel_payload { - /// The number of spacepoints in the event - edm::spacepoint_collection::const_view::size_type n_spacepoints; - /// The seed finding configuration - const seedfinder_config& config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The populated spacepoint grid - const traccc::details::spacepoint_grid_types::const_view& grid; - /// A prefix sum describing the grid contents - const vecmem::data::vector_view& - grid_prefix_sum; - /// The doublet counter collection to fill - doublet_counter_collection_types::view& doublet_counter; - /// The number of middle-bottom doublets found - unsigned int& nMidBot; - /// The number of middle-top doublets found - unsigned int& nMidTop; - }; + using count_doublets_kernel_payload = + triplet_seeding_count_doublets_kernel_payload; /// Doublet counting kernel launcher /// @@ -147,23 +109,8 @@ class triplet_seeding_algorithm const count_doublets_kernel_payload& payload) const = 0; /// Payload for the @c find_doublets_kernel function - struct find_doublets_kernel_payload { - /// The number of doublets counted earlier - device::doublet_counter_collection_types::const_view::size_type - n_doublets; - /// The seed finding configuration - const seedfinder_config& config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The populated spacepoint grid - const traccc::details::spacepoint_grid_types::const_view& grid; - /// The doublet counter collection - const doublet_counter_collection_types::const_view& doublet_counter; - /// The middle-bottom doublet collection to fill - device_doublet_collection_types::view& mb_doublets; - /// The middle-top doublet collection to fill - device_doublet_collection_types::view& mt_doublets; - }; + using find_doublets_kernel_payload = + triplet_seeding_find_doublets_kernel_payload; /// Doublet finding kernel launcher /// @@ -173,26 +120,8 @@ class triplet_seeding_algorithm const find_doublets_kernel_payload& payload) const = 0; /// Payload for the @c count_triplets_kernel function - struct count_triplets_kernel_payload { - /// The number of middle-bottom doublets found earlier - unsigned int nMidBot; - /// The seed finding configuration - const seedfinder_config& config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The populated spacepoint grid - const traccc::details::spacepoint_grid_types::const_view& grid; - /// The doublet counter collection - const doublet_counter_collection_types::const_view& doublet_counter; - /// The middle-bottom doublet collection - const device_doublet_collection_types::const_view& mb_doublets; - /// The middle-top doublet collection - const device_doublet_collection_types::const_view& mt_doublets; - /// The triplet counter per middle spacepoint to fill - triplet_counter_spM_collection_types::view& spM_counter; - /// The triplet counter per middle-bottom doublet to fill - triplet_counter_collection_types::view& midBot_counter; - }; + using count_triplets_kernel_payload = + triplet_seeding_count_triplets_kernel_payload; /// Triplet counting kernel launcher /// @@ -202,17 +131,8 @@ class triplet_seeding_algorithm const count_triplets_kernel_payload& payload) const = 0; /// Payload for the @c triplet_counts_reduction_kernel function - struct triplet_counts_reduction_kernel_payload { - /// The number of doublets found earlier - device::doublet_counter_collection_types::const_view::size_type - n_doublets; - /// The doublet counter collection - const doublet_counter_collection_types::const_view& doublet_counter; - /// The triplet counter per middle spacepoint - triplet_counter_spM_collection_types::view& spM_counter; - /// The total number of triplets found - unsigned int& nTriplets; - }; + using triplet_counts_reduction_kernel_payload = + triplet_seeding_triplet_counts_reduction_kernel_payload; /// Triplet count reduction kernel launcher /// @@ -222,28 +142,8 @@ class triplet_seeding_algorithm const triplet_counts_reduction_kernel_payload& payload) const = 0; /// Payload for the @c find_triplets_kernel function - struct find_triplets_kernel_payload { - /// The number of middle-bottom doublets found earlier - unsigned int nMidBot; - /// The seed finding configuration - const seedfinder_config& finding_config; - /// The seed filtering configuration - const seedfilter_config& filter_config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The populated spacepoint grid - const traccc::details::spacepoint_grid_types::const_view& grid; - /// The doublet counter collection - const doublet_counter_collection_types::const_view& doublet_counter; - /// The middle-top doublet collection - const device_doublet_collection_types::const_view& mt_doublets; - /// The triplet counter per middle spacepoint - const triplet_counter_spM_collection_types::const_view& spM_tc; - /// The triplet counter per middle-bottom doublet - const triplet_counter_collection_types::const_view& midBot_tc; - /// The triplet collection to fill - device_triplet_collection_types::view& triplets; - }; + using find_triplets_kernel_payload = + triplet_seeding_find_triplets_kernel_payload; /// Triplet finding kernel launcher /// @@ -253,20 +153,8 @@ class triplet_seeding_algorithm const find_triplets_kernel_payload& payload) const = 0; /// Payload for the @c update_triplet_weights_kernel function - struct update_triplet_weights_kernel_payload { - /// The number of triplets found earlier - device_triplet_collection_types::const_view::size_type n_triplets; - /// The seed filtering configuration - const seedfilter_config& config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The triplet counter per middle spacepoint - const triplet_counter_spM_collection_types::const_view& spM_tc; - /// The triplet counter per middle-bottom doublet - const triplet_counter_collection_types::const_view& midBot_tc; - /// The triplet collection to update - device_triplet_collection_types::view& triplets; - }; + using update_triplet_weights_kernel_payload = + triplet_seeding_update_triplet_weights_kernel_payload; /// Triplet weight updater/filler kernel launcher /// @@ -276,27 +164,8 @@ class triplet_seeding_algorithm const update_triplet_weights_kernel_payload& payload) const = 0; /// Payload for the @c select_seeds_kernel function - struct select_seeds_kernel_payload { - /// The number of doublets found earlier - device::doublet_counter_collection_types::const_view::size_type - n_doublets; - /// The seed finding configuration - const seedfinder_config& finder_config; - /// The seed filtering configuration - const seedfilter_config& filter_config; - /// All spacepoints in the event - const edm::spacepoint_collection::const_view& spacepoints; - /// The populated spacepoint grid - const traccc::details::spacepoint_grid_types::const_view& grid; - /// The triplet counter per middle spacepoint - const triplet_counter_spM_collection_types::const_view& spM_tc; - /// The triplet counter per middle-bottom doublet - const triplet_counter_collection_types::const_view& midBot_tc; - /// The triplet collection - const device_triplet_collection_types::const_view& triplets; - /// The seed collection to fill - edm::seed_collection::view& seeds; - }; + using select_seeds_kernel_payload = + triplet_seeding_select_seeds_kernel_payload; /// Seed selection/filling kernel launcher /// @@ -307,8 +176,8 @@ class triplet_seeding_algorithm /// @} - /// Possibly suspend execution until all asynchronous operations are done - virtual void await(vecmem::abstract_event& event) const = 0; + /// Suspend execution until all asynchronous operations are done + virtual task await(vecmem::abstract_event& event) const = 0; private: /// Internal data type diff --git a/device/common/include/traccc/seeding/device/triplet_seeding_kernel_payloads.hpp b/device/common/include/traccc/seeding/device/triplet_seeding_kernel_payloads.hpp new file mode 100644 index 0000000000..d7f7350546 --- /dev/null +++ b/device/common/include/traccc/seeding/device/triplet_seeding_kernel_payloads.hpp @@ -0,0 +1,120 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Local include(s). +#include "traccc/device/prefix_sum_element.hpp" +#include "traccc/edm/device/device_doublet.hpp" +#include "traccc/edm/device/device_triplet.hpp" +#include "traccc/edm/device/doublet_counter.hpp" +#include "traccc/edm/device/triplet_counter.hpp" + +// Project include(s). +#include "traccc/edm/seed_collection.hpp" +#include "traccc/edm/spacepoint_collection.hpp" +#include "traccc/seeding/detail/seeding_config.hpp" +#include "traccc/seeding/detail/spacepoint_grid.hpp" + +// VecMem include(s). +#include + +namespace traccc::device { + +struct triplet_seeding_count_grid_capacities_kernel_payload { + edm::spacepoint_collection::const_view::size_type n_spacepoints; + const seedfinder_config& config; + const traccc::details::spacepoint_grid_types::host::axis_p0_type& phi_axis; + const traccc::details::spacepoint_grid_types::host::axis_p1_type& z_axis; + const edm::spacepoint_collection::const_view& spacepoints; + vecmem::data::vector_view& grid_capacities; +}; + +struct triplet_seeding_populate_grid_kernel_payload { + edm::spacepoint_collection::const_view::size_type n_spacepoints; + const seedfinder_config& config; + const edm::spacepoint_collection::const_view& spacepoints; + traccc::details::spacepoint_grid_types::view& grid; + vecmem::data::vector_view& grid_prefix_sum; +}; + +struct triplet_seeding_count_doublets_kernel_payload { + edm::spacepoint_collection::const_view::size_type n_spacepoints; + const seedfinder_config& config; + const edm::spacepoint_collection::const_view& spacepoints; + const traccc::details::spacepoint_grid_types::const_view& grid; + const vecmem::data::vector_view& + grid_prefix_sum; + doublet_counter_collection_types::view& doublet_counter; + unsigned int& nMidBot; + unsigned int& nMidTop; +}; + +struct triplet_seeding_find_doublets_kernel_payload { + device::doublet_counter_collection_types::const_view::size_type n_doublets; + const seedfinder_config& config; + const edm::spacepoint_collection::const_view& spacepoints; + const traccc::details::spacepoint_grid_types::const_view& grid; + const doublet_counter_collection_types::const_view& doublet_counter; + device_doublet_collection_types::view& mb_doublets; + device_doublet_collection_types::view& mt_doublets; +}; + +struct triplet_seeding_count_triplets_kernel_payload { + unsigned int nMidBot; + const seedfinder_config& config; + const edm::spacepoint_collection::const_view& spacepoints; + const traccc::details::spacepoint_grid_types::const_view& grid; + const doublet_counter_collection_types::const_view& doublet_counter; + const device_doublet_collection_types::const_view& mb_doublets; + const device_doublet_collection_types::const_view& mt_doublets; + triplet_counter_spM_collection_types::view& spM_counter; + triplet_counter_collection_types::view& midBot_counter; +}; + +struct triplet_seeding_triplet_counts_reduction_kernel_payload { + device::doublet_counter_collection_types::const_view::size_type n_doublets; + const doublet_counter_collection_types::const_view& doublet_counter; + triplet_counter_spM_collection_types::view& spM_counter; + unsigned int& nTriplets; +}; + +struct triplet_seeding_find_triplets_kernel_payload { + unsigned int nMidBot; + const seedfinder_config& finding_config; + const seedfilter_config& filter_config; + const edm::spacepoint_collection::const_view& spacepoints; + const traccc::details::spacepoint_grid_types::const_view& grid; + const doublet_counter_collection_types::const_view& doublet_counter; + const device_doublet_collection_types::const_view& mt_doublets; + const triplet_counter_spM_collection_types::const_view& spM_tc; + const triplet_counter_collection_types::const_view& midBot_tc; + device_triplet_collection_types::view& triplets; +}; + +struct triplet_seeding_update_triplet_weights_kernel_payload { + device_triplet_collection_types::const_view::size_type n_triplets; + const seedfilter_config& config; + const edm::spacepoint_collection::const_view& spacepoints; + const triplet_counter_spM_collection_types::const_view& spM_tc; + const triplet_counter_collection_types::const_view& midBot_tc; + device_triplet_collection_types::view& triplets; +}; + +struct triplet_seeding_select_seeds_kernel_payload { + device::doublet_counter_collection_types::const_view::size_type n_doublets; + const seedfinder_config& finder_config; + const seedfilter_config& filter_config; + const edm::spacepoint_collection::const_view& spacepoints; + const traccc::details::spacepoint_grid_types::const_view& grid; + const triplet_counter_spM_collection_types::const_view& spM_tc; + const triplet_counter_collection_types::const_view& midBot_tc; + const device_triplet_collection_types::const_view& triplets; + edm::seed_collection::view& seeds; +}; + +} // namespace traccc::device diff --git a/device/common/src/clusterization/clusterization_algorithm.cpp b/device/common/src/clusterization/clusterization_algorithm.cpp index 833deaaf6a..907d866ca6 100644 --- a/device/common/src/clusterization/clusterization_algorithm.cpp +++ b/device/common/src/clusterization/clusterization_algorithm.cpp @@ -32,42 +32,44 @@ clusterization_algorithm::clusterization_algorithm( ->wait(); } -edm::measurement_collection::buffer +task::buffer> clusterization_algorithm::operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr) const { - return this->operator()(cells, det_descr, - clustering_discard_disjoint_set{}); + co_return co_await this->operator()(cells, det_descr, + clustering_discard_disjoint_set{}); } -edm::measurement_collection::buffer +task::buffer> clusterization_algorithm::operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, clustering_discard_disjoint_set&&) const { static constexpr bool KEEP_DISJOINT_SET = false; - auto [res, djs] = this->execute_impl(cells, det_descr, KEEP_DISJOINT_SET); + auto [res, djs] = + co_await this->execute_impl(cells, det_descr, KEEP_DISJOINT_SET); assert(!djs.has_value()); - return std::move(res); + co_return std::move(res); } -std::pair::buffer, - edm::silicon_cluster_collection::buffer> +task::buffer, + edm::silicon_cluster_collection::buffer>> clusterization_algorithm::operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, clustering_keep_disjoint_set&&) const { static constexpr bool KEEP_DISJOINT_SET = true; - auto [res, djs] = this->execute_impl(cells, det_descr, KEEP_DISJOINT_SET); + auto [res, djs] = + co_await this->execute_impl(cells, det_descr, KEEP_DISJOINT_SET); assert(djs.has_value()); - return {std::move(res), std::move(*djs)}; + co_return {std::move(res), std::move(*djs)}; } -std::pair::buffer, - std::optional> +task::buffer, + std::optional>> clusterization_algorithm::execute_impl( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, @@ -82,7 +84,7 @@ clusterization_algorithm::execute_impl( vecmem::async_size size = copy().get_size(cells, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); num_cells = size.unsafe_get(); } else { num_cells = copy().get_size(cells); @@ -91,10 +93,10 @@ clusterization_algorithm::execute_impl( // If there are no cells, return right away. if (num_cells == 0) { if (keep_disjoint_set) { - return {edm::measurement_collection::buffer{}, - edm::silicon_cluster_collection::buffer{}}; + co_return {edm::measurement_collection::buffer{}, + edm::silicon_cluster_collection::buffer{}}; } else { - return {}; + co_return {}; } } @@ -140,7 +142,7 @@ clusterization_algorithm::execute_impl( copy().get_size(measurements, *(mr().host)); // Here we could give control back to the caller, once our code // allows for it. (coroutines...)<-WIP - await(size); + co_await await(size); num_measurements = size.unsafe_get(); } else { num_measurements = copy().get_size(measurements); @@ -165,7 +167,7 @@ clusterization_algorithm::execute_impl( } // Return the reconstructed measurements. - return {std::move(measurements), std::move(cluster_data)}; + co_return {std::move(measurements), std::move(cluster_data)}; } } // namespace traccc::device diff --git a/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp b/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp index 7f8acaf719..b98b3d1ea5 100644 --- a/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp +++ b/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp @@ -41,7 +41,7 @@ auto seed_parameter_estimation_algorithm::operator()( vecmem::async_size size = copy().get_size(seeds, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); n_seeds = size.unsafe_get(); } else { n_seeds = copy().get_size(seeds); @@ -49,7 +49,7 @@ auto seed_parameter_estimation_algorithm::operator()( // If there are no seeds, return right away. if (n_seeds == 0) { - return {}; + co_return {}; } // Set up the output buffer. @@ -61,7 +61,7 @@ auto seed_parameter_estimation_algorithm::operator()( measurements, spacepoints, seeds, result}); // Return the result. - return result; + co_return result; } } // namespace traccc::device diff --git a/device/common/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp b/device/common/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp index 4d99633182..d36b1775b8 100644 --- a/device/common/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp +++ b/device/common/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -28,7 +28,7 @@ auto silicon_pixel_spacepoint_formation_algorithm::operator()( vecmem::async_size size = copy().get_size(measurements, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); n_measurements = size.unsafe_get(); } else { n_measurements = copy().get_size(measurements); @@ -36,7 +36,7 @@ auto silicon_pixel_spacepoint_formation_algorithm::operator()( // If there are no measurements, return right away. if (n_measurements == 0) { - return {}; + co_return {}; } // Create the result buffer. @@ -48,7 +48,7 @@ auto silicon_pixel_spacepoint_formation_algorithm::operator()( form_spacepoints_kernel({n_measurements, det, measurements, spacepoints}); // Return the reconstructed spacepoints. - return spacepoints; + co_return spacepoints; } } // namespace traccc::device diff --git a/device/common/src/seeding/triplet_seeding_algorithm.cpp b/device/common/src/seeding/triplet_seeding_algorithm.cpp index 5550486038..e105237739 100644 --- a/device/common/src/seeding/triplet_seeding_algorithm.cpp +++ b/device/common/src/seeding/triplet_seeding_algorithm.cpp @@ -66,7 +66,7 @@ auto triplet_seeding_algorithm::operator()( vecmem::async_size size = copy().get_size(spacepoints, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); n_spacepoints = size.unsafe_get(); } else { n_spacepoints = copy().get_size(spacepoints); @@ -74,7 +74,7 @@ auto triplet_seeding_algorithm::operator()( // If there are no spacepoints, return right away. if (n_spacepoints == 0) { - return {}; + co_return {}; } // Set up the container that will be filled with the required capacities for @@ -119,7 +119,7 @@ auto triplet_seeding_algorithm::operator()( copy().get_size(grid_prefix_sum_buffer, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); n_spacepoints = size.unsafe_get(); } else { n_spacepoints = copy().get_size(grid_prefix_sum_buffer); @@ -154,7 +154,7 @@ auto triplet_seeding_algorithm::operator()( copy().get_size(doublet_counter_buffer, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); n_doublets = size.unsafe_get(); } else { n_doublets = copy().get_size(doublet_counter_buffer); @@ -172,7 +172,7 @@ auto triplet_seeding_algorithm::operator()( // Exit already here if we won't find any triplets anyway. if ((globalCounter_host->m_nMidBot == 0) || (globalCounter_host->m_nMidTop == 0)) { - return {}; + co_return {}; } // Set up the doublet buffers. @@ -216,7 +216,7 @@ auto triplet_seeding_algorithm::operator()( copy().get_size(triplet_counter_midBot_buffer, *(mr().host)); // Here we could give control back to the caller, once our code allows // for it. (coroutines...)<-WIP - await(size); + co_await await(size); n_midBotTriplets = size.unsafe_get(); } else { n_midBotTriplets = copy().get_size(triplet_counter_midBot_buffer); @@ -229,7 +229,7 @@ auto triplet_seeding_algorithm::operator()( // If no triplets could be found, exit already here. if (globalCounter_host->m_nTriplets == 0) { - return {}; + co_return {}; } // Set up the triplet buffer. @@ -262,7 +262,7 @@ auto triplet_seeding_algorithm::operator()( triplet_counter_midBot_buffer, triplet_buffer, seed_buffer}); // Return the seed buffer. - return seed_buffer; + co_return seed_buffer; } } // namespace traccc::device diff --git a/device/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index 55c5697a31..3b7187306a 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -35,11 +35,17 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED "src/sanity/ordered_on.cuh" # Seed finding code. "include/traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp" - "src/seeding/seed_parameter_estimation_algorithm.cu" + "src/seeding/seed_parameter_estimation_algorithm.cpp" + "src/seeding/seed_parameter_estimation_kernel.cu" + "src/seeding/seed_parameter_estimation_kernel.hpp" "include/traccc/cuda/seeding/triplet_seeding_algorithm.hpp" - "src/seeding/triplet_seeding_algorithm.cu" + "src/seeding/triplet_seeding_algorithm.cpp" + "src/seeding/triplet_seeding_kernel.cu" + "src/seeding/triplet_seeding_kernel.hpp" "include/traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" - "src/seeding/silicon_pixel_spacepoint_formation_algorithm.cu" + "src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp" + "src/seeding/silicon_pixel_spacepoint_formation_kernel.cu" + "src/seeding/silicon_pixel_spacepoint_formation_kernel.hpp" #GBTS seed finding code "include/traccc/cuda/gbts_seeding/gbts_seeding_algorithm.hpp" "src/gbts_seeding/gbts_seeding_algorithm.cu" @@ -48,7 +54,9 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED "src/gbts_seeding/kernels/GbtsGraphProcessingKernels.cuh" # Clusterization "include/traccc/cuda/clusterization/clusterization_algorithm.hpp" - "src/clusterization/clusterization_algorithm.cu" + "src/clusterization/clusterization_algorithm.cpp" + "src/clusterization/clusterization_kernel.cu" + "src/clusterization/clusterization_kernel.hpp" "include/traccc/cuda/clusterization/measurement_sorting_algorithm.hpp" "src/clusterization/measurement_sorting_algorithm.cu" "src/clusterization/kernels/ccl_kernel.cu" diff --git a/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp b/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp index aef6b0a730..455204392a 100644 --- a/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp @@ -71,7 +71,7 @@ class clusterization_algorithm : public device::clusterization_algorithm, /// @} - void await(vecmem::abstract_event& event) const override; + task await(vecmem::abstract_event& event) const override; private: await_function_t m_await_function; diff --git a/device/cuda/include/traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp b/device/cuda/include/traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp index bb96f9fafc..35c7d34c22 100644 --- a/device/cuda/include/traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp @@ -12,6 +12,7 @@ // Project include(s). #include "traccc/seeding/device/seed_parameter_estimation_algorithm.hpp" +#include "traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp" namespace traccc::cuda { @@ -48,12 +49,11 @@ struct seed_parameter_estimation_algorithm /// @param payload The payload for the kernel /// void estimate_seed_params_kernel( - const struct estimate_seed_params_kernel_payload& payload) - const override; + const estimate_seed_params_kernel_payload& payload) const override; /// @} - void await(vecmem::abstract_event& event) const override; + task await(vecmem::abstract_event& event) const override; private: await_function_t m_await_function; diff --git a/device/cuda/include/traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp b/device/cuda/include/traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp index 6602c9dd45..eaba311c20 100644 --- a/device/cuda/include/traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp @@ -54,7 +54,7 @@ class silicon_pixel_spacepoint_formation_algorithm /// @} - void await(vecmem::abstract_event& event) const override; + task await(vecmem::abstract_event& event) const override; private: await_function_t m_await_function; diff --git a/device/cuda/include/traccc/cuda/seeding/triplet_seeding_algorithm.hpp b/device/cuda/include/traccc/cuda/seeding/triplet_seeding_algorithm.hpp index e1d5458c63..b394dfc9e4 100644 --- a/device/cuda/include/traccc/cuda/seeding/triplet_seeding_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/seeding/triplet_seeding_algorithm.hpp @@ -110,7 +110,7 @@ class triplet_seeding_algorithm : public device::triplet_seeding_algorithm, /// @} - void await(vecmem::abstract_event& event) const override; + task await(vecmem::abstract_event& event) const override; private: await_function_t m_await_function; diff --git a/device/cuda/include/traccc/cuda/utils/algorithm_base.hpp b/device/cuda/include/traccc/cuda/utils/algorithm_base.hpp index fad3a774e7..fea0293e03 100644 --- a/device/cuda/include/traccc/cuda/utils/algorithm_base.hpp +++ b/device/cuda/include/traccc/cuda/utils/algorithm_base.hpp @@ -9,6 +9,7 @@ // Local include(s). #include "traccc/cuda/utils/stream.hpp" +#include "traccc/execution/task.hpp" // VecMem include(s). #include @@ -45,16 +46,16 @@ class algorithm_base { }; // class algorithm_base using await_function_t = - std::function; - -// Default await function, same as await_event_sync -void default_await_function(const cuda::stream& stream, + std::function(const cuda::stream&, vecmem::abstract_event&)>; + +// Default await coroutine same as await_event_sync +task default_await_function(const cuda::stream& stream, + vecmem::abstract_event& event); +// Await coroutine that synchronizes the stream +task await_stream_sync(const cuda::stream& stream, + vecmem::abstract_event& event); +// Await coroutine that waits on the event +task await_event_sync(const cuda::stream& stream, vecmem::abstract_event& event); -// Await function that synchronizes the stream -void await_stream_sync(const cuda::stream& stream, - vecmem::abstract_event& event); -// Await function that waits on the event -void await_event_sync(const cuda::stream& stream, - vecmem::abstract_event& event); } // namespace traccc::cuda diff --git a/device/cuda/src/clusterization/clusterization_algorithm.cpp b/device/cuda/src/clusterization/clusterization_algorithm.cpp new file mode 100644 index 0000000000..28e68183b9 --- /dev/null +++ b/device/cuda/src/clusterization/clusterization_algorithm.cpp @@ -0,0 +1,48 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2022-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "traccc/cuda/clusterization/clusterization_algorithm.hpp" + +#include "../utils/utils.hpp" +#include "clusterization_kernel.hpp" + +namespace traccc::cuda { + +clusterization_algorithm::clusterization_algorithm( + const traccc::memory_resource& mr, vecmem::copy& copy, cuda::stream& str, + const config_type& config, std::unique_ptr logger, + await_function_t await_func) + : device::clusterization_algorithm(mr, copy, config, std::move(logger)), + cuda::algorithm_base(str), + m_await_function(await_func) {} + +bool clusterization_algorithm::input_is_valid( + const edm::silicon_cell_collection::const_view& cells) const { + return input_is_valid_on_device(mr().main, copy(), stream(), cells); +} + +void clusterization_algorithm::ccl_kernel( + const ccl_kernel_payload& payload) const { + launch_ccl_kernel(payload, details::get_stream(stream())); +} + +void clusterization_algorithm::cluster_maker_kernel( + unsigned int num_cells, + const vecmem::data::vector_view& disjoint_set, + edm::silicon_cluster_collection::view& cluster_data) const { + launch_reify_cluster_data_kernel(num_cells, disjoint_set, cluster_data, + details::get_stream(stream()), + warp_size()); +} + +task clusterization_algorithm::await( + vecmem::abstract_event& event) const { + co_await m_await_function(stream(), event); +} + +} // namespace traccc::cuda diff --git a/device/cuda/src/clusterization/clusterization_algorithm.cu b/device/cuda/src/clusterization/clusterization_kernel.cu similarity index 51% rename from device/cuda/src/clusterization/clusterization_algorithm.cu rename to device/cuda/src/clusterization/clusterization_kernel.cu index 90da7d746f..e531f55bdd 100644 --- a/device/cuda/src/clusterization/clusterization_algorithm.cu +++ b/device/cuda/src/clusterization/clusterization_kernel.cu @@ -5,43 +5,33 @@ * Mozilla Public License Version 2.0 */ -// CUDA Library include(s). +// Local include(s). #include "../sanity/contiguous_on.cuh" #include "../sanity/ordered_on.cuh" #include "../utils/cuda_error_handling.hpp" #include "./kernels/ccl_kernel.cuh" #include "./kernels/reify_cluster_data.cuh" -#include "traccc/clusterization/device/ccl_kernel_definitions.hpp" -#include "traccc/cuda/clusterization/clusterization_algorithm.hpp" +#include "clusterization_kernel.hpp" + +// Project include(s). #include "traccc/utils/projections.hpp" #include "traccc/utils/relations.hpp" -// Vecmem include(s). -#include -#include - namespace traccc::cuda { -clusterization_algorithm::clusterization_algorithm( - const traccc::memory_resource& mr, vecmem::copy& copy, cuda::stream& str, - const config_type& config, std::unique_ptr logger, - await_function_t await_func) - : device::clusterization_algorithm(mr, copy, config, std::move(logger)), - cuda::algorithm_base(str), - m_await_function(await_func) {} - -bool clusterization_algorithm::input_is_valid( - const edm::silicon_cell_collection::const_view& cells) const { +bool input_is_valid_on_device( + vecmem::memory_resource& mr, const vecmem::copy& copy, stream& stream, + const edm::silicon_cell_collection::const_view& cells) { return (is_contiguous_on( - cell_module_projection(), mr().main, copy(), stream(), cells) && + cell_module_projection(), mr, copy, stream, cells) && is_ordered_on( - channel0_major_cell_order_relation(), mr().main, copy(), - stream(), cells)); + channel0_major_cell_order_relation(), mr, copy, stream, cells)); } -void clusterization_algorithm::ccl_kernel( - const ccl_kernel_payload& payload) const { +void launch_ccl_kernel( + const traccc::device::clusterization_ccl_kernel_payload& payload, + cudaStream_t stream) { const unsigned int num_blocks = (payload.n_cells + (payload.config.target_partition_size()) - 1) / @@ -49,7 +39,7 @@ void clusterization_algorithm::ccl_kernel( kernels::ccl_kernel<<>>( + stream>>>( payload.config, payload.cells, payload.det_descr, payload.measurements, payload.cell_links, payload.f_backup, payload.gf_backup, payload.adjc_backup, payload.adjv_backup, payload.backup_mutex, @@ -57,21 +47,17 @@ void clusterization_algorithm::ccl_kernel( TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void clusterization_algorithm::cluster_maker_kernel( +void launch_reify_cluster_data_kernel( unsigned int num_cells, const vecmem::data::vector_view& disjoint_set, - edm::silicon_cluster_collection::view& cluster_data) const { + edm::silicon_cluster_collection::view& cluster_data, cudaStream_t stream, + unsigned int warp_size) { - const unsigned int num_threads = warp_size() * 16u; + const unsigned int num_threads = warp_size * 16u; const unsigned int num_blocks = (num_cells + num_threads - 1) / num_threads; - kernels::reify_cluster_data<<>>( + kernels::reify_cluster_data<<>>( disjoint_set, cluster_data); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void clusterization_algorithm::await(vecmem::abstract_event& event) const { - m_await_function(stream(), event); -} - } // namespace traccc::cuda diff --git a/device/cuda/src/clusterization/clusterization_kernel.hpp b/device/cuda/src/clusterization/clusterization_kernel.hpp new file mode 100644 index 0000000000..7b0693d062 --- /dev/null +++ b/device/cuda/src/clusterization/clusterization_kernel.hpp @@ -0,0 +1,44 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2022-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/clusterization/device/clusterization_kernel_payload.hpp" +#include "traccc/cuda/utils/stream.hpp" +#include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/edm/silicon_cluster_collection.hpp" + +// VecMem include(s). +#include +#include +#include + +// CUDA include(s). +#include + +namespace traccc::cuda { + +/// Host-side interface for the sanity check whether the input cells are valid +/// for clusterization. +bool input_is_valid_on_device( + vecmem::memory_resource& mr, const vecmem::copy& copy, stream& stream, + const edm::silicon_cell_collection::const_view& cells); + +/// Host-side interface for the main CCL kernel. +void launch_ccl_kernel( + const traccc::device::clusterization_ccl_kernel_payload& payload, + cudaStream_t stream); + +/// Host-side interface for the cluster reification kernel. +void launch_reify_cluster_data_kernel( + unsigned int num_cells, + const vecmem::data::vector_view& disjoint_set, + edm::silicon_cluster_collection::view& cluster_data, cudaStream_t stream, + unsigned int warp_size); + +} // namespace traccc::cuda diff --git a/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp b/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp new file mode 100644 index 0000000000..b29f22851d --- /dev/null +++ b/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp @@ -0,0 +1,39 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp" + +#include "../utils/utils.hpp" +#include "seed_parameter_estimation_kernel.hpp" + +// Project include(s). +#include "traccc/seeding/device/estimate_track_params.hpp" + +namespace traccc::cuda { + +seed_parameter_estimation_algorithm::seed_parameter_estimation_algorithm( + const track_params_estimation_config& config, + const traccc::memory_resource& mr, vecmem::copy& copy, cuda::stream& str, + std::unique_ptr logger, await_function_t await_func) + : device::seed_parameter_estimation_algorithm(config, mr, copy, + std::move(logger)), + cuda::algorithm_base(str), + m_await_function(await_func) {} + +void seed_parameter_estimation_algorithm::estimate_seed_params_kernel( + const estimate_seed_params_kernel_payload& payload) const { + + launch_estimate_track_params_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +task seed_parameter_estimation_algorithm::await( + vecmem::abstract_event& event) const { + co_await m_await_function(stream(), event); +} +} // namespace traccc::cuda diff --git a/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cu b/device/cuda/src/seeding/seed_parameter_estimation_kernel.cu similarity index 64% rename from device/cuda/src/seeding/seed_parameter_estimation_algorithm.cu rename to device/cuda/src/seeding/seed_parameter_estimation_kernel.cu index f5521171ef..39dd6baf90 100644 --- a/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cu +++ b/device/cuda/src/seeding/seed_parameter_estimation_kernel.cu @@ -10,7 +10,7 @@ #include "../utils/global_index.hpp" #include "../utils/magnetic_field_types.hpp" #include "../utils/utils.hpp" -#include "traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp" +#include "seed_parameter_estimation_kernel.hpp" // Project include(s). #include "traccc/seeding/device/estimate_track_params.hpp" @@ -34,33 +34,19 @@ __global__ void estimate_track_params( } // namespace kernels -seed_parameter_estimation_algorithm::seed_parameter_estimation_algorithm( - const track_params_estimation_config& config, - const traccc::memory_resource& mr, vecmem::copy& copy, cuda::stream& str, - std::unique_ptr logger, await_function_t await_func) - : device::seed_parameter_estimation_algorithm(config, mr, copy, - std::move(logger)), - cuda::algorithm_base(str), - m_await_function(await_func) {} - -void seed_parameter_estimation_algorithm::estimate_seed_params_kernel( - const struct estimate_seed_params_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 4; +void launch_estimate_track_params_kernel( + const traccc::device::estimate_seed_params_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 4; const unsigned int n_blocks = (payload.n_seeds + n_threads - 1) / n_threads; magnetic_field_visitor>( payload.bfield, [&](const bfield_view_t& bfield) { kernels::estimate_track_params - <<>>( + <<>>( payload.config, payload.measurements, payload.spacepoints, payload.seeds, bfield, payload.params); }); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } - -void seed_parameter_estimation_algorithm::await( - vecmem::abstract_event& event) const { - m_await_function(stream(), event); -} } // namespace traccc::cuda diff --git a/device/cuda/src/seeding/seed_parameter_estimation_kernel.hpp b/device/cuda/src/seeding/seed_parameter_estimation_kernel.hpp new file mode 100644 index 0000000000..21915b80f6 --- /dev/null +++ b/device/cuda/src/seeding/seed_parameter_estimation_kernel.hpp @@ -0,0 +1,29 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). + +// Project include(s). +#include "traccc/seeding/device/estimate_track_params.hpp" +#include "traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp" + +// CUDA include(s). +#include + +namespace traccc::cuda { + +/// Host-side interface for seed parameter estimation kernel +/// +/// @param payload The payload for the kernel +/// @param stream The CUDA stream to launch the kernel in +/// @param warp_size The warp size of the GPU being used +/// +void launch_estimate_track_params_kernel( + const traccc::device::estimate_seed_params_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size); + +} // namespace traccc::cuda diff --git a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp new file mode 100644 index 0000000000..670aedb069 --- /dev/null +++ b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -0,0 +1,38 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" + +#include "../utils/utils.hpp" +#include "silicon_pixel_spacepoint_formation_kernel.hpp" + +namespace traccc::cuda { + +silicon_pixel_spacepoint_formation_algorithm:: + silicon_pixel_spacepoint_formation_algorithm( + const traccc::memory_resource& mr, vecmem::copy& copy, + cuda::stream& str, std::unique_ptr logger, + await_function_t await_func) + : device::silicon_pixel_spacepoint_formation_algorithm(mr, copy, + std::move(logger)), + cuda::algorithm_base(str), + m_await_function(await_func) {} + +void silicon_pixel_spacepoint_formation_algorithm::form_spacepoints_kernel( + const form_spacepoints_kernel_payload& payload) const { + + launch_form_spacepoints_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +task silicon_pixel_spacepoint_formation_algorithm::await( + vecmem::abstract_event& event) const { + co_await m_await_function(stream(), event); +} + +} // namespace traccc::cuda diff --git a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cu b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.cu similarity index 57% rename from device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cu rename to device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.cu index a0894ee3b2..196eee12ce 100644 --- a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cu +++ b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.cu @@ -8,8 +8,7 @@ // Local include(s). #include "../utils/cuda_error_handling.hpp" #include "../utils/global_index.hpp" -#include "../utils/utils.hpp" -#include "traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" +#include "silicon_pixel_spacepoint_formation_kernel.hpp" // Project include(s). #include "traccc/geometry/detector.hpp" @@ -33,35 +32,22 @@ __global__ void __launch_bounds__(1024, 1) form_spacepoints( } // namespace kernels -silicon_pixel_spacepoint_formation_algorithm:: - silicon_pixel_spacepoint_formation_algorithm( - const traccc::memory_resource& mr, vecmem::copy& copy, - cuda::stream& str, std::unique_ptr logger, - await_function_t await_func) - : device::silicon_pixel_spacepoint_formation_algorithm(mr, copy, - std::move(logger)), - cuda::algorithm_base(str), - m_await_function(await_func) {} +void launch_form_spacepoints_kernel( + const traccc::device::silicon_pixel_spacepoint_formation_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size) { -void silicon_pixel_spacepoint_formation_algorithm::form_spacepoints_kernel( - const form_spacepoints_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 8; + const unsigned int n_threads = warp_size * 8; const unsigned int n_blocks = (payload.n_measurements + n_threads - 1) / n_threads; detector_buffer_visitor( payload.detector, [&]( const typename detector_traits_t::view& det) { kernels::form_spacepoints - <<>>( - det, payload.measurements, payload.spacepoints); + <<>>(det, payload.measurements, + payload.spacepoints); }); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void silicon_pixel_spacepoint_formation_algorithm::await( - vecmem::abstract_event& event) const { - m_await_function(stream(), event); -} - } // namespace traccc::cuda diff --git a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.hpp b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.hpp new file mode 100644 index 0000000000..d424170a38 --- /dev/null +++ b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.hpp @@ -0,0 +1,29 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp" + +// CUDA include(s). +#include + +namespace traccc::cuda { + +/// Host-side interface for silicon pixel spacepoint formation kernel. +/// +/// @param payload The payload for the kernel +/// @param stream The CUDA stream to launch the kernel in +/// @param warp_size The warp size of the GPU being used +/// +void launch_form_spacepoints_kernel( + const traccc::device::silicon_pixel_spacepoint_formation_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size); + +} // namespace traccc::cuda diff --git a/device/cuda/src/seeding/triplet_seeding_algorithm.cpp b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp new file mode 100644 index 0000000000..692d07f307 --- /dev/null +++ b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp @@ -0,0 +1,96 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "traccc/cuda/seeding/triplet_seeding_algorithm.hpp" + +#include "../utils/utils.hpp" +#include "triplet_seeding_kernel.hpp" + +namespace traccc::cuda { + +triplet_seeding_algorithm::triplet_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config, const traccc::memory_resource& mr, + vecmem::copy& copy, cuda::stream& str, std::unique_ptr logger, + await_function_t await_func) + : device::triplet_seeding_algorithm(finder_config, grid_config, + filter_config, mr, copy, + std::move(logger)), + cuda::algorithm_base{str}, + m_await_function(await_func) {} + +void triplet_seeding_algorithm::count_grid_capacities_kernel( + const count_grid_capacities_kernel_payload& payload) const { + + launch_count_grid_capacities_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::populate_grid_kernel( + const populate_grid_kernel_payload& payload) const { + + launch_populate_grid_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::count_doublets_kernel( + const count_doublets_kernel_payload& payload) const { + + launch_count_doublets_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::find_doublets_kernel( + const find_doublets_kernel_payload& payload) const { + + launch_find_doublets_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::count_triplets_kernel( + const count_triplets_kernel_payload& payload) const { + + launch_count_triplets_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::triplet_counts_reduction_kernel( + const triplet_counts_reduction_kernel_payload& payload) const { + + launch_triplet_counts_reduction_kernel( + payload, details::get_stream(stream()), warp_size()); +} + +void triplet_seeding_algorithm::find_triplets_kernel( + const find_triplets_kernel_payload& payload) const { + + launch_find_triplets_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::update_triplet_weights_kernel( + const update_triplet_weights_kernel_payload& payload) const { + + launch_update_triplet_weights_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +void triplet_seeding_algorithm::select_seeds_kernel( + const select_seeds_kernel_payload& payload) const { + + launch_select_seeds_kernel(payload, details::get_stream(stream()), + warp_size()); +} + +task triplet_seeding_algorithm::await( + vecmem::abstract_event& event) const { + co_await m_await_function(stream(), event); +} + +} // namespace traccc::cuda diff --git a/device/cuda/src/seeding/triplet_seeding_algorithm.cu b/device/cuda/src/seeding/triplet_seeding_kernel.cu similarity index 73% rename from device/cuda/src/seeding/triplet_seeding_algorithm.cu rename to device/cuda/src/seeding/triplet_seeding_kernel.cu index d58a6ad582..94703b6ba3 100644 --- a/device/cuda/src/seeding/triplet_seeding_algorithm.cu +++ b/device/cuda/src/seeding/triplet_seeding_kernel.cu @@ -9,7 +9,7 @@ #include "../utils/cuda_error_handling.hpp" #include "../utils/global_index.hpp" #include "../utils/utils.hpp" -#include "traccc/cuda/seeding/triplet_seeding_algorithm.hpp" +#include "triplet_seeding_kernel.hpp" // Project include(s). #include "traccc/seeding/detail/spacepoint_grid.hpp" @@ -161,143 +161,126 @@ __global__ void select_seeds( } // namespace kernels -triplet_seeding_algorithm::triplet_seeding_algorithm( - const seedfinder_config& finder_config, - const spacepoint_grid_config& grid_config, - const seedfilter_config& filter_config, const traccc::memory_resource& mr, - vecmem::copy& copy, cuda::stream& str, std::unique_ptr logger, - await_function_t await_func) - : device::triplet_seeding_algorithm(finder_config, grid_config, - filter_config, mr, copy, - std::move(logger)), - cuda::algorithm_base{str}, - m_await_function(await_func) {} - -void triplet_seeding_algorithm::count_grid_capacities_kernel( - const count_grid_capacities_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 8; +void launch_count_grid_capacities_kernel( + const traccc::device::triplet_seeding_count_grid_capacities_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size) { + + const unsigned int n_threads = warp_size * 8; const unsigned int n_blocks = (payload.n_spacepoints + n_threads - 1) / n_threads; - kernels::count_grid_capacities<<>>( + kernels::count_grid_capacities<<>>( payload.config, payload.phi_axis, payload.z_axis, payload.spacepoints, payload.grid_capacities); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::populate_grid_kernel( - const populate_grid_kernel_payload& payload) const { +void launch_populate_grid_kernel( + const traccc::device::triplet_seeding_populate_grid_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size) { - const unsigned int n_threads = warp_size() * 8; + const unsigned int n_threads = warp_size * 8; const unsigned int n_blocks = (payload.n_spacepoints + n_threads - 1) / n_threads; - kernels::populate_grid<<>>( + kernels::populate_grid<<>>( payload.config, payload.spacepoints, payload.grid, payload.grid_prefix_sum); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::count_doublets_kernel( - const count_doublets_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_count_doublets_kernel( + const traccc::device::triplet_seeding_count_doublets_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.n_spacepoints + n_threads - 1) / n_threads; - kernels::count_doublets<<>>( + kernels::count_doublets<<>>( payload.config, payload.spacepoints, payload.grid, payload.grid_prefix_sum, payload.doublet_counter, payload.nMidBot, payload.nMidTop); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::find_doublets_kernel( - const find_doublets_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_find_doublets_kernel( + const traccc::device::triplet_seeding_find_doublets_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.n_doublets + n_threads - 1) / n_threads; - kernels::find_doublets<<>>( + kernels::find_doublets<<>>( payload.config, payload.spacepoints, payload.grid, payload.doublet_counter, payload.mb_doublets, payload.mt_doublets); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::count_triplets_kernel( - const count_triplets_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_count_triplets_kernel( + const traccc::device::triplet_seeding_count_triplets_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.nMidBot + n_threads - 1) / n_threads; - kernels::count_triplets<<>>( + kernels::count_triplets<<>>( payload.config, payload.spacepoints, payload.grid, payload.doublet_counter, payload.mb_doublets, payload.mt_doublets, payload.spM_counter, payload.midBot_counter); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::triplet_counts_reduction_kernel( - const triplet_counts_reduction_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_triplet_counts_reduction_kernel( + const traccc::device:: + triplet_seeding_triplet_counts_reduction_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.n_doublets + n_threads - 1) / n_threads; - kernels::reduce_triplet_counts<<>>( + kernels::reduce_triplet_counts<<>>( payload.doublet_counter, payload.spM_counter, payload.nTriplets); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::find_triplets_kernel( - const find_triplets_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_find_triplets_kernel( + const traccc::device::triplet_seeding_find_triplets_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.nMidBot + n_threads - 1) / n_threads; - kernels::find_triplets<<>>( + kernels::find_triplets<<>>( payload.finding_config, payload.filter_config, payload.spacepoints, payload.grid, payload.doublet_counter, payload.mt_doublets, payload.spM_tc, payload.midBot_tc, payload.triplets); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::update_triplet_weights_kernel( - const update_triplet_weights_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_update_triplet_weights_kernel( + const traccc::device::triplet_seeding_update_triplet_weights_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.n_triplets + n_threads - 1) / n_threads; kernels::update_triplet_weights<<< n_blocks, n_threads, - sizeof(scalar) * payload.config.compatSeedLimit * n_threads, - details::get_stream(stream())>>>(payload.config, payload.spacepoints, - payload.spM_tc, payload.midBot_tc, - payload.triplets); + sizeof(scalar) * payload.config.compatSeedLimit * n_threads, stream>>>( + payload.config, payload.spacepoints, payload.spM_tc, payload.midBot_tc, + payload.triplets); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::select_seeds_kernel( - const select_seeds_kernel_payload& payload) const { - - const unsigned int n_threads = warp_size() * 2; +void launch_select_seeds_kernel( + const traccc::device::triplet_seeding_select_seeds_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size) { + const unsigned int n_threads = warp_size * 2; const unsigned int n_blocks = (payload.n_doublets + n_threads - 1) / n_threads; kernels:: select_seeds<<>>( - payload.finder_config, payload.filter_config, payload.spacepoints, - payload.grid, payload.spM_tc, payload.midBot_tc, payload.triplets, - payload.seeds); + stream>>>(payload.finder_config, payload.filter_config, + payload.spacepoints, payload.grid, + payload.spM_tc, payload.midBot_tc, + payload.triplets, payload.seeds); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } -void triplet_seeding_algorithm::await(vecmem::abstract_event& event) const { - m_await_function(stream(), event); -} - } // namespace traccc::cuda diff --git a/device/cuda/src/seeding/triplet_seeding_kernel.hpp b/device/cuda/src/seeding/triplet_seeding_kernel.hpp new file mode 100644 index 0000000000..60589ffc31 --- /dev/null +++ b/device/cuda/src/seeding/triplet_seeding_kernel.hpp @@ -0,0 +1,59 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2026 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/seeding/device/triplet_seeding_kernel_payloads.hpp" + +// CUDA include(s). +#include + +namespace traccc::cuda { + +void launch_count_grid_capacities_kernel( + const traccc::device::triplet_seeding_count_grid_capacities_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_populate_grid_kernel( + const traccc::device::triplet_seeding_populate_grid_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_count_doublets_kernel( + const traccc::device::triplet_seeding_count_doublets_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_find_doublets_kernel( + const traccc::device::triplet_seeding_find_doublets_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_count_triplets_kernel( + const traccc::device::triplet_seeding_count_triplets_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_triplet_counts_reduction_kernel( + const traccc::device:: + triplet_seeding_triplet_counts_reduction_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_find_triplets_kernel( + const traccc::device::triplet_seeding_find_triplets_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_update_triplet_weights_kernel( + const traccc::device::triplet_seeding_update_triplet_weights_kernel_payload& + payload, + cudaStream_t stream, unsigned int warp_size); + +void launch_select_seeds_kernel( + const traccc::device::triplet_seeding_select_seeds_kernel_payload& payload, + cudaStream_t stream, unsigned int warp_size); + +} // namespace traccc::cuda diff --git a/device/cuda/src/utils/algorithm_base.cpp b/device/cuda/src/utils/algorithm_base.cpp index ea9a5136ac..eb1bd1aa65 100644 --- a/device/cuda/src/utils/algorithm_base.cpp +++ b/device/cuda/src/utils/algorithm_base.cpp @@ -9,6 +9,7 @@ #include "traccc/cuda/utils/algorithm_base.hpp" #include "../utils/utils.hpp" +#include "traccc/execution/task.hpp" namespace traccc::cuda { @@ -25,17 +26,22 @@ unsigned int algorithm_base::warp_size() const { return m_warp_size; } -void default_await_function(const cuda::stream& stream, - vecmem::abstract_event& event) { - await_event_sync(stream, event); +task default_await_function(const cuda::stream&, + vecmem::abstract_event& event) { + event.wait(); + co_return; } -void await_stream_sync(const cuda::stream& stream, vecmem::abstract_event&) { +task await_stream_sync(const cuda::stream& stream, + vecmem::abstract_event&) { stream.synchronize(); + co_return; } -void await_event_sync(const cuda::stream&, vecmem::abstract_event& event) { +task await_event_sync(const cuda::stream&, + vecmem::abstract_event& event) { event.wait(); + co_return; } } // namespace traccc::cuda diff --git a/examples/run/CMakeLists.txt b/examples/run/CMakeLists.txt index ccee7a13a3..ca1a698801 100644 --- a/examples/run/CMakeLists.txt +++ b/examples/run/CMakeLists.txt @@ -19,9 +19,11 @@ add_library(traccc_examples_common STATIC "common/throughput_st.ipp" "common/await_strategy.hpp" "common/event_sync_strategy.hpp" + "common/task_arena_scheduler.hpp" + "common/task_arena_scheduler.cpp" ) target_link_libraries(traccc_examples_common - PUBLIC traccc::core traccc::options + PUBLIC traccc::core traccc::options STDEXEC::stdexec TBB::tbb PRIVATE traccc::io) # Add all the subdirectories that can be built. diff --git a/examples/run/common/await_strategy.hpp b/examples/run/common/await_strategy.hpp index a1aa61ab53..82134cedd1 100644 --- a/examples/run/common/await_strategy.hpp +++ b/examples/run/common/await_strategy.hpp @@ -4,8 +4,14 @@ namespace traccc { /// Enumeration of await strategies for synchronous or suspending operations enum class await_strategy { - sync_event, ///< Synchronous waiting on an event - sync_stream ///< Synchronous waiting on a stream + sync_event, ///< Synchronous waiting on an event + sync_stream, ///< Synchronous waiting on a stream + callback, ///< Suspending on a stream with a callback + poll, ///< Suspending with polling on an event + defer_sync_event, ///< Suspending and deferring event synchronization to a + ///< service threadpool + defer_sync_stream ///< Suspending and deferring stream synchronization to a + ///< service }; } // namespace traccc diff --git a/examples/run/common/task_arena_scheduler.cpp b/examples/run/common/task_arena_scheduler.cpp new file mode 100644 index 0000000000..241c0936f2 --- /dev/null +++ b/examples/run/common/task_arena_scheduler.cpp @@ -0,0 +1,30 @@ +// Local include(s). +#include "task_arena_scheduler.hpp" + +// TBB include(s). +#include + +namespace traccc { + +task_arena_scheduler::task_arena_scheduler(tbb::task_arena& task_arena) + : m_arena(&task_arena) {} + +task_arena_scheduler::sender task_arena_scheduler::schedule() const noexcept { + return sender{m_arena}; +} + +task_arena_scheduler::env::env(tbb::task_arena* arena) noexcept + : m_arena(arena) {} + +task_arena_scheduler::sender::sender(tbb::task_arena* arena) noexcept + : m_arena(arena) {} + +task_arena_scheduler::env task_arena_scheduler::sender::get_env() + const noexcept { + return env{m_arena}; +} + +static_assert(stdexec::scheduler, + "task_arena_scheduler should model scheduler"); + +} // namespace traccc diff --git a/examples/run/common/task_arena_scheduler.hpp b/examples/run/common/task_arena_scheduler.hpp new file mode 100644 index 0000000000..af8fc8daca --- /dev/null +++ b/examples/run/common/task_arena_scheduler.hpp @@ -0,0 +1,83 @@ +#pragma once + +// TBB include(s). +#include + +// Stdexec include(s). +#include + +namespace traccc { + +/// Wrapper around a TBB task arena to be used as a scheduler for stdexec. +class task_arena_scheduler { + public: + using scheduler_concept = stdexec::scheduler_t; + + /// Construct a task_arena_scheduler that uses the given TBB task arena. + /// @param task_arena The TBB task arena to use for scheduling. + /// + /// @note The task_arena_scheduler does not take ownership of the task + /// arena, the task arena should remain valid for the lifetime of the + /// scheduler. + /// + task_arena_scheduler(tbb::task_arena& task_arena); + + class env { + public: + env(tbb::task_arena* arena) noexcept; + + template + auto query( + const stdexec::get_completion_scheduler_t&) const noexcept { + return task_arena_scheduler{*m_arena}; + } + + private: + tbb::task_arena* m_arena; /// non-owning pointer to the task arena + }; + + template + class operation { + public: + using operation_state_concept = stdexec::operation_state_t; + + operation(Receiver&& receiver, tbb::task_arena* arena) noexcept + : m_receiver(std::forward(receiver)), m_arena(arena) {} + + void start() & noexcept { + m_arena->enqueue( + [this]() { stdexec::set_value(std::move(m_receiver)); }); + } + + private: + std::remove_cvref_t m_receiver; + tbb::task_arena* m_arena; + }; + + class sender { + public: + using sender_concept = stdexec::sender_t; + using completion_signatures = + stdexec::completion_signatures; + + sender(tbb::task_arena* arena) noexcept; + env get_env() const noexcept; + + template + auto connect(Receiver&& receiver) { + return operation(std::forward(receiver), + m_arena); + } + + private: + tbb::task_arena* m_arena; + }; + + sender schedule() const noexcept; + bool operator==(const task_arena_scheduler& other) const = default; + + private: + tbb::task_arena* m_arena = nullptr; +}; + +} // namespace traccc diff --git a/examples/run/common/throughput_mt.ipp b/examples/run/common/throughput_mt.ipp index 3f0f6d9852..b6eb7c6524 100644 --- a/examples/run/common/throughput_mt.ipp +++ b/examples/run/common/throughput_mt.ipp @@ -11,9 +11,11 @@ #include "await_strategy.hpp" #include "event_sync_strategy.hpp" #include "make_magnetic_field.hpp" +#include "task_arena_scheduler.hpp" #include "traccc/examples/utils/threadpool.hpp" // Project include(s) +#include "traccc/execution/task.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/geometry/host_detector.hpp" #include "traccc/seeding/detail/track_params_estimation_config.hpp" @@ -55,6 +57,9 @@ #include #include +// Stdexec include(s). +#include + // Indicators include(s). #include @@ -207,11 +212,13 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { case opts::threading::await_strategy::sync_stream: return await_strategy::sync_stream; case opts::threading::await_strategy::callback: + return await_strategy::callback; case opts::threading::await_strategy::poll: + return await_strategy::poll; case opts::threading::await_strategy::defer_sync_event: + return await_strategy::defer_sync_event; case opts::threading::await_strategy::defer_sync_stream: - throw std::invalid_argument( - "Suspending await strategies are not supported"); + return await_strategy::defer_sync_stream; default: throw std::invalid_argument("Unknown await strategy"); } @@ -235,21 +242,24 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { } // Set up a lambda that calls the correct function on the algorithms. - std::function + std::function(std::vector&, int, + const edm::silicon_cell_collection::host&)> process_event; if (throughput_opts.reco_stage == opts::throughput::stage::seeding) { - process_event = [&](int thread, - const edm::silicon_cell_collection::host& cells) - -> std::size_t { - return algs.at(static_cast(thread)) - .seeding(cells) - .size(); + process_event = [](std::vector& algs_, int slot_, + const edm::silicon_cell_collection::host& cells_) + -> task { + auto result = co_await algs_.at(static_cast(slot_)) + .seeding(cells_); + co_return result.size(); }; } else if (throughput_opts.reco_stage == opts::throughput::stage::full) { - process_event = [&](int thread, - const edm::silicon_cell_collection::host& cells) - -> std::size_t { - return algs.at(static_cast(thread))(cells).size(); + process_event = [](std::vector& algs_, int slot_, + const edm::silicon_cell_collection::host& cells_) + -> task { + auto result = + co_await algs_.at(static_cast(slot_))(cells_); + co_return result.size(); }; } else { throw std::invalid_argument("Unknown reconstruction stage"); @@ -261,7 +271,8 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { tbb::global_control::max_allowed_parallelism, threading_opts.threads + 1); tbb::task_arena arena{static_cast(threading_opts.threads), 0}; - tbb::task_group group; + auto scheduler = task_arena_scheduler{arena}; + auto scope = exec::async_scope{}; // Seed the random number generator. if (throughput_opts.random_seed == 0u) { @@ -300,19 +311,25 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { // Get a free concurrent slot. size_t slot = std::numeric_limits::max(); concurrent_slots.pop(slot); + auto payload = [](auto& algs_, auto& input_, auto& progress_bar_, + auto& rec_track_params_, auto& queue_, + size_t event_, size_t slot_, + auto& process_event_) -> task { + auto result = co_await process_event_( + algs_, static_cast(slot_), input_.at(event_)); + rec_track_params_.fetch_add(result); + progress_bar_.tick(); + queue_.push(slot_); + }; // Launch the processing of the event. - arena.execute([&, event, slot]() { - group.run([&, event, slot]() { - rec_track_params.fetch_add( - process_event(static_cast(slot), input[event])); - progress_bar.tick(); - concurrent_slots.push(slot); - }); - }); + scope.spawn(stdexec::starts_on( + scheduler, + payload(algs, input, progress_bar, rec_track_params, + concurrent_slots, event, slot, process_event))); } // Wait for all tasks to finish. - group.wait(); + stdexec::sync_wait(scope.on_empty()); } // Reset the dummy counter. @@ -342,19 +359,25 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { // Get a free slot. size_t slot = std::numeric_limits::max(); concurrent_slots.pop(slot); + auto payload = [](auto& algs_, auto& input_, auto& progress_bar_, + auto& rec_track_params_, auto& queue_, + size_t event_, size_t slot_, + auto& process_event_) -> task { + auto result = co_await process_event_( + algs_, static_cast(slot_), input_.at(event_)); + rec_track_params_.fetch_add(result); + progress_bar_.tick(); + queue_.push(slot_); + }; // Launch the processing of the event. - arena.execute([&, event, slot]() { - group.run([&, event, slot]() { - rec_track_params.fetch_add( - process_event(static_cast(slot), input[event])); - progress_bar.tick(); - concurrent_slots.push(slot); - }); - }); + scope.spawn(stdexec::starts_on( + scheduler, + payload(algs, input, progress_bar, rec_track_params, + concurrent_slots, event, slot, process_event))); } // Wait for all tasks to finish. - group.wait(); + stdexec::sync_wait(scope.on_empty()); } // Delete the algorithms explicitly before their parent object would go out diff --git a/examples/run/common/throughput_st.ipp b/examples/run/common/throughput_st.ipp index 8fe118b58f..92425b647c 100644 --- a/examples/run/common/throughput_st.ipp +++ b/examples/run/common/throughput_st.ipp @@ -12,6 +12,7 @@ #include "traccc/examples/utils/threadpool.hpp" // Project include(s) +#include "traccc/execution/task.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/geometry/host_detector.hpp" #include "traccc/seeding/detail/track_params_estimation_config.hpp" @@ -46,6 +47,10 @@ // Indicators include(s). #include +// Stdexec include(s). +#include +#include + // System include(s). #include #include @@ -159,18 +164,30 @@ int throughput_st(std::string_view description, int argc, char* argv[]) { } // Set up a lambda that calls the correct function on the algorithm. - std::function + std::function(FULL_CHAIN_ALG*, + const edm::silicon_cell_collection::host&)> process_event; if (throughput_opts.reco_stage == opts::throughput::stage::seeding) { - process_event = [&](const edm::silicon_cell_collection::host& cells) - -> std::size_t { return alg->seeding(cells).size(); }; + process_event = [](FULL_CHAIN_ALG* alg_, + const edm::silicon_cell_collection::host& cells_) + -> task { + auto result = co_await alg_->seeding(cells_); + co_return result.size(); + }; } else if (throughput_opts.reco_stage == opts::throughput::stage::full) { - process_event = [&](const edm::silicon_cell_collection::host& cells) - -> std::size_t { return (*alg)(cells).size(); }; + process_event = [](FULL_CHAIN_ALG* alg_, + const edm::silicon_cell_collection::host& cells_) + -> task { + auto result = co_await (*alg_)(cells_); + co_return result.size(); + }; } else { throw std::invalid_argument("Unknown reconstruction stage"); } + // Set up a scheduler executing the tasks on the current thread. + stdexec::inline_scheduler scheduler; + // Dummy count uses output of tp algorithm to ensure the compiler // optimisations don't skip any step std::size_t rec_track_params = 0; @@ -200,7 +217,12 @@ int throughput_st(std::string_view description, int argc, char* argv[]) { input_opts.events; // Process one event. - rec_track_params += process_event(input[event]); + auto result = stdexec::sync_wait(stdexec::starts_on( + scheduler, process_event(alg.get(), input[event]))); + if (!result.has_value()) { + throw std::runtime_error("Task execution failed"); + } + rec_track_params += std::get<0>(result.value()); progress_bar.tick(); } } @@ -231,7 +253,12 @@ int throughput_st(std::string_view description, int argc, char* argv[]) { input_opts.events; // Process one event. - rec_track_params += process_event(input[event]); + auto result = stdexec::sync_wait(stdexec::starts_on( + scheduler, process_event(alg.get(), input[event]))); + if (!result.has_value()) { + throw std::runtime_error("Task execution failed"); + } + rec_track_params += std::get<0>(result.value()); progress_bar.tick(); } } diff --git a/examples/run/cpu/CMakeLists.txt b/examples/run/cpu/CMakeLists.txt index d8b44e2a47..89a6b64c5c 100644 --- a/examples/run/cpu/CMakeLists.txt +++ b/examples/run/cpu/CMakeLists.txt @@ -4,27 +4,28 @@ # # Mozilla Public License Version 2.0 -traccc_add_executable( seeding_example "seeding_example.cpp" - LINK_LIBRARIES vecmem::core traccc::core traccc::io - traccc::performance traccc::options detray::detectors detray::io - traccc_examples_common ) - -traccc_add_executable( seq_example "seq_example.cpp" - LINK_LIBRARIES vecmem::core traccc::core traccc::io - traccc::performance traccc::options detray::detectors detray::io - traccc_examples_common ) - -traccc_add_executable( truth_finding_example "truth_finding_example.cpp" - LINK_LIBRARIES vecmem::core detray::detectors traccc::core traccc::io - traccc::performance traccc::options traccc_examples_common ) - -traccc_add_executable( truth_fitting_example "truth_fitting_example.cpp" - LINK_LIBRARIES vecmem::core detray::io detray::detectors traccc::core - traccc::io traccc::performance traccc::options traccc_examples_common ) - -traccc_add_executable( misaligned_truth_fitting_example "misaligned_truth_fitting_example.cpp" - LINK_LIBRARIES vecmem::core detray::io detray::detectors traccc::core - traccc::io traccc::performance traccc::options traccc_examples_common ) +# Disable examples for now +# traccc_add_executable( seeding_example "seeding_example.cpp" +# LINK_LIBRARIES vecmem::core traccc::core traccc::io +# traccc::performance traccc::options detray::detectors detray::io +# traccc_examples_common ) +# +# traccc_add_executable( seq_example "seq_example.cpp" +# LINK_LIBRARIES vecmem::core traccc::core traccc::io +# traccc::performance traccc::options detray::detectors detray::io +# traccc_examples_common ) +# +# traccc_add_executable( truth_finding_example "truth_finding_example.cpp" +# LINK_LIBRARIES vecmem::core detray::detectors traccc::core traccc::io +# traccc::performance traccc::options traccc_examples_common ) +# +# traccc_add_executable( truth_fitting_example "truth_fitting_example.cpp" +# LINK_LIBRARIES vecmem::core detray::io detray::detectors traccc::core +# traccc::io traccc::performance traccc::options traccc_examples_common ) +# +# traccc_add_executable( misaligned_truth_fitting_example "misaligned_truth_fitting_example.cpp" +# LINK_LIBRARIES vecmem::core detray::io detray::detectors traccc::core +# traccc::io traccc::performance traccc::options traccc_examples_common ) # # Set up the "throughput applications". diff --git a/examples/run/cpu/full_chain_algorithm.cpp b/examples/run/cpu/full_chain_algorithm.cpp index 8cabf468ee..96b9d77633 100644 --- a/examples/run/cpu/full_chain_algorithm.cpp +++ b/examples/run/cpu/full_chain_algorithm.cpp @@ -8,6 +8,9 @@ // Local include(s). #include "full_chain_algorithm.hpp" +// Project include(s). +#include "traccc/execution/task.hpp" + namespace traccc { full_chain_algorithm::full_chain_algorithm( @@ -82,20 +85,21 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( *m_detector, m_field, measurements_view, track_params_view); // Run the track fitting, and return its results. - return m_fitting(*m_detector, m_field, - edm::track_container::const_data( - track_candidates)) + co_return m_fitting( + *m_detector, m_field, + edm::track_container::const_data(track_candidates)) .tracks; } // If not, just return an empty object. else { // Return an empty object. - return output_type{m_mr.get()}; + co_return edm::track_collection::host{m_mr.get()}; } } -bound_track_parameters_collection_types::host full_chain_algorithm::seeding( +task +full_chain_algorithm::seeding( const edm::silicon_cell_collection::host& cells) const { // Create a data object for the detector description. @@ -121,14 +125,14 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( m_seeding(spacepoints_data); const edm::seed_collection::const_data seeds_data = vecmem::get_data(seeds); - return m_track_parameter_estimation(measurements_view, spacepoints_data, - seeds_data, m_field_vec); + co_return m_track_parameter_estimation( + measurements_view, spacepoints_data, seeds_data, m_field_vec); } // If not, just return an empty object. else { // Return an empty object. - return {}; + co_return {}; } } diff --git a/examples/run/cpu/full_chain_algorithm.hpp b/examples/run/cpu/full_chain_algorithm.hpp index c3b4987ef4..c8434ca0ca 100644 --- a/examples/run/cpu/full_chain_algorithm.hpp +++ b/examples/run/cpu/full_chain_algorithm.hpp @@ -18,6 +18,7 @@ #include "traccc/edm/silicon_cell_collection.hpp" #include "traccc/edm/track_collection.hpp" #include "traccc/edm/track_parameters.hpp" +#include "traccc/execution/task.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" @@ -46,7 +47,7 @@ namespace traccc { /// At least as much as is implemented in the project at any given moment. /// class full_chain_algorithm - : public algorithm::host( + : public algorithm::host>( const edm::silicon_cell_collection::host&)>, public messaging { @@ -93,7 +94,7 @@ class full_chain_algorithm /// Reconstruct track parameters in the entire detector /// /// @param cells The cells for every detector module in the event - /// @return The track parameters reconstructed + /// @return A task returning the track parameters reconstructed /// output_type operator()( const edm::silicon_cell_collection::host& cells) const override; @@ -101,9 +102,9 @@ class full_chain_algorithm /// Reconstruct track seeds in the entire detector /// /// @param cells The cells for every detector module in the event - /// @return The track seeds reconstructed + /// @return A task returning the track seeds reconstructed /// - bound_track_parameters_collection_types::host seeding( + task seeding( const edm::silicon_cell_collection::host& cells) const; private: diff --git a/examples/run/cuda/CMakeLists.txt b/examples/run/cuda/CMakeLists.txt index f8209e53c7..302b857431 100644 --- a/examples/run/cuda/CMakeLists.txt +++ b/examples/run/cuda/CMakeLists.txt @@ -10,31 +10,36 @@ include( traccc-compiler-options-cuda ) # External(s). find_package( CUDAToolkit REQUIRED ) -traccc_add_executable( seq_example_cuda "seq_example_cuda.cpp" - LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda_utils traccc::cuda - traccc::options detray::detectors detray::io - traccc_examples_common ) -traccc_add_executable( seeding_example_cuda "seeding_example_cuda.cpp" - LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda_utils traccc::cuda - traccc::options traccc_examples_common ) -traccc_add_executable( truth_finding_example_cuda "truth_finding_example_cuda.cpp" - LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda_utils traccc::cuda - traccc::options traccc_examples_common ) -traccc_add_executable( truth_fitting_example_cuda "truth_fitting_example_cuda.cpp" - LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda_utils traccc::cuda - traccc::options traccc_examples_common ) -# +# Disable examples for now +# traccc_add_executable( seq_example_cuda "seq_example_cuda.cpp" +# LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance +# traccc::core traccc::device_common traccc::cuda_utils traccc::cuda +# traccc::options detray::detectors detray::io +# traccc_examples_common ) +# traccc_add_executable( seeding_example_cuda "seeding_example_cuda.cpp" +# LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance +# traccc::core traccc::device_common traccc::cuda_utils traccc::cuda +# traccc::options traccc_examples_common ) +# traccc_add_executable( truth_finding_example_cuda "truth_finding_example_cuda.cpp" +# LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance +# traccc::core traccc::device_common traccc::cuda_utils traccc::cuda +# traccc::options traccc_examples_common ) +# traccc_add_executable( truth_fitting_example_cuda "truth_fitting_example_cuda.cpp" +# LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance +# traccc::core traccc::device_common traccc::cuda_utils traccc::cuda +# traccc::options traccc_examples_common ) + + # Set up the "throughput applications". # add_library( traccc_examples_cuda STATIC "full_chain_algorithm.hpp" "full_chain_algorithm.cpp" "device_config.hpp" - "device_config.cpp") + "device_config.cpp" + "await_exec.hpp" + "await_exec.cpp" + ) target_link_libraries( traccc_examples_cuda PUBLIC CUDA::cudart vecmem::core vecmem::cuda detray::core detray::detectors traccc::core traccc::device_common traccc::cuda_utils traccc::cuda diff --git a/examples/run/cuda/await_exec.cpp b/examples/run/cuda/await_exec.cpp new file mode 100644 index 0000000000..e64be07362 --- /dev/null +++ b/examples/run/cuda/await_exec.cpp @@ -0,0 +1,108 @@ +// Local include(s). +#include "await_exec.hpp" + +// Project include(s). +#include "traccc/cuda/utils/algorithm_base.hpp" +#include "traccc/cuda/utils/stream.hpp" +#include "traccc/execution/task.hpp" + +// CUDA includes(s). +#include + +// Stdexec include(s). +#include +#include + +#define CUDA_ERROR_CHECK(EXP) \ + do { \ + const cudaError_t errorCode = EXP; \ + if (errorCode != cudaSuccess) { \ + throw std::runtime_error(std::string("Failed to run " #EXP " (") + \ + cudaGetErrorString(errorCode) + ")"); \ + } \ + } while (false) + +namespace traccc::cuda { + +/// Wrapper sender suspending execution until all operations on a CUDA +/// stream are complete. +class stream_await_sender { + public: + // associated operation state + template + class stream_await_operation; + + using sender_concept = stdexec::sender_t; + using completion_signatures = + stdexec::completion_signatures; + + stream_await_sender(const cudaStream_t stream) : m_stream(stream) {} + stdexec::env<> get_env() const noexcept { return {}; } + + template + auto connect(Receiver&& receiver) const { + return stream_await_operation>( + std::forward(receiver), m_stream); + } + + private: + cudaStream_t m_stream; +}; + +/// Operation state associated with @c stream_await_sender +/// +template +class stream_await_sender::stream_await_operation { + public: + using operation_state_concept = stdexec::operation_state_t; + + stream_await_operation(Receiver&& recv, const cudaStream_t stream) + : m_receiver(std::forward(recv)), m_stream(stream) {} + + void start() & noexcept { + + auto error = cudaLaunchHostFunc(m_stream, callback, &m_receiver); + // resume immediately if the callback could not be registered + if (error != cudaSuccess) { + stdexec::set_value(std::move(m_receiver), error); + } + } + + private: + std::remove_cvref_t m_receiver; + cudaStream_t m_stream; + + static void CUDART_CB callback(void* userData) noexcept { + auto& recv = *static_cast(userData); + stdexec::set_value(recv, cudaSuccess); + } +}; + +static_assert(stdexec::sender); + +task await_poll::operator()(const cuda::stream&, + vecmem::abstract_event& event) const { + auto query_once = stdexec::just() | + stdexec::then([&event]() { return event.is_ready(); }); + co_await exec::repeat_effect_until(stdexec::starts_on( + threadpool_scheduler{threadpool}, std::move(query_once))); +} + +task await_callback(const cuda::stream& stream, vecmem::abstract_event&) { + auto cuda_stream = static_cast(stream.cudaStream()); + CUDA_ERROR_CHECK(co_await stream_await_sender{cuda_stream}); + co_return; +} + +task await_defer_event_sync::operator()( + const cuda::stream& stream, vecmem::abstract_event& event) const { + co_await stdexec::starts_on(threadpool_scheduler{threadpool}, + await_event_sync(stream, event)); +} + +task await_defer_stream_sync::operator()( + const cuda::stream& stream, vecmem::abstract_event& event) const { + co_await stdexec::starts_on(threadpool_scheduler{threadpool}, + await_stream_sync(stream, event)); +} +} // namespace traccc::cuda diff --git a/examples/run/cuda/await_exec.hpp b/examples/run/cuda/await_exec.hpp new file mode 100644 index 0000000000..f9cbc71fb3 --- /dev/null +++ b/examples/run/cuda/await_exec.hpp @@ -0,0 +1,40 @@ +#pragma once + +// Project include(s). +#include "traccc/cuda/utils/stream.hpp" +#include "traccc/examples/utils/threadpool.hpp" +#include "traccc/execution/task.hpp" + +// Vecmem include(s). +#include + +namespace traccc::cuda { + +/// Await CUDA steam completion by registering a callback on the stream +task await_callback(const cuda::stream& stream, + vecmem::abstract_event& event); + +/// Await CUDA event completion by polling in a service threadpool +struct await_poll { + traccc::threadpool& threadpool; + task operator()(const traccc::cuda::stream& stream, + vecmem::abstract_event& event) const; +}; + +/// Await CUDA event completion by deferring synchronization to a service +/// threadpool +struct await_defer_event_sync { + traccc::threadpool& threadpool; + task operator()(const traccc::cuda::stream& stream, + vecmem::abstract_event& event) const; +}; + +/// Await CUDA stream completion by deferring synchronization to a service +/// threadpool +struct await_defer_stream_sync { + traccc::threadpool& threadpool; + task operator()(const traccc::cuda::stream& stream, + vecmem::abstract_event& event) const; +}; + +} // namespace traccc::cuda diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 7425ebae54..7339aa7bd2 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -10,10 +10,12 @@ #include "../common/await_strategy.hpp" #include "../common/event_sync_strategy.hpp" +#include "await_exec.hpp" // Project include(s). #include "traccc/cuda/utils/algorithm_base.hpp" #include "traccc/cuda/utils/make_magnetic_field.hpp" +#include "traccc/execution/task.hpp" #include "traccc/seeding/detail/track_params_estimation_config.hpp" // Vecmem include(s). @@ -38,13 +40,21 @@ namespace traccc::cuda { -await_function_t get_await_function(await_strategy await_mode, - std::optional&) { +await_function_t get_await_function( + await_strategy await_mode, std::optional& threadpool) { switch (await_mode) { case await_strategy::sync_stream: return await_stream_sync; case await_strategy::sync_event: return await_event_sync; + case await_strategy::callback: + return await_callback; + case await_strategy::poll: + return await_poll{threadpool.value()}; + case traccc::await_strategy::defer_sync_event: + return await_defer_event_sync{threadpool.value()}; + case traccc::await_strategy::defer_sync_stream: + return await_defer_stream_sync{threadpool.value()}; default: throw std::invalid_argument("Unknown await strategy"); } @@ -203,18 +213,18 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // Run the clusterization (asynchronously). const auto unsorted_measurements = - m_clusterization(cells_buffer, m_device_det_descr); + co_await m_clusterization(cells_buffer, m_device_det_descr); const measurement_sorting_algorithm::output_type measurements = m_measurement_sorting(unsorted_measurements); // If we have a Detray detector, run the seeding, track finding and fitting. if (m_detector != nullptr) { // Run the seed-finding (asynchronously). - const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector, measurements); - const seed_parameter_estimation_algorithm::output_type track_params = - m_track_parameter_estimation(m_field, measurements, spacepoints, - m_seeding(spacepoints)); + const auto spacepoints = + co_await m_spacepoint_formation(m_device_detector, measurements); + const auto seeds = co_await m_seeding(spacepoints); + const auto track_params = co_await m_track_parameter_estimation( + m_field, measurements, spacepoints, seeds); // Run the track finding (asynchronously). const finding_algorithm::output_type track_candidates = @@ -224,10 +234,10 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( const auto host_tracks = m_copy.to(track_candidates.tracks, m_cached_pinned_host_mr, nullptr, vecmem::copy::type::device_to_host); - output_type result{m_host_mr}; + edm::track_collection::host result{m_host_mr}; vecmem::copy host_copy; host_copy(host_tracks, result)->wait(); - return result; + co_return result; } // If not, copy the measurements back to the host, and return a dummy @@ -240,11 +250,12 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_copy(measurements, measurements_host)->wait(); // Return an empty object. - return output_type{m_host_mr}; + co_return edm::track_collection::host{m_host_mr}; } } -bound_track_parameters_collection_types::host full_chain_algorithm::seeding( +task +full_chain_algorithm::seeding( const edm::silicon_cell_collection::host& cells) const { // Create device copy of input collections @@ -254,7 +265,7 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( // Run the clusterization (asynchronously). const auto unsorted_measurements = - m_clusterization(cells_buffer, m_device_det_descr); + co_await m_clusterization(cells_buffer, m_device_det_descr); const measurement_sorting_algorithm::output_type measurements = m_measurement_sorting(unsorted_measurements); @@ -262,11 +273,11 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( if (m_detector != nullptr) { // Run the seed-finding (asynchronously). - const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector, measurements); - const seed_parameter_estimation_algorithm::output_type track_params = - m_track_parameter_estimation(m_field, measurements, spacepoints, - m_seeding(spacepoints)); + const auto spacepoints = + co_await m_spacepoint_formation(m_device_detector, measurements); + const auto seeds = co_await m_seeding(spacepoints); + const auto track_params = co_await m_track_parameter_estimation( + m_field, measurements, spacepoints, seeds); // Copy a limited amount of result data back to the host. const auto host_seeds = m_copy.to(track_params, m_cached_pinned_host_mr, @@ -274,7 +285,7 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( bound_track_parameters_collection_types::host result{&m_host_mr}; vecmem::copy host_copy; host_copy(host_seeds, result)->wait(); - return result; + co_return result; } // If not, copy the measurements back to the host, and return a dummy @@ -287,7 +298,7 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( m_copy(measurements, measurements_host)->wait(); // Return an empty object. - return {}; + co_return {}; } } diff --git a/examples/run/cuda/full_chain_algorithm.hpp b/examples/run/cuda/full_chain_algorithm.hpp index 85df3fb5c7..a7d0b29796 100644 --- a/examples/run/cuda/full_chain_algorithm.hpp +++ b/examples/run/cuda/full_chain_algorithm.hpp @@ -26,6 +26,7 @@ #include "traccc/edm/silicon_cell_collection.hpp" #include "traccc/edm/track_collection.hpp" #include "traccc/edm/track_parameters.hpp" +#include "traccc/execution/task.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/geometry/detector_buffer.hpp" #include "traccc/geometry/host_detector.hpp" @@ -56,7 +57,7 @@ await_function_t get_await_function( /// At least as much as is implemented in the project at any given moment. /// class full_chain_algorithm - : public algorithm::host( + : public algorithm::host>( const edm::silicon_cell_collection::host&)>, public messaging { @@ -113,7 +114,7 @@ class full_chain_algorithm /// Reconstruct track parameters in the entire detector /// /// @param cells The cells for every detector module in the event - /// @return The track parameters reconstructed + /// @return A task returning the track parameters reconstructed /// output_type operator()( const edm::silicon_cell_collection::host& cells) const override; @@ -121,9 +122,9 @@ class full_chain_algorithm /// Reconstruct track seeds in the entire detector /// /// @param cells The cells for every detector module in the event - /// @return The track seeds reconstructed + /// @return A task returning the track seeds reconstructed /// - bound_track_parameters_collection_types::host seeding( + task seeding( const edm::silicon_cell_collection::host& cells) const; private: diff --git a/examples/utils/CMakeLists.txt b/examples/utils/CMakeLists.txt index 0f214c4104..f1d05036a1 100644 --- a/examples/utils/CMakeLists.txt +++ b/examples/utils/CMakeLists.txt @@ -9,4 +9,4 @@ traccc_add_library( traccc_utils utils TYPE SHARED "src/threadpool.cpp" ) -target_link_libraries(traccc_utils PUBLIC TBB::tbb) +target_link_libraries(traccc_utils PUBLIC TBB::tbb STDEXEC::stdexec) diff --git a/examples/utils/include/traccc/examples/utils/threadpool.hpp b/examples/utils/include/traccc/examples/utils/threadpool.hpp index 952a31b1db..377267b8c6 100644 --- a/examples/utils/include/traccc/examples/utils/threadpool.hpp +++ b/examples/utils/include/traccc/examples/utils/threadpool.hpp @@ -10,6 +10,9 @@ // Tbb include(s). #include +// Stdexec include(s). +#include + // System include(s). #include #include @@ -67,4 +70,78 @@ class threadpool { }; std::ostream& operator<<(std::ostream& os, threadpool::wait_policy policy); + +/// Wrapper around a threadpool to be used as a scheduler for stdexec. +class threadpool_scheduler { + public: + using scheduler_concept = stdexec::scheduler_t; + + /// Construct a threadpool_scheduler that uses the given threadpool. + /// @param threadpool The threadpool to use for scheduling. + /// + /// @note The threadpool_scheduler does not take ownership of the + /// threadpool, the threadpool should remain valid for the lifetime of the + /// scheduler. + /// + threadpool_scheduler(threadpool& pool); + + class env { + public: + env(threadpool* pool) noexcept; + + template + auto query( + const stdexec::get_completion_scheduler_t&) const noexcept { + return threadpool_scheduler{*m_threadpool}; + } + + private: + threadpool* m_threadpool; /// non-owning pointer to the threadpool + }; + + template + class operation { + public: + using operation_state_concept = stdexec::operation_state_t; + + operation(Receiver&& receiver, threadpool* pool) noexcept + : m_receiver(std::forward(receiver)), + m_threadpool(pool) {} + + void start() & noexcept { + m_threadpool->enqueue( + [this]() { stdexec::set_value(std::move(m_receiver)); }); + } + + private: + std::remove_cvref_t m_receiver; + threadpool* m_threadpool; + }; + + class sender { + public: + using sender_concept = stdexec::sender_t; + using completion_signatures = + stdexec::completion_signatures; + + sender(threadpool* pool) noexcept; + env get_env() const noexcept; + + template + auto connect(Receiver&& receiver) { + return operation(std::forward(receiver), + m_threadpool); + } + + private: + threadpool* m_threadpool; /// non-owning pointer to the threadpool + }; + + sender schedule() const noexcept; + bool operator==(const threadpool_scheduler& other) const = default; + + private: + threadpool* m_threadpool = nullptr; +}; + } // namespace traccc diff --git a/examples/utils/src/threadpool.cpp b/examples/utils/src/threadpool.cpp index b0dfc2a01c..d72af379ed 100644 --- a/examples/utils/src/threadpool.cpp +++ b/examples/utils/src/threadpool.cpp @@ -92,4 +92,25 @@ std::ostream& operator<<(std::ostream& os, threadpool::wait_policy policy) { return os; } +threadpool_scheduler::threadpool_scheduler(threadpool& pool) + : m_threadpool(&pool) {} + +threadpool_scheduler::sender threadpool_scheduler::schedule() const noexcept { + return sender{m_threadpool}; +} + +threadpool_scheduler::env::env(threadpool* pool) noexcept + : m_threadpool(pool) {} + +threadpool_scheduler::sender::sender(threadpool* pool) noexcept + : m_threadpool(pool) {} + +threadpool_scheduler::env threadpool_scheduler::sender::get_env() + const noexcept { + return env{m_threadpool}; +} + +static_assert(stdexec::scheduler, + "threadpool_scheduler should model scheduler"); + } // namespace traccc diff --git a/extern/stdexec/CMakeLists.txt b/extern/stdexec/CMakeLists.txt new file mode 100644 index 0000000000..0595daa1c6 --- /dev/null +++ b/extern/stdexec/CMakeLists.txt @@ -0,0 +1,31 @@ +# TRACCC library, part of the ACTS project (R&D line) +# +# (c) 2022-2026 CERN for the benefit of the ACTS project +# +# Mozilla Public License Version 2.0 + +cmake_minimum_required( VERSION 3.25 ) +include( FetchContent ) + +# Tell the user what's happening. +message( STATUS "Building stdexec as part of the TRACCC project" ) + +# Declare where to get stdexec from. +set( TRACCC_STDEXEC_SOURCE + "URL;https://github.com/NVIDIA/stdexec/archive/refs/tags/nvhpc-25.09.tar.gz;URL_MD5;e385824bb560a3ac36ce859a2e3ded1f" + CACHE STRING "Source for stdexec, when built as part of this project" ) + +mark_as_advanced( TRACCC_STDEXEC_SOURCE ) +FetchContent_Declare( stdexec SYSTEM ${TRACCC_STDEXEC_SOURCE} ) + +# Turn off build tests, examples, and docs. +set( BUILD_TESTING FALSE ) +set( STDEXEC_BUILD_EXAMPLES FALSE ) +set( STDEXEC_BUILD_DOCS FALSE ) +set( STDEXEC_ENABLE_TBB FALSE ) +set( STDEXEC_ENABLE_TASKFLOW FALSE ) +set( STDEXEC_ENABLE_CUDA FALSE ) # This option adds nvexec target with nvidia specific extensions and wrappers. + # Some parts of the stdexec still can't be compiled with nvcc. + +# Get it into the current directory. +FetchContent_MakeAvailable( stdexec ) diff --git a/extern/stdexec/README.md b/extern/stdexec/README.md new file mode 100644 index 0000000000..63b1053eba --- /dev/null +++ b/extern/stdexec/README.md @@ -0,0 +1,4 @@ +# Build Recipe for stdexec + +This directory holds a build recipe for building +[stdexec](https://github.com/NVIDIA/stdexec) for this project. diff --git a/spack.yaml b/spack.yaml index 447e0adc71..0c4fb0d896 100644 --- a/spack.yaml +++ b/spack.yaml @@ -18,6 +18,7 @@ spack: - "alpaka" - "boost@1.85.0: +log+program_options" - "indicators" + - "stdexec" # SYCL dependencies - "intel-oneapi-dpl" # Examples and test dependencies diff --git a/tests/cuda/test_cca.cpp b/tests/cuda/test_cca.cpp index ce9b2f85b6..4cb94e7b3b 100644 --- a/tests/cuda/test_cca.cpp +++ b/tests/cuda/test_cca.cpp @@ -19,6 +19,10 @@ #include "traccc/cuda/utils/stream.hpp" #include "traccc/geometry/silicon_detector_description.hpp" +// Stdexec include(s). +#include +#include + namespace { vecmem::host_memory_resource host_mr; @@ -58,9 +62,15 @@ cca_function_t get_f_with(traccc::clustering_config cfg) { copy.setup(cells_buffer)->wait(); copy(vecmem::get_data(cells), cells_buffer)->wait(); - auto [measurements_buffer, cluster_buffer] = + auto clustering_result = stdexec::sync_wait(stdexec::starts_on( + stdexec::inline_scheduler{}, cc(cells_buffer, dd_buffer, - traccc::device::clustering_keep_disjoint_set{}); + traccc::device::clustering_keep_disjoint_set{}))); + if (!clustering_result.has_value()) { + throw std::runtime_error("Clusterization algorithm failed"); + } + auto [measurements_buffer, cluster_buffer] = + std::move(std::get<0>(clustering_result.value())); traccc::edm::measurement_collection::host measurements{host_mr}; copy(measurements_buffer, measurements)->wait(); diff --git a/tests/cuda/test_clusterization.cpp b/tests/cuda/test_clusterization.cpp index 1c881f61e2..68ce3bc9ed 100644 --- a/tests/cuda/test_clusterization.cpp +++ b/tests/cuda/test_clusterization.cpp @@ -17,6 +17,10 @@ #include #include +// Stdexec include(s). +#include +#include + // GTest include(s). #include @@ -60,8 +64,12 @@ TEST(CUDAClustering, SingleModule) { traccc::cuda::clusterization_algorithm ca_cuda(mr, copy, stream, default_ccl_test_config()); + auto clusterization_results = stdexec::sync_wait(stdexec::starts_on( + stdexec::inline_scheduler{}, + ca_cuda(vecmem::get_data(cells), vecmem::get_data(dd)))); + ASSERT_TRUE(clusterization_results.has_value()); auto measurements_buffer = - ca_cuda(vecmem::get_data(cells), vecmem::get_data(dd)); + std::move(std::get<0>(clusterization_results.value())); edm::measurement_collection::const_device measurements( measurements_buffer); diff --git a/tests/cuda/test_spacepoint_formation.cpp b/tests/cuda/test_spacepoint_formation.cpp index 7dfa99bc26..3a96876e5c 100644 --- a/tests/cuda/test_spacepoint_formation.cpp +++ b/tests/cuda/test_spacepoint_formation.cpp @@ -19,6 +19,10 @@ #include #include +// Stdexec include(s). +#include +#include + // GTest include(s). #include @@ -94,8 +98,11 @@ TEST(CUDASpacepointFormation, cuda) { // Run spacepoint formation traccc::cuda::silicon_pixel_spacepoint_formation_algorithm sp_formation( mr, copy, stream); - auto spacepoints_buffer = - sp_formation(device_det, vecmem::get_data(measurements)); + auto result = stdexec::sync_wait(stdexec::starts_on( + stdexec::inline_scheduler{}, + sp_formation(device_det, vecmem::get_data(measurements)))); + ASSERT_TRUE(result.has_value()); + auto spacepoints_buffer = std::move(std::get<0>(result.value())); edm::spacepoint_collection::device spacepoints(spacepoints_buffer);