From e36f9345bd82d61691dba26e38f69fe767c4a3ad Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Wed, 4 Feb 2026 11:29:25 +0100 Subject: [PATCH 01/15] build with stdexec --- CMakeLists.txt | 19 +++++++++++++++++++ device/common/CMakeLists.txt | 2 +- extern/stdexec/CMakeLists.txt | 31 +++++++++++++++++++++++++++++++ extern/stdexec/README.md | 4 ++++ spack.yaml | 1 + 5 files changed, 56 insertions(+), 1 deletion(-) create mode 100644 extern/stdexec/CMakeLists.txt create mode 100644 extern/stdexec/README.md 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/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index 8866f4373c..442773c9e3 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -86,4 +86,4 @@ traccc_add_library( traccc_device_common device_common "include/traccc/fitting/device/impl/fill_fitting_sort_keys.ipp" ) target_link_libraries( traccc_device_common - PUBLIC traccc::core vecmem::core ) + PUBLIC traccc::core vecmem::core STDEXEC::stdexec ) 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 From bbdb41c37fc802360810cbb27a745b718489bbac Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Thu, 5 Feb 2026 12:54:58 +0100 Subject: [PATCH 02/15] split seed_parameter_estimation_algorithm cuda source so it compiles with task header --- device/common/CMakeLists.txt | 1 + .../seed_parameter_estimation_algorithm.hpp | 25 ++++-------- ...ed_parameter_estimation_kernel_payload.hpp | 40 +++++++++++++++++++ device/cuda/CMakeLists.txt | 4 +- .../seed_parameter_estimation_algorithm.hpp | 4 +- .../seed_parameter_estimation_algorithm.cpp | 39 ++++++++++++++++++ ...cu => seed_parameter_estimation_kernel.cu} | 26 +++--------- .../seed_parameter_estimation_kernel.hpp | 29 ++++++++++++++ 8 files changed, 127 insertions(+), 41 deletions(-) create mode 100644 device/common/include/traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp create mode 100644 device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp rename device/cuda/src/seeding/{seed_parameter_estimation_algorithm.cu => seed_parameter_estimation_kernel.cu} (64%) create mode 100644 device/cuda/src/seeding/seed_parameter_estimation_kernel.hpp diff --git a/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index 442773c9e3..858fba2e2f 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -64,6 +64,7 @@ traccc_add_library( traccc_device_common device_common "include/traccc/seeding/device/silicon_pixel_spacepoint_formation_algorithm.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/seeding/device/seed_parameter_estimation_algorithm.hpp b/device/common/include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp index 52a25329c6..6c32c83377 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,7 @@ // Local include(s). #include "traccc/device/algorithm_base.hpp" +#include "traccc/seeding/device/seed_parameter_estimation_kernel_payload.hpp" // Project include(s) #include "traccc/bfield/magnetic_field.hpp" @@ -21,6 +22,9 @@ #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" +// Stdexec include(s). +#include + namespace traccc::device { /// Seed track parameter estimation algorithm @@ -73,30 +77,15 @@ 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; /// @} 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/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index 55c5697a31..351d95f197 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -35,7 +35,9 @@ 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" "include/traccc/cuda/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" 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..da64a204e2 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,8 +49,7 @@ 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; /// @} 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..7e308e3388 --- /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()); +} + +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_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 From 79925dbc7e99fc12597b4f2e57eb94888ae8a2b4 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Thu, 5 Feb 2026 17:13:06 +0100 Subject: [PATCH 03/15] make seed_parameter_estimation_algorithm execution a task, update execution in throughput examples, disable the other examples --- .../seed_parameter_estimation_algorithm.hpp | 14 ++-- .../seed_parameter_estimation_algorithm.cpp | 4 +- examples/run/CMakeLists.txt | 4 +- examples/run/common/task_arena_scheduler.cpp | 30 +++++++ examples/run/common/task_arena_scheduler.hpp | 83 +++++++++++++++++++ examples/run/common/throughput_mt.ipp | 82 +++++++++++------- examples/run/common/throughput_st.ipp | 41 +++++++-- examples/run/cpu/CMakeLists.txt | 43 +++++----- examples/run/cpu/full_chain_algorithm.cpp | 20 +++-- examples/run/cpu/full_chain_algorithm.hpp | 11 ++- examples/run/cuda/CMakeLists.txt | 38 +++++---- examples/run/cuda/full_chain_algorithm.cpp | 26 +++--- examples/run/cuda/full_chain_algorithm.hpp | 11 ++- 13 files changed, 294 insertions(+), 113 deletions(-) create mode 100644 examples/run/common/task_arena_scheduler.cpp create mode 100644 examples/run/common/task_arena_scheduler.hpp 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 6c32c83377..61db9e9b67 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 @@ -33,11 +33,12 @@ namespace traccc::device { /// synchronisation statement is required before destroying this buffer. /// struct seed_parameter_estimation_algorithm - : public algorithm::const_view&, - const edm::spacepoint_collection::const_view&, - const edm::seed_collection::const_view&)>, + : public algorithm< + exec::task( + const magnetic_field&, + const edm::measurement_collection::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&)>, public messaging, public algorithm_base { @@ -63,7 +64,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, diff --git a/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp b/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp index 7f8acaf719..e90487f3ef 100644 --- a/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp +++ b/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp @@ -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/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/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..228b301a66 100644 --- a/examples/run/common/throughput_mt.ipp +++ b/examples/run/common/throughput_mt.ipp @@ -11,6 +11,7 @@ #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) @@ -55,6 +56,10 @@ #include #include +// Stdexec include(s). +#include +#include + // Indicators include(s). #include @@ -235,21 +240,25 @@ 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_) + -> exec::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_) + -> exec::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 +270,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 +310,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_) -> exec::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 +358,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_) -> exec::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..a8d881cee3 100644 --- a/examples/run/common/throughput_st.ipp +++ b/examples/run/common/throughput_st.ipp @@ -46,6 +46,11 @@ // Indicators include(s). #include +// Stdexec include(s). +#include +#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_) + -> exec::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_) + -> exec::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..4c82acbab8 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" +// Stdexec include(s). +#include + 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( +exec::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..a9dc57c57f 100644 --- a/examples/run/cpu/full_chain_algorithm.hpp +++ b/examples/run/cpu/full_chain_algorithm.hpp @@ -34,6 +34,9 @@ #include #include +// Stdexec include(s). +#include + // System include(s). #include #include @@ -46,7 +49,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 +96,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 +104,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( + exec::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..a1d31d7c54 100644 --- a/examples/run/cuda/CMakeLists.txt +++ b/examples/run/cuda/CMakeLists.txt @@ -10,24 +10,26 @@ 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 diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 7425ebae54..cbf8ae9ba8 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -22,6 +22,9 @@ // CUDA include(s). #include +// Stdexec include(s). +#include + // System include(s). #include #include @@ -212,9 +215,8 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // 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 track_params = co_await m_track_parameter_estimation( + m_field, measurements, spacepoints, m_seeding(spacepoints)); // Run the track finding (asynchronously). const finding_algorithm::output_type track_candidates = @@ -224,10 +226,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 +242,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( +exec::task +full_chain_algorithm::seeding( const edm::silicon_cell_collection::host& cells) const { // Create device copy of input collections @@ -264,9 +267,8 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( // 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 track_params = co_await m_track_parameter_estimation( + m_field, measurements, spacepoints, m_seeding(spacepoints)); // 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 +276,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 +289,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..2f916cebad 100644 --- a/examples/run/cuda/full_chain_algorithm.hpp +++ b/examples/run/cuda/full_chain_algorithm.hpp @@ -42,6 +42,9 @@ #include #include +// Stdexec include(s). +#include + // System include(s). #include #include @@ -56,7 +59,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 +116,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 +124,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( + exec::task seeding( const edm::silicon_cell_collection::host& cells) const; private: From 30f691dd3fe5fa8704f9f2079bf01cba9382ce3b Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Thu, 5 Feb 2026 17:52:23 +0100 Subject: [PATCH 04/15] split clusterization_algorithm --- device/common/CMakeLists.txt | 1 + .../device/clusterization_algorithm.hpp | 33 ++---------- .../device/clusterization_kernel_payload.hpp | 53 +++++++++++++++++++ device/cuda/CMakeLists.txt | 4 +- .../clusterization_algorithm.cpp | 47 ++++++++++++++++ ..._algorithm.cu => clusterization_kernel.cu} | 50 +++++++---------- .../clusterization/clusterization_kernel.hpp | 44 +++++++++++++++ 7 files changed, 171 insertions(+), 61 deletions(-) create mode 100644 device/common/include/traccc/clusterization/device/clusterization_kernel_payload.hpp create mode 100644 device/cuda/src/clusterization/clusterization_algorithm.cpp rename device/cuda/src/clusterization/{clusterization_algorithm.cu => clusterization_kernel.cu} (51%) create mode 100644 device/cuda/src/clusterization/clusterization_kernel.hpp diff --git a/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index 858fba2e2f..b96ef214e9 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" diff --git a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp index 89636c7842..69ec0dbf82 100644 --- a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp +++ b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp @@ -9,6 +9,7 @@ // 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" @@ -26,6 +27,9 @@ #include #include +// StdExec include(s). +#include + // System include(s). #include #include @@ -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 /// 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/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index 351d95f197..f27f3b7fc7 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -50,7 +50,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/src/clusterization/clusterization_algorithm.cpp b/device/cuda/src/clusterization/clusterization_algorithm.cpp new file mode 100644 index 0000000000..3a59b8b089 --- /dev/null +++ b/device/cuda/src/clusterization/clusterization_algorithm.cpp @@ -0,0 +1,47 @@ +/** 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()); +} + +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_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 From 663ebfdc98fb9d2583d304fe0e2fadf7123bd5d0 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Thu, 5 Feb 2026 18:15:28 +0100 Subject: [PATCH 05/15] split silicon_pixel_spacepoint_formation_algorithm --- device/common/CMakeLists.txt | 1 + ...n_pixel_spacepoint_formation_algorithm.hpp | 18 +++------ ...el_spacepoint_formation_kernel_payload.hpp | 34 +++++++++++++++++ device/cuda/CMakeLists.txt | 4 +- ...n_pixel_spacepoint_formation_algorithm.cpp | 38 +++++++++++++++++++ ...icon_pixel_spacepoint_formation_kernel.cu} | 30 ++++----------- ...icon_pixel_spacepoint_formation_kernel.hpp | 29 ++++++++++++++ 7 files changed, 119 insertions(+), 35 deletions(-) create mode 100644 device/common/include/traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp create mode 100644 device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp rename device/cuda/src/seeding/{silicon_pixel_spacepoint_formation_algorithm.cu => silicon_pixel_spacepoint_formation_kernel.cu} (57%) create mode 100644 device/cuda/src/seeding/silicon_pixel_spacepoint_formation_kernel.hpp diff --git a/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index b96ef214e9..6f4f7b242f 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -63,6 +63,7 @@ traccc_add_library( traccc_device_common device_common "include/traccc/seeding/device/triplet_seeding_algorithm.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" 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..9a1808f29e 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,7 @@ // Local include(s). #include "traccc/device/algorithm_base.hpp" +#include "traccc/seeding/device/silicon_pixel_spacepoint_formation_kernel_payload.hpp" // Project include(s). #include "traccc/edm/measurement_collection.hpp" @@ -18,6 +19,9 @@ #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" +// Stdexec include(s). +#include + namespace traccc::device { /// Algorithm forming space points out of measurements @@ -61,18 +65,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 /// 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/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index f27f3b7fc7..764e2f0361 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -41,7 +41,9 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED "include/traccc/cuda/seeding/triplet_seeding_algorithm.hpp" "src/seeding/triplet_seeding_algorithm.cu" "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" 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..4234de4816 --- /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()); +} + +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_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 From 3047e8938aad7a78df95da7b49eb3bbc612deb82 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Thu, 5 Feb 2026 22:42:25 +0100 Subject: [PATCH 06/15] split triplet_seeding_algorithm --- device/common/CMakeLists.txt | 1 + .../device/triplet_seeding_algorithm.hpp | 174 +++--------------- .../triplet_seeding_kernel_payloads.hpp | 120 ++++++++++++ device/cuda/CMakeLists.txt | 4 +- .../src/seeding/triplet_seeding_algorithm.cpp | 95 ++++++++++ ...algorithm.cu => triplet_seeding_kernel.cu} | 131 ++++++------- .../src/seeding/triplet_seeding_kernel.hpp | 59 ++++++ 7 files changed, 357 insertions(+), 227 deletions(-) create mode 100644 device/common/include/traccc/seeding/device/triplet_seeding_kernel_payloads.hpp create mode 100644 device/cuda/src/seeding/triplet_seeding_algorithm.cpp rename device/cuda/src/seeding/{triplet_seeding_algorithm.cu => triplet_seeding_kernel.cu} (73%) create mode 100644 device/cuda/src/seeding/triplet_seeding_kernel.hpp diff --git a/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index 6f4f7b242f..b4662c22b0 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -61,6 +61,7 @@ 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" 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..eee4e06856 100644 --- a/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp +++ b/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp @@ -20,10 +20,14 @@ #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" +// Stdexec include(s). +#include + // System include(s). #include @@ -73,22 +77,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 +88,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 +99,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 +110,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 +121,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 +132,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 +143,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 +154,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 +165,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 /// 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/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index 764e2f0361..3b7187306a 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -39,7 +39,9 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED "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.cpp" "src/seeding/silicon_pixel_spacepoint_formation_kernel.cu" 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..0b9fdc34b3 --- /dev/null +++ b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp @@ -0,0 +1,95 @@ +/** 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()); +} + +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_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 From 4552a56c76461ac0cadeb904b3b3225847f4d3df Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Thu, 5 Feb 2026 23:39:53 +0100 Subject: [PATCH 07/15] make spacepoint formation algorithm execution a task --- .../silicon_pixel_spacepoint_formation_algorithm.hpp | 5 +++-- .../silicon_pixel_spacepoint_formation_algorithm.cpp | 4 ++-- examples/run/cuda/full_chain_algorithm.cpp | 8 ++++---- tests/cuda/test_spacepoint_formation.cpp | 11 +++++++++-- 4 files changed, 18 insertions(+), 10 deletions(-) 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 9a1808f29e..12f99ae0f1 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 @@ -30,7 +30,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, @@ -52,7 +52,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()( 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..5dece21b5b 100644 --- a/device/common/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp +++ b/device/common/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -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/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index cbf8ae9ba8..13bd657d06 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -213,8 +213,8 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // 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 auto spacepoints = + co_await m_spacepoint_formation(m_device_detector, measurements); const auto track_params = co_await m_track_parameter_estimation( m_field, measurements, spacepoints, m_seeding(spacepoints)); @@ -265,8 +265,8 @@ 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 auto spacepoints = + co_await m_spacepoint_formation(m_device_detector, measurements); const auto track_params = co_await m_track_parameter_estimation( m_field, measurements, spacepoints, m_seeding(spacepoints)); 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); From 0d926f0db4c7c3a820e87c29897462419d6d50d6 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 6 Feb 2026 00:10:34 +0100 Subject: [PATCH 08/15] make triplet_seeding_algorithm execution a task --- .../traccc/seeding/device/triplet_seeding_algorithm.hpp | 5 +++-- device/common/src/seeding/triplet_seeding_algorithm.cpp | 8 ++++---- examples/run/cuda/full_chain_algorithm.cpp | 6 ++++-- 3 files changed, 11 insertions(+), 8 deletions(-) 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 eee4e06856..08bbd4d9f3 100644 --- a/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp +++ b/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp @@ -39,7 +39,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 { @@ -67,7 +67,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; diff --git a/device/common/src/seeding/triplet_seeding_algorithm.cpp b/device/common/src/seeding/triplet_seeding_algorithm.cpp index 5550486038..4f9d91ee30 100644 --- a/device/common/src/seeding/triplet_seeding_algorithm.cpp +++ b/device/common/src/seeding/triplet_seeding_algorithm.cpp @@ -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 @@ -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. @@ -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/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 13bd657d06..b22f2b17a7 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -215,8 +215,9 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // Run the seed-finding (asynchronously). 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, m_seeding(spacepoints)); + m_field, measurements, spacepoints, seeds); // Run the track finding (asynchronously). const finding_algorithm::output_type track_candidates = @@ -267,8 +268,9 @@ full_chain_algorithm::seeding( // Run the seed-finding (asynchronously). 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, m_seeding(spacepoints)); + 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, From fc15d72b3e33fd422b512a1b4f410e599d116ee9 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 6 Feb 2026 00:56:49 +0100 Subject: [PATCH 09/15] make clusterization_algorithm execution a task --- .../device/clusterization_algorithm.hpp | 35 ++++++++++--------- .../clusterization_algorithm.cpp | 34 +++++++++--------- examples/run/cuda/full_chain_algorithm.cpp | 4 +-- tests/cuda/test_cca.cpp | 14 ++++++-- tests/cuda/test_clusterization.cpp | 10 +++++- 5 files changed, 60 insertions(+), 37 deletions(-) diff --git a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp index 69ec0dbf82..8fccda8276 100644 --- a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp +++ b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp @@ -45,19 +45,21 @@ 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< + exec::task::buffer>( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>, public algorithm< - std::pair::buffer, - edm::silicon_cluster_collection::buffer>( + exec::task::buffer>( const edm::silicon_cell_collection::const_view&, const silicon_detector_description::const_view&, - clustering_keep_disjoint_set&&)>, + clustering_discard_disjoint_set&&)>, + public algorithm::buffer, + edm::silicon_cluster_collection::buffer>>( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&, + clustering_keep_disjoint_set&&)>, public messaging, public algorithm_base { @@ -85,18 +87,18 @@ class clusterization_algorithm /// @return a measurement collection (buffer) /// /// @{ - edm::measurement_collection::buffer operator()( + exec::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()( + exec::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> + exec::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; @@ -141,8 +143,9 @@ class clusterization_algorithm private: /// Main algorithmic implementation of the clusterization algorithm - std::pair::buffer, - std::optional> + exec::task< + std::pair::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/src/clusterization/clusterization_algorithm.cpp b/device/common/src/clusterization/clusterization_algorithm.cpp index 833deaaf6a..a53038d426 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 +exec::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 +exec::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> +exec::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> +exec::task::buffer, + std::optional>> clusterization_algorithm::execute_impl( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, @@ -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 {}; } } @@ -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/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index b22f2b17a7..992a0d00c2 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -206,7 +206,7 @@ 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); @@ -258,7 +258,7 @@ 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); 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); From 6d3f6c1ff04cbd89ad54925b613e4a7f8bb3d749 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 6 Feb 2026 09:07:04 +0100 Subject: [PATCH 10/15] make await function a coroutine --- .../device/clusterization_algorithm.hpp | 4 ++-- .../seed_parameter_estimation_algorithm.hpp | 4 ++-- ...n_pixel_spacepoint_formation_algorithm.hpp | 4 ++-- .../device/triplet_seeding_algorithm.hpp | 4 ++-- .../clusterization_algorithm.cpp | 4 ++-- .../seed_parameter_estimation_algorithm.cpp | 2 +- ...n_pixel_spacepoint_formation_algorithm.cpp | 2 +- .../src/seeding/triplet_seeding_algorithm.cpp | 8 +++---- .../clusterization_algorithm.hpp | 2 +- .../seed_parameter_estimation_algorithm.hpp | 2 +- ...n_pixel_spacepoint_formation_algorithm.hpp | 2 +- .../seeding/triplet_seeding_algorithm.hpp | 2 +- .../traccc/cuda/utils/algorithm_base.hpp | 23 +++++++++++-------- .../clusterization_algorithm.cpp | 5 ++-- .../seed_parameter_estimation_algorithm.cpp | 4 ++-- ...n_pixel_spacepoint_formation_algorithm.cpp | 4 ++-- .../src/seeding/triplet_seeding_algorithm.cpp | 5 ++-- device/cuda/src/utils/algorithm_base.cpp | 15 ++++++++---- 18 files changed, 53 insertions(+), 43 deletions(-) diff --git a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp index 8fccda8276..2322d8bee0 100644 --- a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp +++ b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp @@ -138,8 +138,8 @@ 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 exec::task await(vecmem::abstract_event& event) const = 0; private: /// Main algorithmic implementation of the clusterization algorithm 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 61db9e9b67..adcc2a8caa 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 @@ -91,8 +91,8 @@ struct seed_parameter_estimation_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 exec::task await(vecmem::abstract_event& event) const = 0; private: /// Internal data type 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 12f99ae0f1..c472385baa 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 @@ -78,8 +78,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 exec::task await(vecmem::abstract_event& event) const = 0; }; // class silicon_pixel_spacepoint_formation_algorithm 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 08bbd4d9f3..3a241c5d4f 100644 --- a/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp +++ b/device/common/include/traccc/seeding/device/triplet_seeding_algorithm.hpp @@ -178,8 +178,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 exec::task await(vecmem::abstract_event& event) const = 0; private: /// Internal data type diff --git a/device/common/src/clusterization/clusterization_algorithm.cpp b/device/common/src/clusterization/clusterization_algorithm.cpp index a53038d426..1aca0c14f8 100644 --- a/device/common/src/clusterization/clusterization_algorithm.cpp +++ b/device/common/src/clusterization/clusterization_algorithm.cpp @@ -84,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); @@ -142,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); diff --git a/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp b/device/common/src/seeding/seed_parameter_estimation_algorithm.cpp index e90487f3ef..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); 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 5dece21b5b..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); diff --git a/device/common/src/seeding/triplet_seeding_algorithm.cpp b/device/common/src/seeding/triplet_seeding_algorithm.cpp index 4f9d91ee30..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); @@ -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); @@ -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); diff --git a/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp b/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp index aef6b0a730..4bf1a74687 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; + exec::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 da64a204e2..b5575a560f 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 @@ -53,7 +53,7 @@ struct seed_parameter_estimation_algorithm /// @} - void await(vecmem::abstract_event& event) const override; + exec::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..97bde804ee 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; + exec::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..b401d7f7ce 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; + exec::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..fc032ec3d9 100644 --- a/device/cuda/include/traccc/cuda/utils/algorithm_base.hpp +++ b/device/cuda/include/traccc/cuda/utils/algorithm_base.hpp @@ -16,6 +16,9 @@ // System include(s). #include +// Stdexec include(s). +#include + namespace traccc::cuda { /// Base class for all CUDA algorithms @@ -44,17 +47,17 @@ class algorithm_base { }; // class algorithm_base -using await_function_t = - std::function; +using await_function_t = std::function( + const cuda::stream&, vecmem::abstract_event&)>; -// Default await function, same as await_event_sync -void default_await_function(const cuda::stream& stream, +// Default await coroutine same as await_event_sync +exec::task default_await_function(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); +// Await coroutine that synchronizes the stream +exec::task await_stream_sync(const cuda::stream& stream, + vecmem::abstract_event& event); +// Await coroutine that waits on the event +exec::task 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 index 3a59b8b089..b4df01dbcd 100644 --- a/device/cuda/src/clusterization/clusterization_algorithm.cpp +++ b/device/cuda/src/clusterization/clusterization_algorithm.cpp @@ -40,8 +40,9 @@ void clusterization_algorithm::cluster_maker_kernel( warp_size()); } -void clusterization_algorithm::await(vecmem::abstract_event& event) const { - m_await_function(stream(), event); +exec::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/seeding/seed_parameter_estimation_algorithm.cpp b/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp index 7e308e3388..02a5763492 100644 --- a/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp +++ b/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp @@ -32,8 +32,8 @@ void seed_parameter_estimation_algorithm::estimate_seed_params_kernel( warp_size()); } -void seed_parameter_estimation_algorithm::await( +exec::task seed_parameter_estimation_algorithm::await( vecmem::abstract_event& event) const { - m_await_function(stream(), event); + co_await m_await_function(stream(), event); } } // 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 index 4234de4816..f50a9b33bd 100644 --- a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp +++ b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -30,9 +30,9 @@ void silicon_pixel_spacepoint_formation_algorithm::form_spacepoints_kernel( warp_size()); } -void silicon_pixel_spacepoint_formation_algorithm::await( +exec::task silicon_pixel_spacepoint_formation_algorithm::await( vecmem::abstract_event& event) const { - m_await_function(stream(), event); + co_await m_await_function(stream(), event); } } // namespace traccc::cuda diff --git a/device/cuda/src/seeding/triplet_seeding_algorithm.cpp b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp index 0b9fdc34b3..57da4055a1 100644 --- a/device/cuda/src/seeding/triplet_seeding_algorithm.cpp +++ b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp @@ -88,8 +88,9 @@ void triplet_seeding_algorithm::select_seeds_kernel( warp_size()); } -void triplet_seeding_algorithm::await(vecmem::abstract_event& event) const { - m_await_function(stream(), event); +exec::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/utils/algorithm_base.cpp b/device/cuda/src/utils/algorithm_base.cpp index ea9a5136ac..f3a3f567ad 100644 --- a/device/cuda/src/utils/algorithm_base.cpp +++ b/device/cuda/src/utils/algorithm_base.cpp @@ -25,17 +25,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); +exec::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&) { +exec::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) { +exec::task await_event_sync(const cuda::stream&, + vecmem::abstract_event& event) { event.wait(); + co_return; } } // namespace traccc::cuda From acc91a34ff792c42219a2eb8b7204e1e30968115 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 6 Feb 2026 11:32:39 +0100 Subject: [PATCH 11/15] add await policy suspending task until cuda stream is done --- examples/run/common/await_strategy.hpp | 5 +- examples/run/common/throughput_mt.ipp | 1 + examples/run/cuda/CMakeLists.txt | 5 +- examples/run/cuda/await_exec.cpp | 95 ++++++++++++++++++++++ examples/run/cuda/await_exec.hpp | 20 +++++ examples/run/cuda/full_chain_algorithm.cpp | 3 + 6 files changed, 126 insertions(+), 3 deletions(-) create mode 100644 examples/run/cuda/await_exec.cpp create mode 100644 examples/run/cuda/await_exec.hpp diff --git a/examples/run/common/await_strategy.hpp b/examples/run/common/await_strategy.hpp index a1aa61ab53..40658e6cb7 100644 --- a/examples/run/common/await_strategy.hpp +++ b/examples/run/common/await_strategy.hpp @@ -4,8 +4,9 @@ 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 }; } // namespace traccc diff --git a/examples/run/common/throughput_mt.ipp b/examples/run/common/throughput_mt.ipp index 228b301a66..168a416b38 100644 --- a/examples/run/common/throughput_mt.ipp +++ b/examples/run/common/throughput_mt.ipp @@ -212,6 +212,7 @@ 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: case opts::threading::await_strategy::defer_sync_event: case opts::threading::await_strategy::defer_sync_stream: diff --git a/examples/run/cuda/CMakeLists.txt b/examples/run/cuda/CMakeLists.txt index a1d31d7c54..302b857431 100644 --- a/examples/run/cuda/CMakeLists.txt +++ b/examples/run/cuda/CMakeLists.txt @@ -36,7 +36,10 @@ 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..1410a6b858 --- /dev/null +++ b/examples/run/cuda/await_exec.cpp @@ -0,0 +1,95 @@ +// Local include(s). +#include "await_exec.hpp" + +// Project include(s). +#include "traccc/cuda/utils/stream.hpp" + +// CUDA includes(s). +#include +#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< + stdexec::set_value_t(void), stdexec::set_error_t(std::exception_ptr)>; + + 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 { + try { + CUDA_ERROR_CHECK( + cudaLaunchHostFunc(m_stream, callback, &m_receiver)); + } catch (...) { + stdexec::set_error(std::move(m_receiver), std::current_exception()); + } + } + + private: + std::remove_cvref_t m_receiver; + cudaStream_t m_stream; + + static void callback(void* userData) { + auto& recv = *static_cast(userData); + try { + CUDA_ERROR_CHECK(cudaGetLastError()); + stdexec::set_value(std::move(recv)); + } catch (...) { + stdexec::set_error(std::move(recv), std::current_exception()); + return; + } + } +}; + +static_assert(stdexec::sender); + +exec::task await_callback(const cuda::stream& stream, + vecmem::abstract_event&) { + auto cuda_stream = static_cast(stream.cudaStream()); + co_await stream_await_sender{cuda_stream}; + co_return; +} + +} // 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..770e368aa7 --- /dev/null +++ b/examples/run/cuda/await_exec.hpp @@ -0,0 +1,20 @@ +#pragma once + +// Project include(s). +#include "traccc/cuda/utils/stream.hpp" + +// Vecmem include(s). +#include + +// Stdexec include(s). +#include + +namespace traccc::cuda { + +/// Await coroutine that returns a task which suspends execution with callback +/// until all asynchronous operations on the given stream are complete. +/// +exec::task await_callback(const cuda::stream& stream, + vecmem::abstract_event& event); + +} // namespace traccc::cuda diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 992a0d00c2..034efae9d5 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -10,6 +10,7 @@ #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" @@ -48,6 +49,8 @@ await_function_t get_await_function(await_strategy await_mode, return await_stream_sync; case await_strategy::sync_event: return await_event_sync; + case await_strategy::callback: + return await_callback; default: throw std::invalid_argument("Unknown await strategy"); } From da07f1f47fbd9f929f126268bcec32949a4c3695 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Wed, 25 Feb 2026 16:04:16 +0100 Subject: [PATCH 12/15] hide task type behind an alias --- core/CMakeLists.txt | 4 ++- core/include/traccc/execution/task.hpp | 10 ++++++ device/common/CMakeLists.txt | 2 +- .../device/clusterization_algorithm.hpp | 35 +++++++++---------- .../seed_parameter_estimation_algorithm.hpp | 17 ++++----- ...n_pixel_spacepoint_formation_algorithm.hpp | 8 ++--- .../device/triplet_seeding_algorithm.hpp | 8 ++--- .../clusterization_algorithm.cpp | 12 +++---- .../clusterization_algorithm.hpp | 2 +- .../seed_parameter_estimation_algorithm.hpp | 2 +- ...n_pixel_spacepoint_formation_algorithm.hpp | 2 +- .../seeding/triplet_seeding_algorithm.hpp | 2 +- .../traccc/cuda/utils/algorithm_base.hpp | 20 +++++------ .../clusterization_algorithm.cpp | 2 +- .../seed_parameter_estimation_algorithm.cpp | 2 +- ...n_pixel_spacepoint_formation_algorithm.cpp | 2 +- .../src/seeding/triplet_seeding_algorithm.cpp | 2 +- device/cuda/src/utils/algorithm_base.cpp | 13 +++---- examples/run/common/throughput_mt.ipp | 15 ++++---- examples/run/common/throughput_st.ipp | 10 +++--- examples/run/cpu/full_chain_algorithm.cpp | 6 ++-- examples/run/cpu/full_chain_algorithm.hpp | 8 ++--- examples/run/cuda/await_exec.cpp | 5 ++- examples/run/cuda/await_exec.hpp | 8 ++--- examples/run/cuda/full_chain_algorithm.cpp | 6 ++-- examples/run/cuda/full_chain_algorithm.hpp | 8 ++--- 26 files changed, 101 insertions(+), 110 deletions(-) create mode 100644 core/include/traccc/execution/task.hpp 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 b4662c22b0..4e96372dde 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -90,4 +90,4 @@ traccc_add_library( traccc_device_common device_common "include/traccc/fitting/device/impl/fill_fitting_sort_keys.ipp" ) target_link_libraries( traccc_device_common - PUBLIC traccc::core vecmem::core STDEXEC::stdexec ) + PUBLIC traccc::core vecmem::core ) diff --git a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp index 2322d8bee0..a8c0d349f1 100644 --- a/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp +++ b/device/common/include/traccc/clusterization/device/clusterization_algorithm.hpp @@ -12,6 +12,7 @@ #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" @@ -27,9 +28,6 @@ #include #include -// StdExec include(s). -#include - // System include(s). #include #include @@ -46,20 +44,20 @@ namespace traccc::device { /// class clusterization_algorithm : public algorithm< - exec::task::buffer>( + task::buffer>( const edm::silicon_cell_collection::const_view&, const silicon_detector_description::const_view&)>, public algorithm< - exec::task::buffer>( + task::buffer>( const edm::silicon_cell_collection::const_view&, const silicon_detector_description::const_view&, clustering_discard_disjoint_set&&)>, - public algorithm::buffer, - edm::silicon_cluster_collection::buffer>>( - const edm::silicon_cell_collection::const_view&, - const silicon_detector_description::const_view&, - clustering_keep_disjoint_set&&)>, + public algorithm< + task::buffer, + edm::silicon_cluster_collection::buffer>>( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&, + clustering_keep_disjoint_set&&)>, public messaging, public algorithm_base { @@ -87,18 +85,18 @@ class clusterization_algorithm /// @return a measurement collection (buffer) /// /// @{ - exec::task::buffer> operator()( + task::buffer> operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr) const override; - exec::task::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; - exec::task::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; @@ -139,13 +137,12 @@ class clusterization_algorithm /// @} /// Suspend execution until all asynchronous operations are done - virtual exec::task await(vecmem::abstract_event& event) const = 0; + virtual task await(vecmem::abstract_event& event) const = 0; private: /// Main algorithmic implementation of the clusterization algorithm - exec::task< - 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/seeding/device/seed_parameter_estimation_algorithm.hpp b/device/common/include/traccc/seeding/device/seed_parameter_estimation_algorithm.hpp index adcc2a8caa..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,7 @@ // 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) @@ -22,9 +23,6 @@ #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" -// Stdexec include(s). -#include - namespace traccc::device { /// Seed track parameter estimation algorithm @@ -33,12 +31,11 @@ namespace traccc::device { /// synchronisation statement is required before destroying this buffer. /// struct seed_parameter_estimation_algorithm - : public algorithm< - exec::task( - const magnetic_field&, - const edm::measurement_collection::const_view&, - const edm::spacepoint_collection::const_view&, - const edm::seed_collection::const_view&)>, + : public algorithm( + const magnetic_field&, + const edm::measurement_collection::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&)>, public messaging, public algorithm_base { @@ -92,7 +89,7 @@ struct seed_parameter_estimation_algorithm /// @} /// Suspend execution until all asynchronous operations are done - virtual exec::task await(vecmem::abstract_event& event) const = 0; + virtual task await(vecmem::abstract_event& event) const = 0; private: /// Internal data type 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 c472385baa..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,7 @@ // 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). @@ -19,9 +20,6 @@ #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" -// Stdexec include(s). -#include - namespace traccc::device { /// Algorithm forming space points out of measurements @@ -30,7 +28,7 @@ namespace traccc::device { /// measurements made on every detector module, into 3D spacepoint coordinates. /// class silicon_pixel_spacepoint_formation_algorithm - : public algorithm( + : public algorithm( const detector_buffer&, const edm::measurement_collection::const_view&)>, public messaging, @@ -79,7 +77,7 @@ class silicon_pixel_spacepoint_formation_algorithm /// @} /// Suspend execution until all asynchronous operations are done - virtual exec::task await(vecmem::abstract_event& event) const = 0; + virtual task await(vecmem::abstract_event& event) const = 0; }; // class silicon_pixel_spacepoint_formation_algorithm 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 3a241c5d4f..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,6 +14,7 @@ #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" @@ -25,9 +26,6 @@ #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" -// Stdexec include(s). -#include - // System include(s). #include @@ -39,7 +37,7 @@ namespace traccc::device { /// synchronisation statement is required before destroying this buffer. /// class triplet_seeding_algorithm - : public algorithm( + : public algorithm( const edm::spacepoint_collection::const_view&)>, public messaging, public algorithm_base { @@ -179,7 +177,7 @@ class triplet_seeding_algorithm /// @} /// Suspend execution until all asynchronous operations are done - virtual exec::task await(vecmem::abstract_event& event) const = 0; + virtual task await(vecmem::abstract_event& event) const = 0; private: /// Internal data type diff --git a/device/common/src/clusterization/clusterization_algorithm.cpp b/device/common/src/clusterization/clusterization_algorithm.cpp index 1aca0c14f8..907d866ca6 100644 --- a/device/common/src/clusterization/clusterization_algorithm.cpp +++ b/device/common/src/clusterization/clusterization_algorithm.cpp @@ -32,7 +32,7 @@ clusterization_algorithm::clusterization_algorithm( ->wait(); } -exec::task::buffer> +task::buffer> clusterization_algorithm::operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr) const { @@ -41,7 +41,7 @@ clusterization_algorithm::operator()( clustering_discard_disjoint_set{}); } -exec::task::buffer> +task::buffer> clusterization_algorithm::operator()( const edm::silicon_cell_collection::const_view& cells, const silicon_detector_description::const_view& det_descr, @@ -54,8 +54,8 @@ clusterization_algorithm::operator()( co_return std::move(res); } -exec::task::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, @@ -68,8 +68,8 @@ clusterization_algorithm::operator()( co_return {std::move(res), std::move(*djs)}; } -exec::task::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, diff --git a/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp b/device/cuda/include/traccc/cuda/clusterization/clusterization_algorithm.hpp index 4bf1a74687..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, /// @} - exec::task 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 b5575a560f..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 @@ -53,7 +53,7 @@ struct seed_parameter_estimation_algorithm /// @} - exec::task 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 97bde804ee..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 /// @} - exec::task 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 b401d7f7ce..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, /// @} - exec::task 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 fc032ec3d9..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 @@ -16,9 +17,6 @@ // System include(s). #include -// Stdexec include(s). -#include - namespace traccc::cuda { /// Base class for all CUDA algorithms @@ -47,17 +45,17 @@ class algorithm_base { }; // class algorithm_base -using await_function_t = std::function( - const cuda::stream&, vecmem::abstract_event&)>; +using await_function_t = + std::function(const cuda::stream&, vecmem::abstract_event&)>; // Default await coroutine same as await_event_sync -exec::task default_await_function(const cuda::stream& stream, - vecmem::abstract_event& event); +task default_await_function(const cuda::stream& stream, + vecmem::abstract_event& event); // Await coroutine that synchronizes the stream -exec::task await_stream_sync(const cuda::stream& stream, - vecmem::abstract_event& event); +task await_stream_sync(const cuda::stream& stream, + vecmem::abstract_event& event); // Await coroutine that waits on the event -exec::task await_event_sync(const cuda::stream& stream, - vecmem::abstract_event& event); +task 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 index b4df01dbcd..28e68183b9 100644 --- a/device/cuda/src/clusterization/clusterization_algorithm.cpp +++ b/device/cuda/src/clusterization/clusterization_algorithm.cpp @@ -40,7 +40,7 @@ void clusterization_algorithm::cluster_maker_kernel( warp_size()); } -exec::task clusterization_algorithm::await( +task clusterization_algorithm::await( vecmem::abstract_event& event) const { co_await m_await_function(stream(), event); } diff --git a/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp b/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp index 02a5763492..b29f22851d 100644 --- a/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp +++ b/device/cuda/src/seeding/seed_parameter_estimation_algorithm.cpp @@ -32,7 +32,7 @@ void seed_parameter_estimation_algorithm::estimate_seed_params_kernel( warp_size()); } -exec::task seed_parameter_estimation_algorithm::await( +task seed_parameter_estimation_algorithm::await( vecmem::abstract_event& event) const { co_await m_await_function(stream(), event); } diff --git a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp index f50a9b33bd..670aedb069 100644 --- a/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp +++ b/device/cuda/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -30,7 +30,7 @@ void silicon_pixel_spacepoint_formation_algorithm::form_spacepoints_kernel( warp_size()); } -exec::task silicon_pixel_spacepoint_formation_algorithm::await( +task silicon_pixel_spacepoint_formation_algorithm::await( vecmem::abstract_event& event) const { co_await m_await_function(stream(), event); } diff --git a/device/cuda/src/seeding/triplet_seeding_algorithm.cpp b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp index 57da4055a1..692d07f307 100644 --- a/device/cuda/src/seeding/triplet_seeding_algorithm.cpp +++ b/device/cuda/src/seeding/triplet_seeding_algorithm.cpp @@ -88,7 +88,7 @@ void triplet_seeding_algorithm::select_seeds_kernel( warp_size()); } -exec::task triplet_seeding_algorithm::await( +task triplet_seeding_algorithm::await( vecmem::abstract_event& event) const { co_await m_await_function(stream(), event); } diff --git a/device/cuda/src/utils/algorithm_base.cpp b/device/cuda/src/utils/algorithm_base.cpp index f3a3f567ad..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,20 +26,20 @@ unsigned int algorithm_base::warp_size() const { return m_warp_size; } -exec::task default_await_function(const cuda::stream&, - vecmem::abstract_event& event) { +task default_await_function(const cuda::stream&, + vecmem::abstract_event& event) { event.wait(); co_return; } -exec::task 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; } -exec::task 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; } diff --git a/examples/run/common/throughput_mt.ipp b/examples/run/common/throughput_mt.ipp index 168a416b38..0b8ea3610c 100644 --- a/examples/run/common/throughput_mt.ipp +++ b/examples/run/common/throughput_mt.ipp @@ -15,6 +15,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" @@ -58,7 +59,6 @@ // Stdexec include(s). #include -#include // Indicators include(s). #include @@ -241,14 +241,13 @@ 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::vector&, int, - const edm::silicon_cell_collection::host&)> + std::function(std::vector&, int, + const edm::silicon_cell_collection::host&)> process_event; if (throughput_opts.reco_stage == opts::throughput::stage::seeding) { process_event = [](std::vector& algs_, int slot_, const edm::silicon_cell_collection::host& cells_) - -> exec::task { + -> task { auto result = co_await algs_.at(static_cast(slot_)) .seeding(cells_); co_return result.size(); @@ -256,7 +255,7 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { } else if (throughput_opts.reco_stage == opts::throughput::stage::full) { process_event = [](std::vector& algs_, int slot_, const edm::silicon_cell_collection::host& cells_) - -> exec::task { + -> task { auto result = co_await algs_.at(static_cast(slot_))(cells_); co_return result.size(); @@ -314,7 +313,7 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { auto payload = [](auto& algs_, auto& input_, auto& progress_bar_, auto& rec_track_params_, auto& queue_, size_t event_, size_t slot_, - auto& process_event_) -> exec::task { + auto& process_event_) -> task { auto result = co_await process_event_( algs_, static_cast(slot_), input_.at(event_)); rec_track_params_.fetch_add(result); @@ -362,7 +361,7 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { auto payload = [](auto& algs_, auto& input_, auto& progress_bar_, auto& rec_track_params_, auto& queue_, size_t event_, size_t slot_, - auto& process_event_) -> exec::task { + auto& process_event_) -> task { auto result = co_await process_event_( algs_, static_cast(slot_), input_.at(event_)); rec_track_params_.fetch_add(result); diff --git a/examples/run/common/throughput_st.ipp b/examples/run/common/throughput_st.ipp index a8d881cee3..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" @@ -48,7 +49,6 @@ // Stdexec include(s). #include -#include #include // System include(s). @@ -164,20 +164,20 @@ 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( - FULL_CHAIN_ALG*, const edm::silicon_cell_collection::host&)> + std::function(FULL_CHAIN_ALG*, + const edm::silicon_cell_collection::host&)> process_event; if (throughput_opts.reco_stage == opts::throughput::stage::seeding) { process_event = [](FULL_CHAIN_ALG* alg_, const edm::silicon_cell_collection::host& cells_) - -> exec::task { + -> task { auto result = co_await alg_->seeding(cells_); co_return result.size(); }; } else if (throughput_opts.reco_stage == opts::throughput::stage::full) { process_event = [](FULL_CHAIN_ALG* alg_, const edm::silicon_cell_collection::host& cells_) - -> exec::task { + -> task { auto result = co_await (*alg_)(cells_); co_return result.size(); }; diff --git a/examples/run/cpu/full_chain_algorithm.cpp b/examples/run/cpu/full_chain_algorithm.cpp index 4c82acbab8..96b9d77633 100644 --- a/examples/run/cpu/full_chain_algorithm.cpp +++ b/examples/run/cpu/full_chain_algorithm.cpp @@ -8,8 +8,8 @@ // Local include(s). #include "full_chain_algorithm.hpp" -// Stdexec include(s). -#include +// Project include(s). +#include "traccc/execution/task.hpp" namespace traccc { @@ -98,7 +98,7 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( } } -exec::task +task full_chain_algorithm::seeding( const edm::silicon_cell_collection::host& cells) const { diff --git a/examples/run/cpu/full_chain_algorithm.hpp b/examples/run/cpu/full_chain_algorithm.hpp index a9dc57c57f..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" @@ -34,9 +35,6 @@ #include #include -// Stdexec include(s). -#include - // System include(s). #include #include @@ -49,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 { @@ -106,7 +104,7 @@ class full_chain_algorithm /// @param cells The cells for every detector module in the event /// @return A task returning the track seeds reconstructed /// - exec::task seeding( + task seeding( const edm::silicon_cell_collection::host& cells) const; private: diff --git a/examples/run/cuda/await_exec.cpp b/examples/run/cuda/await_exec.cpp index 1410a6b858..521259a4ad 100644 --- a/examples/run/cuda/await_exec.cpp +++ b/examples/run/cuda/await_exec.cpp @@ -3,13 +3,13 @@ // Project include(s). #include "traccc/cuda/utils/stream.hpp" +#include "traccc/execution/task.hpp" // CUDA includes(s). #include #include // Stdexec include(s). -#include #include #define CUDA_ERROR_CHECK(EXP) \ @@ -85,8 +85,7 @@ class stream_await_sender::stream_await_operation { static_assert(stdexec::sender); -exec::task await_callback(const cuda::stream& stream, - vecmem::abstract_event&) { +task await_callback(const cuda::stream& stream, vecmem::abstract_event&) { auto cuda_stream = static_cast(stream.cudaStream()); co_await stream_await_sender{cuda_stream}; co_return; diff --git a/examples/run/cuda/await_exec.hpp b/examples/run/cuda/await_exec.hpp index 770e368aa7..7fd8a9a6a2 100644 --- a/examples/run/cuda/await_exec.hpp +++ b/examples/run/cuda/await_exec.hpp @@ -2,19 +2,17 @@ // Project include(s). #include "traccc/cuda/utils/stream.hpp" +#include "traccc/execution/task.hpp" // Vecmem include(s). #include -// Stdexec include(s). -#include - namespace traccc::cuda { /// Await coroutine that returns a task which suspends execution with callback /// until all asynchronous operations on the given stream are complete. /// -exec::task await_callback(const cuda::stream& stream, - vecmem::abstract_event& event); +task await_callback(const cuda::stream& stream, + vecmem::abstract_event& event); } // namespace traccc::cuda diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 034efae9d5..37f99fa5fa 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -15,6 +15,7 @@ // 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). @@ -23,9 +24,6 @@ // CUDA include(s). #include -// Stdexec include(s). -#include - // System include(s). #include #include @@ -250,7 +248,7 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( } } -exec::task +task full_chain_algorithm::seeding( const edm::silicon_cell_collection::host& cells) const { diff --git a/examples/run/cuda/full_chain_algorithm.hpp b/examples/run/cuda/full_chain_algorithm.hpp index 2f916cebad..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" @@ -42,9 +43,6 @@ #include #include -// Stdexec include(s). -#include - // System include(s). #include #include @@ -59,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 { @@ -126,7 +124,7 @@ class full_chain_algorithm /// @param cells The cells for every detector module in the event /// @return A task returning the track seeds reconstructed /// - exec::task seeding( + task seeding( const edm::silicon_cell_collection::host& cells) const; private: From 1af3ff7e4da4618b67c325578f675d8e4cdfb8b1 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Mon, 2 Mar 2026 11:28:21 +0100 Subject: [PATCH 13/15] move cuda error check outside of sender --- examples/run/cuda/await_exec.cpp | 26 ++++++++++---------------- 1 file changed, 10 insertions(+), 16 deletions(-) diff --git a/examples/run/cuda/await_exec.cpp b/examples/run/cuda/await_exec.cpp index 521259a4ad..e84153942c 100644 --- a/examples/run/cuda/await_exec.cpp +++ b/examples/run/cuda/await_exec.cpp @@ -32,8 +32,8 @@ class stream_await_sender { class stream_await_operation; using sender_concept = stdexec::sender_t; - using completion_signatures = stdexec::completion_signatures< - stdexec::set_value_t(void), stdexec::set_error_t(std::exception_ptr)>; + using completion_signatures = + stdexec::completion_signatures; stream_await_sender(const cudaStream_t stream) : m_stream(stream) {} stdexec::env<> get_env() const noexcept { return {}; } @@ -59,11 +59,11 @@ class stream_await_sender::stream_await_operation { : m_receiver(std::forward(recv)), m_stream(stream) {} void start() & noexcept { - try { - CUDA_ERROR_CHECK( - cudaLaunchHostFunc(m_stream, callback, &m_receiver)); - } catch (...) { - stdexec::set_error(std::move(m_receiver), std::current_exception()); + + 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); } } @@ -71,15 +71,9 @@ class stream_await_sender::stream_await_operation { std::remove_cvref_t m_receiver; cudaStream_t m_stream; - static void callback(void* userData) { + static void callback(void* userData) noexcept { auto& recv = *static_cast(userData); - try { - CUDA_ERROR_CHECK(cudaGetLastError()); - stdexec::set_value(std::move(recv)); - } catch (...) { - stdexec::set_error(std::move(recv), std::current_exception()); - return; - } + stdexec::set_value(recv, cudaSuccess); } }; @@ -87,7 +81,7 @@ static_assert(stdexec::sender); task await_callback(const cuda::stream& stream, vecmem::abstract_event&) { auto cuda_stream = static_cast(stream.cudaStream()); - co_await stream_await_sender{cuda_stream}; + CUDA_ERROR_CHECK(co_await stream_await_sender{cuda_stream}); co_return; } From 9cfbeda3bb926464be1ee90b5e7955e006517b30 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Sun, 19 Apr 2026 14:53:48 +0200 Subject: [PATCH 14/15] add polling and defer sync --- examples/run/common/await_strategy.hpp | 11 ++- examples/run/common/throughput_mt.ipp | 5 +- examples/run/cuda/await_exec.cpp | 22 +++++- examples/run/cuda/await_exec.hpp | 28 ++++++- examples/run/cuda/full_chain_algorithm.cpp | 10 ++- examples/utils/CMakeLists.txt | 2 +- .../traccc/examples/utils/threadpool.hpp | 77 +++++++++++++++++++ examples/utils/src/threadpool.cpp | 21 +++++ 8 files changed, 164 insertions(+), 12 deletions(-) diff --git a/examples/run/common/await_strategy.hpp b/examples/run/common/await_strategy.hpp index 40658e6cb7..82134cedd1 100644 --- a/examples/run/common/await_strategy.hpp +++ b/examples/run/common/await_strategy.hpp @@ -4,9 +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 - callback ///< Suspending on a stream with a callback + 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/throughput_mt.ipp b/examples/run/common/throughput_mt.ipp index 0b8ea3610c..b6eb7c6524 100644 --- a/examples/run/common/throughput_mt.ipp +++ b/examples/run/common/throughput_mt.ipp @@ -214,10 +214,11 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { 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"); } diff --git a/examples/run/cuda/await_exec.cpp b/examples/run/cuda/await_exec.cpp index e84153942c..b318dff130 100644 --- a/examples/run/cuda/await_exec.cpp +++ b/examples/run/cuda/await_exec.cpp @@ -2,14 +2,15 @@ #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 -#include // Stdexec include(s). +#include #include #define CUDA_ERROR_CHECK(EXP) \ @@ -79,10 +80,29 @@ class stream_await_sender::stream_await_operation { 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 index 7fd8a9a6a2..f9cbc71fb3 100644 --- a/examples/run/cuda/await_exec.hpp +++ b/examples/run/cuda/await_exec.hpp @@ -2,6 +2,7 @@ // Project include(s). #include "traccc/cuda/utils/stream.hpp" +#include "traccc/examples/utils/threadpool.hpp" #include "traccc/execution/task.hpp" // Vecmem include(s). @@ -9,10 +10,31 @@ namespace traccc::cuda { -/// Await coroutine that returns a task which suspends execution with callback -/// until all asynchronous operations on the given stream are complete. -/// +/// 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 37f99fa5fa..7339aa7bd2 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -40,8 +40,8 @@ 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; @@ -49,6 +49,12 @@ await_function_t get_await_function(await_strategy await_mode, 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"); } 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 From cf22e443d6e4e3f401b12453aaf093874e540b29 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Tue, 18 Aug 2026 12:50:51 +0200 Subject: [PATCH 15/15] fix missing CUDART_CB in callback --- examples/run/cuda/await_exec.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/run/cuda/await_exec.cpp b/examples/run/cuda/await_exec.cpp index b318dff130..e64be07362 100644 --- a/examples/run/cuda/await_exec.cpp +++ b/examples/run/cuda/await_exec.cpp @@ -72,7 +72,7 @@ class stream_await_sender::stream_await_operation { std::remove_cvref_t m_receiver; cudaStream_t m_stream; - static void callback(void* userData) noexcept { + static void CUDART_CB callback(void* userData) noexcept { auto& recv = *static_cast(userData); stdexec::set_value(recv, cudaSuccess); }