From c29320f2f5e3af2472f4bd2a832c6a1060a3d225 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 8 Jan 2025 15:30:05 +0100 Subject: [PATCH 1/5] Harmonized the different thread_id classes with each other. Made all of them into "private headers", and added automated tests that they would fulfill the appropriate concept. --- .../include/traccc/alpaka/utils/thread_id.hpp | 53 --------------- .../clusterization_algorithm.cpp | 6 +- device/alpaka/src/utils/thread_id.hpp | 66 +++++++++++++++++++ .../clusterization/device/impl/ccl_kernel.ipp | 3 +- .../include/traccc/cuda/utils/thread_id.hpp | 41 ------------ .../clusterization_algorithm.cu | 4 +- .../src/clusterization/kernels/ccl_kernel.cu | 4 +- device/cuda/src/finding/finding_algorithm.cu | 2 +- .../cuda/src/finding/kernels/find_tracks.cuh | 8 +-- .../find_tracks_default_detector.cu | 6 +- .../specializations/find_tracks_src.cuh | 15 +++-- device/cuda/src/utils/thread_id.hpp | 51 ++++++++++++++ device/sycl/src/utils/thread_id.hpp | 30 ++++++--- tests/cuda/test_sort.cu | 6 +- 14 files changed, 166 insertions(+), 129 deletions(-) delete mode 100644 device/alpaka/include/traccc/alpaka/utils/thread_id.hpp create mode 100644 device/alpaka/src/utils/thread_id.hpp delete mode 100644 device/cuda/include/traccc/cuda/utils/thread_id.hpp create mode 100644 device/cuda/src/utils/thread_id.hpp diff --git a/device/alpaka/include/traccc/alpaka/utils/thread_id.hpp b/device/alpaka/include/traccc/alpaka/utils/thread_id.hpp deleted file mode 100644 index 475f452c9b..0000000000 --- a/device/alpaka/include/traccc/alpaka/utils/thread_id.hpp +++ /dev/null @@ -1,53 +0,0 @@ -/** - * traccc library, part of the ACTS project (R&D line) - * - * (c) 2024 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -#include - -#include "traccc/definitions/qualifiers.hpp" - -namespace traccc::alpaka { -template -struct thread_id1 { - TRACCC_DEVICE thread_id1(const Acc& acc) : m_acc(acc) {} - - auto inline TRACCC_DEVICE getLocalThreadId() const { - return ::alpaka::getIdx<::alpaka::Block, ::alpaka::Threads>(m_acc)[0u]; - } - - auto inline TRACCC_DEVICE getLocalThreadIdX() const { - return getLocalThreadId(); - } - - auto inline TRACCC_DEVICE getGlobalThreadId() const { - return getLocalThreadId() + getBlockIdX() * getBlockDimX(); - } - - auto inline TRACCC_DEVICE getGlobalThreadIdX() const { - return getLocalThreadId() + getBlockIdX() * getBlockDimX(); - } - - auto inline TRACCC_DEVICE getBlockIdX() const { - return ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Blocks>(m_acc)[0u]; - } - - auto inline TRACCC_DEVICE getBlockDimX() const { - return ::alpaka::getWorkDiv<::alpaka::Block, ::alpaka::Threads>( - m_acc)[0u]; - } - - auto inline TRACCC_DEVICE getGridDimX() const { - return ::alpaka::getWorkDiv<::alpaka::Grid, ::alpaka::Blocks>( - m_acc)[0u]; - } - - private: - const Acc& m_acc; -}; -} // namespace traccc::alpaka diff --git a/device/alpaka/src/clusterization/clusterization_algorithm.cpp b/device/alpaka/src/clusterization/clusterization_algorithm.cpp index 68b60f6dc6..3fca603798 100644 --- a/device/alpaka/src/clusterization/clusterization_algorithm.cpp +++ b/device/alpaka/src/clusterization/clusterization_algorithm.cpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -9,10 +9,10 @@ #include "traccc/alpaka/clusterization/clusterization_algorithm.hpp" #include "../utils/barrier.hpp" +#include "../utils/thread_id.hpp" #include "../utils/utils.hpp" // Project include(s) -#include "traccc/alpaka/utils/thread_id.hpp" #include "traccc/clusterization/clustering_config.hpp" #include "traccc/clusterization/device/ccl_kernel.hpp" @@ -36,7 +36,7 @@ struct CCLKernel { measurement_collection_types::view measurements_view, vecmem::data::vector_view cell_links) const { - traccc::alpaka::thread_id1 thread_id(acc); + details::thread_id1 thread_id(acc); auto& partition_start = ::alpaka::declareSharedVar(acc); diff --git a/device/alpaka/src/utils/thread_id.hpp b/device/alpaka/src/utils/thread_id.hpp new file mode 100644 index 0000000000..832ca5a432 --- /dev/null +++ b/device/alpaka/src/utils/thread_id.hpp @@ -0,0 +1,66 @@ +/** + * traccc library, part of the ACTS project (R&D line) + * + * (c) 2024-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Local include(s). +#include "utils.hpp" + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" +#include "traccc/device/concepts/thread_id.hpp" + +// Alpaka include(s). +#include + +namespace traccc::alpaka::details { + +/// An Alpaka thread identifier type +template +struct thread_id1 { + TRACCC_HOST_DEVICE explicit thread_id1(const Acc& acc) : m_acc(acc) {} + + auto inline TRACCC_HOST_DEVICE getLocalThreadId() const { + return ::alpaka::getIdx<::alpaka::Block, ::alpaka::Threads>(m_acc)[0u]; + } + + auto inline TRACCC_HOST_DEVICE getLocalThreadIdX() const { + return getLocalThreadId(); + } + + auto inline TRACCC_HOST_DEVICE getGlobalThreadId() const { + return getLocalThreadId() + getBlockIdX() * getBlockDimX(); + } + + auto inline TRACCC_HOST_DEVICE getGlobalThreadIdX() const { + return getLocalThreadId() + getBlockIdX() * getBlockDimX(); + } + + auto inline TRACCC_HOST_DEVICE getBlockIdX() const { + return ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Blocks>(m_acc)[0u]; + } + + auto inline TRACCC_HOST_DEVICE getBlockDimX() const { + return ::alpaka::getWorkDiv<::alpaka::Block, ::alpaka::Threads>( + m_acc)[0u]; + } + + auto inline TRACCC_HOST_DEVICE getGridDimX() const { + return ::alpaka::getWorkDiv<::alpaka::Grid, ::alpaka::Blocks>( + m_acc)[0u]; + } + + private: + const Acc& m_acc; +}; + +/// Verify that @c traccc::alpaka::details::thread_id1 fulfills the +/// @c traccc::device::concepts::thread_id1 concept. +static_assert(traccc::device::concepts::thread_id1>); + +} // namespace traccc::alpaka::details diff --git a/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp b/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp index 93268a970a..aadfcaf633 100644 --- a/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp +++ b/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp @@ -254,8 +254,7 @@ TRACCC_DEVICE inline void ccl_kernel( */ if (thread_id.getLocalThreadIdX() == 0) { unsigned int start = - static_cast(thread_id.getBlockIdX()) * - cfg.target_partition_size(); + thread_id.getBlockIdX() * cfg.target_partition_size(); assert(start < num_cells); unsigned int end = std::min(num_cells, start + cfg.target_partition_size()); diff --git a/device/cuda/include/traccc/cuda/utils/thread_id.hpp b/device/cuda/include/traccc/cuda/utils/thread_id.hpp deleted file mode 100644 index ede3216652..0000000000 --- a/device/cuda/include/traccc/cuda/utils/thread_id.hpp +++ /dev/null @@ -1,41 +0,0 @@ -/** - * traccc library, part of the ACTS project (R&D line) - * - * (c) 2024 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -#include - -#include "traccc/definitions/qualifiers.hpp" - -namespace traccc::cuda { -struct thread_id1 { - TRACCC_DEVICE thread_id1() {} - - std::size_t inline TRACCC_DEVICE getLocalThreadId() const { - return threadIdx.x; - } - - std::size_t inline TRACCC_DEVICE getLocalThreadIdX() const { - return threadIdx.x; - } - - std::size_t inline TRACCC_DEVICE getGlobalThreadId() const { - return threadIdx.x + blockIdx.x * blockDim.x; - } - - std::size_t inline TRACCC_DEVICE getGlobalThreadIdX() const { - return threadIdx.x + blockIdx.x * blockDim.x; - } - - std::size_t inline TRACCC_DEVICE getBlockIdX() const { return blockIdx.x; } - - std::size_t inline TRACCC_DEVICE getBlockDimX() const { return blockDim.x; } - - std::size_t inline TRACCC_DEVICE getGridDimX() const { return gridDim.x; } -}; -} // namespace traccc::cuda diff --git a/device/cuda/src/clusterization/clusterization_algorithm.cu b/device/cuda/src/clusterization/clusterization_algorithm.cu index 60a0848a86..75972fdb21 100644 --- a/device/cuda/src/clusterization/clusterization_algorithm.cu +++ b/device/cuda/src/clusterization/clusterization_algorithm.cu @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,12 +10,12 @@ #include "../sanity/ordered_on.cuh" #include "../utils/barrier.hpp" #include "../utils/cuda_error_handling.hpp" +#include "../utils/thread_id.hpp" #include "../utils/utils.hpp" #include "./kernels/ccl_kernel.cuh" #include "traccc/clusterization/clustering_config.hpp" #include "traccc/clusterization/device/ccl_kernel_definitions.hpp" #include "traccc/cuda/clusterization/clusterization_algorithm.hpp" -#include "traccc/cuda/utils/thread_id.hpp" #include "traccc/utils/projections.hpp" #include "traccc/utils/relations.hpp" diff --git a/device/cuda/src/clusterization/kernels/ccl_kernel.cu b/device/cuda/src/clusterization/kernels/ccl_kernel.cu index ca10b1a534..ac1313f45b 100644 --- a/device/cuda/src/clusterization/kernels/ccl_kernel.cu +++ b/device/cuda/src/clusterization/kernels/ccl_kernel.cu @@ -10,11 +10,11 @@ #include "../../sanity/ordered_on.cuh" #include "../../utils/barrier.hpp" #include "../../utils/cuda_error_handling.hpp" +#include "../../utils/thread_id.hpp" #include "../../utils/utils.hpp" #include "traccc/clusterization/clustering_config.hpp" #include "traccc/clusterization/device/ccl_kernel_definitions.hpp" #include "traccc/cuda/clusterization/clusterization_algorithm.hpp" -#include "traccc/cuda/utils/thread_id.hpp" #include "traccc/utils/projections.hpp" #include "traccc/utils/relations.hpp" @@ -54,7 +54,7 @@ __global__ void ccl_kernel( static_cast(cfg.max_partition_size()), shared_v + cfg.max_partition_size()}; traccc::cuda::barrier barry_r; - const cuda::thread_id1 thread_id; + const details::thread_id1 thread_id; device::ccl_kernel(cfg, thread_id, cells_view, det_descr_view, partition_start, partition_end, outi, f_view, gf_view, diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu index d46dc2ff86..f9d333659d 100644 --- a/device/cuda/src/finding/finding_algorithm.cu +++ b/device/cuda/src/finding/finding_algorithm.cu @@ -9,6 +9,7 @@ #include "../sanity/contiguous_on.cuh" #include "../utils/barrier.hpp" #include "../utils/cuda_error_handling.hpp" +#include "../utils/thread_id.hpp" #include "../utils/utils.hpp" #include "./kernels/apply_interaction.cuh" #include "./kernels/build_tracks.cuh" @@ -18,7 +19,6 @@ #include "./kernels/propagate_to_next_surface.cuh" #include "./kernels/prune_tracks.cuh" #include "traccc/cuda/finding/finding_algorithm.hpp" -#include "traccc/cuda/utils/thread_id.hpp" #include "traccc/definitions/primitives.hpp" #include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/device/sort_key.hpp" diff --git a/device/cuda/src/finding/kernels/find_tracks.cuh b/device/cuda/src/finding/kernels/find_tracks.cuh index 9420e341f8..6dee2853ca 100644 --- a/device/cuda/src/finding/kernels/find_tracks.cuh +++ b/device/cuda/src/finding/kernels/find_tracks.cuh @@ -1,17 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "../../utils/barrier.hpp" -#include "traccc/cuda/utils/thread_id.hpp" -#include "traccc/edm/track_parameters.hpp" +// Project include(s). #include "traccc/finding/device/find_tracks.hpp" -#include "traccc/geometry/detector.hpp" +#include "traccc/finding/finding_config.hpp" namespace traccc::cuda::kernels { diff --git a/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu b/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu index c2bb3ba910..9695090c63 100644 --- a/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu +++ b/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu @@ -1,12 +1,16 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ +// Local include(s). #include "find_tracks_src.cuh" +// Project include(s). +#include "traccc/geometry/detector.hpp" + namespace traccc::cuda::kernels { template __global__ void find_tracks( const finding_config cfg, diff --git a/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh index 51ffe27ac0..60b4b83a62 100644 --- a/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh +++ b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh @@ -7,12 +7,15 @@ #pragma once -#include "../../../utils/barrier.hpp" -#include "../propagate_to_next_surface.cuh" -#include "traccc/cuda/utils/thread_id.hpp" -#include "traccc/edm/track_parameters.hpp" +// Project include(s). #include "traccc/finding/device/find_tracks.hpp" -#include "traccc/geometry/detector.hpp" + +// Local include(s). +#include "../../../utils/barrier.hpp" +#include "../../../utils/thread_id.hpp" + +// System include(s). +#include namespace traccc::cuda::kernels { @@ -27,7 +30,7 @@ __global__ void find_tracks(const finding_config cfg, &shared_num_candidates[blockDim.x]); cuda::barrier barrier; - cuda::thread_id1 thread_id; + details::thread_id1 thread_id; device::find_tracks( thread_id, barrier, cfg, payload, diff --git a/device/cuda/src/utils/thread_id.hpp b/device/cuda/src/utils/thread_id.hpp new file mode 100644 index 0000000000..c913631103 --- /dev/null +++ b/device/cuda/src/utils/thread_id.hpp @@ -0,0 +1,51 @@ +/** + * traccc library, part of the ACTS project (R&D line) + * + * (c) 2024-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" +#include "traccc/device/concepts/thread_id.hpp" + +namespace traccc::cuda::details { + +/// A CUDA thread identifier type +struct thread_id1 { + TRACCC_DEVICE thread_id1() {} + + inline unsigned int TRACCC_DEVICE getLocalThreadId() const { + return threadIdx.x; + } + + inline unsigned int TRACCC_DEVICE getLocalThreadIdX() const { + return threadIdx.x; + } + + inline unsigned int TRACCC_DEVICE getGlobalThreadId() const { + return threadIdx.x + blockIdx.x * blockDim.x; + } + + inline unsigned int TRACCC_DEVICE getGlobalThreadIdX() const { + return threadIdx.x + blockIdx.x * blockDim.x; + } + + inline unsigned int TRACCC_DEVICE getBlockIdX() const { return blockIdx.x; } + + inline unsigned int TRACCC_DEVICE getBlockDimX() const { + return blockDim.x; + } + + inline unsigned int TRACCC_DEVICE getGridDimX() const { return gridDim.x; } + +}; // struct thread_id1 + +/// Verify that @c traccc::cuda::details::thread_id1 fulfills the +/// @c traccc::device::concepts::thread_id1 concept. +static_assert(traccc::device::concepts::thread_id1); + +} // namespace traccc::cuda::details diff --git a/device/sycl/src/utils/thread_id.hpp b/device/sycl/src/utils/thread_id.hpp index 8e62ae135a..585aeef4b2 100644 --- a/device/sycl/src/utils/thread_id.hpp +++ b/device/sycl/src/utils/thread_id.hpp @@ -1,7 +1,7 @@ /** * traccc library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -33,23 +33,33 @@ requires(DIMENSIONS >= 1 && DIMENSIONS <= 3) struct thread_id { /// @name Function(s) implementing @c traccc::device::concepts::thread_id1 /// @{ - inline auto getLocalThreadId() const { - return m_item.get_local_linear_id(); + inline unsigned int getLocalThreadId() const { + return static_cast(m_item.get_local_linear_id()); } - inline auto getLocalThreadIdX() const { return m_item.get_local_id(0); } + inline unsigned int getLocalThreadIdX() const { + return static_cast(m_item.get_local_id(0)); + } - inline auto getGlobalThreadId() const { - return m_item.get_global_linear_id(); + inline unsigned int getGlobalThreadId() const { + return static_cast(m_item.get_global_linear_id()); } - inline auto getGlobalThreadIdX() const { return m_item.get_global_id(0); } + inline unsigned int getGlobalThreadIdX() const { + return static_cast(m_item.get_global_id(0)); + } - inline auto getBlockIdX() const { return m_item.get_group(0); } + inline unsigned int getBlockIdX() const { + return static_cast(m_item.get_group(0)); + } - inline auto getBlockDimX() const { return m_item.get_local_range(0); } + inline unsigned int getBlockDimX() const { + return static_cast(m_item.get_local_range(0)); + } - inline auto getGridDimX() const { return m_item.get_global_range(0); } + inline unsigned int getGridDimX() const { + return static_cast(m_item.get_global_range(0)); + } /// @} diff --git a/tests/cuda/test_sort.cu b/tests/cuda/test_sort.cu index dc657070a8..f16ed13500 100644 --- a/tests/cuda/test_sort.cu +++ b/tests/cuda/test_sort.cu @@ -1,7 +1,7 @@ /** * traccc library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -12,11 +12,11 @@ #include #include "../../cuda/src/utils/barrier.hpp" -#include "traccc/cuda/utils/thread_id.hpp" +#include "../../cuda/src/utils/thread_id.hpp" #include "traccc/device/sort.hpp" __global__ void testBlockSortKernel(uint32_t *keys, uint32_t n_keys) { - traccc::cuda::thread_id1 thread_id; + traccc::cuda::details::thread_id1 thread_id; traccc::cuda::barrier barrier; traccc::device::blockOddEvenSort(thread_id, barrier, keys, n_keys, std::less()); From ceb4a81369f356ef2f6d74406b0598b8a84cc427 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 8 Jan 2025 23:22:17 +0100 Subject: [PATCH 2/5] Synchronized how thread IDs would be used in the device functions. While also cleaning up the includes of the files a bit. --- .../clusterization/device/impl/ccl_kernel.ipp | 2 +- .../traccc/finding/device/find_tracks.hpp | 10 ++++++---- .../traccc/finding/device/impl/find_tracks.ipp | 4 ++-- .../device/impl/propagate_to_next_surface.ipp | 3 ++- .../device/propagate_to_next_surface.hpp | 1 + .../traccc/fitting/device/fill_sort_keys.hpp | 12 ++++++------ .../include/traccc/fitting/device/fit.hpp | 10 +++++----- .../fitting/device/impl/fill_sort_keys.ipp | 15 ++++++--------- .../include/traccc/fitting/device/impl/fit.ipp | 7 +++---- .../traccc/seeding/device/count_doublets.hpp | 14 +++++++------- .../seeding/device/count_grid_capacities.hpp | 12 ++++++------ .../traccc/seeding/device/count_triplets.hpp | 13 +++++++------ .../seeding/device/estimate_track_params.hpp | 7 +++++-- .../traccc/seeding/device/find_doublets.hpp | 14 +++++++------- .../traccc/seeding/device/find_triplets.hpp | 14 +++++++------- .../traccc/seeding/device/form_spacepoints.hpp | 10 +++++----- .../seeding/device/impl/count_doublets.ipp | 7 +++---- .../device/impl/count_grid_capacities.ipp | 7 +++---- .../seeding/device/impl/count_triplets.ipp | 7 +++---- .../device/impl/estimate_track_params.ipp | 9 ++++----- .../seeding/device/impl/find_doublets.ipp | 17 +++++++---------- .../seeding/device/impl/find_triplets.ipp | 9 ++++----- .../seeding/device/impl/form_spacepoints.ipp | 6 +++--- .../seeding/device/impl/populate_grid.ipp | 4 ++-- .../device/impl/reduce_triplet_counts.ipp | 10 ++++------ .../traccc/seeding/device/impl/select_seeds.ipp | 14 +++++++------- .../device/impl/update_triplet_weights.ipp | 10 ++++------ .../traccc/seeding/device/populate_grid.hpp | 15 ++++++--------- .../seeding/device/reduce_triplet_counts.hpp | 9 ++++++--- .../traccc/seeding/device/select_seeds.hpp | 17 ++++++++--------- .../seeding/device/update_triplet_weights.hpp | 16 ++++++++-------- 31 files changed, 148 insertions(+), 157 deletions(-) diff --git a/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp b/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp index aadfcaf633..0354690207 100644 --- a/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp +++ b/device/common/include/traccc/clusterization/device/impl/ccl_kernel.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp index 8072fd3339..4a71882a68 100644 --- a/device/common/include/traccc/finding/device/find_tracks.hpp +++ b/device/common/include/traccc/finding/device/find_tracks.hpp @@ -7,11 +7,13 @@ #pragma once +// Local include(s). +#include "traccc/device/concepts/barrier.hpp" +#include "traccc/device/concepts/thread_id.hpp" + // Project include(s). #include "traccc/definitions/primitives.hpp" #include "traccc/definitions/qualifiers.hpp" -#include "traccc/device/concepts/barrier.hpp" -#include "traccc/device/concepts/thread_id.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/candidate_link.hpp" @@ -140,8 +142,8 @@ struct find_tracks_shared_payload { template TRACCC_DEVICE inline void find_tracks( - thread_id_t& thread_id, barrier_t& barrier, const finding_config& cfg, - const find_tracks_payload& payload, + const thread_id_t& thread_id, const barrier_t& barrier, + const finding_config& cfg, const find_tracks_payload& payload, const find_tracks_shared_payload& shared_payload); } // namespace traccc::device diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp index 5765b2329a..8a99da65b7 100644 --- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp @@ -19,8 +19,8 @@ namespace traccc::device { template TRACCC_DEVICE inline void find_tracks( - thread_id_t& thread_id, barrier_t& barrier, const finding_config& cfg, - const find_tracks_payload& payload, + const thread_id_t& thread_id, const barrier_t& barrier, + const finding_config& cfg, const find_tracks_payload& payload, const find_tracks_shared_payload& shared_payload) { /* diff --git a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp index fba99d7ab5..14741f8881 100644 --- a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp +++ b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp @@ -11,7 +11,8 @@ #include "traccc/utils/particle.hpp" // Detray include(s). -#include "detray/utils/tuple.hpp" +#include "detray/propagator/constrained_step.hpp" +#include "detray/utils/tuple_helpers.hpp" namespace traccc::device { diff --git a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp index 9fe6560490..825f037746 100644 --- a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp +++ b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp @@ -14,6 +14,7 @@ #include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/candidate_link.hpp" +#include "traccc/finding/finding_config.hpp" // VecMem include(s). #include diff --git a/device/common/include/traccc/fitting/device/fill_sort_keys.hpp b/device/common/include/traccc/fitting/device/fill_sort_keys.hpp index 837e4231e1..c9caf7a877 100644 --- a/device/common/include/traccc/fitting/device/fill_sort_keys.hpp +++ b/device/common/include/traccc/fitting/device/fill_sort_keys.hpp @@ -1,18 +1,18 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// Project include(s). +// Local include(s). +#include "traccc/device/global_index.hpp" #include "traccc/edm/device/sort_key.hpp" -#include "traccc/edm/track_candidate.hpp" -// System include(s). -#include +// Project include(s). +#include "traccc/edm/track_candidate.hpp" namespace traccc::device { @@ -24,7 +24,7 @@ namespace traccc::device { /// @param[out] ids_view The param ids /// TRACCC_HOST_DEVICE inline void fill_sort_keys( - std::size_t globalIndex, + global_index_t globalIndex, const track_candidate_container_types::const_view& track_candidates_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view); diff --git a/device/common/include/traccc/fitting/device/fit.hpp b/device/common/include/traccc/fitting/device/fit.hpp index cd018d11e3..066ce11f9e 100644 --- a/device/common/include/traccc/fitting/device/fit.hpp +++ b/device/common/include/traccc/fitting/device/fit.hpp @@ -1,20 +1,20 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Local include(s). +#include "traccc/device/global_index.hpp" + // Project include(s). #include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/track_candidate.hpp" #include "traccc/edm/track_state.hpp" -// System include(s). -#include - namespace traccc::device { /// Function used for fitting a track for a given track candidates @@ -26,7 +26,7 @@ namespace traccc::device { /// template TRACCC_HOST_DEVICE inline void fit( - std::size_t globalIndex, + global_index_t globalIndex, typename fitter_t::detector_type::view_type det_data, const typename fitter_t::bfield_type field_data, const typename fitter_t::config_type cfg, diff --git a/device/common/include/traccc/fitting/device/impl/fill_sort_keys.ipp b/device/common/include/traccc/fitting/device/impl/fill_sort_keys.ipp index fbb8c68f44..610fd23cb4 100644 --- a/device/common/include/traccc/fitting/device/impl/fill_sort_keys.ipp +++ b/device/common/include/traccc/fitting/device/impl/fill_sort_keys.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,12 +10,12 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void fill_sort_keys( - std::size_t globalIndex, + const global_index_t globalIndex, const track_candidate_container_types::const_view& track_candidates_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view) { - track_candidate_container_types::const_device track_candidates( + const track_candidate_container_types::const_device track_candidates( track_candidates_view); // Keys @@ -29,12 +29,9 @@ TRACCC_HOST_DEVICE inline void fill_sort_keys( } // Key = The number of measurements - keys_device.at(static_cast(globalIndex)) = - static_cast( - track_candidates.at(static_cast(globalIndex)) - .items.size()); - ids_device.at(static_cast(globalIndex)) = - static_cast(globalIndex); + keys_device.at(globalIndex) = static_cast( + track_candidates.at(globalIndex).items.size()); + ids_device.at(globalIndex) = globalIndex; } } // namespace traccc::device diff --git a/device/common/include/traccc/fitting/device/impl/fit.ipp b/device/common/include/traccc/fitting/device/impl/fit.ipp index d3c7e1832b..55611d74c3 100644 --- a/device/common/include/traccc/fitting/device/impl/fit.ipp +++ b/device/common/include/traccc/fitting/device/impl/fit.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -11,7 +11,7 @@ namespace traccc::device { template TRACCC_HOST_DEVICE inline void fit( - std::size_t globalIndex, + const global_index_t globalIndex, typename fitter_t::detector_type::view_type det_data, const typename fitter_t::bfield_type field_data, const typename fitter_t::config_type cfg, @@ -34,8 +34,7 @@ TRACCC_HOST_DEVICE inline void fit( return; } - const unsigned int param_id = - param_ids.at(static_cast(globalIndex)); + const unsigned int param_id = param_ids.at(globalIndex); // Track candidates per track const auto& track_candidates_per_track = diff --git a/device/common/include/traccc/seeding/device/count_doublets.hpp b/device/common/include/traccc/seeding/device/count_doublets.hpp index 9854aa8bbe..b891041de8 100644 --- a/device/common/include/traccc/seeding/device/count_doublets.hpp +++ b/device/common/include/traccc/seeding/device/count_doublets.hpp @@ -1,22 +1,22 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" +// Local include(s). #include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" #include "traccc/edm/device/doublet_counter.hpp" + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" -// System include(s). -#include - namespace traccc::device { /// Function used for calculating the number of spacepoint doublets @@ -35,7 +35,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void count_doublets( - std::size_t globalIndex, const seedfinder_config& config, + global_index_t globalIndex, const seedfinder_config& config, const sp_grid_const_view& sp_view, const vecmem::data::vector_view& sp_ps_view, doublet_counter_collection_types::view doublet_view, unsigned int& nMidBot, diff --git a/device/common/include/traccc/seeding/device/count_grid_capacities.hpp b/device/common/include/traccc/seeding/device/count_grid_capacities.hpp index 9edcb9bcde..72cbbdaa27 100644 --- a/device/common/include/traccc/seeding/device/count_grid_capacities.hpp +++ b/device/common/include/traccc/seeding/device/count_grid_capacities.hpp @@ -1,15 +1,18 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2022 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Local include(s). +#include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" + // Project include(s). #include "traccc/definitions/qualifiers.hpp" -#include "traccc/device/fill_prefix_sum.hpp" #include "traccc/edm/spacepoint.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" @@ -17,9 +20,6 @@ // VecMem include(s). #include -// System include(s). -#include - namespace traccc::device { /// Function used for calculating the capacity for the spacepoint grid @@ -40,7 +40,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void count_grid_capacities( - const std::size_t globalIndex, const seedfinder_config& config, + global_index_t globalIndex, const seedfinder_config& config, const sp_grid::axis_p0_type& phi_axis, const sp_grid::axis_p1_type& z_axis, const spacepoint_collection_types::const_view& spacepoints, vecmem::data::vector_view grid_capacities); diff --git a/device/common/include/traccc/seeding/device/count_triplets.hpp b/device/common/include/traccc/seeding/device/count_triplets.hpp index 5675c722d1..cc4c376769 100644 --- a/device/common/include/traccc/seeding/device/count_triplets.hpp +++ b/device/common/include/traccc/seeding/device/count_triplets.hpp @@ -1,22 +1,23 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" +// Local include(s). #include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" #include "traccc/edm/device/device_doublet.hpp" #include "traccc/edm/device/doublet_counter.hpp" #include "traccc/edm/device/triplet_counter.hpp" + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" -// System include(s). -#include namespace traccc::device { @@ -38,7 +39,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void count_triplets( - std::size_t globalIndex, const seedfinder_config& config, + global_index_t globalIndex, const seedfinder_config& config, const sp_grid_const_view& sp_view, const doublet_counter_collection_types::const_view& dc_view, const device_doublet_collection_types::const_view& mid_bot_doublet_view, diff --git a/device/common/include/traccc/seeding/device/estimate_track_params.hpp b/device/common/include/traccc/seeding/device/estimate_track_params.hpp index 77960bfa79..99c87be010 100644 --- a/device/common/include/traccc/seeding/device/estimate_track_params.hpp +++ b/device/common/include/traccc/seeding/device/estimate_track_params.hpp @@ -1,12 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Local include(s). +#include "traccc/device/global_index.hpp" + // Project include(s). #include "traccc/edm/seed.hpp" #include "traccc/edm/spacepoint.hpp" @@ -25,7 +28,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void estimate_track_params( - const std::size_t globalIndex, + global_index_t globalIndex, const spacepoint_collection_types::const_view& spacepoints_view, const seed_collection_types::const_view& seeds_view, const vector3& bfield, const std::array& stddev, diff --git a/device/common/include/traccc/seeding/device/find_doublets.hpp b/device/common/include/traccc/seeding/device/find_doublets.hpp index d44c1fd651..1f4a5ab639 100644 --- a/device/common/include/traccc/seeding/device/find_doublets.hpp +++ b/device/common/include/traccc/seeding/device/find_doublets.hpp @@ -1,23 +1,23 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" +// Local include(s). #include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" #include "traccc/edm/device/device_doublet.hpp" #include "traccc/edm/device/doublet_counter.hpp" + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" -// System include(s). -#include - namespace traccc::device { /// Function finding all of the spacepoint doublets @@ -34,7 +34,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void find_doublets( - std::size_t globalIndex, const seedfinder_config& config, + global_index_t globalIndex, const seedfinder_config& config, const sp_grid_const_view& sp_view, const doublet_counter_collection_types::const_view& dc_view, device_doublet_collection_types::view mb_doublets_view, diff --git a/device/common/include/traccc/seeding/device/find_triplets.hpp b/device/common/include/traccc/seeding/device/find_triplets.hpp index fc6672832e..6f6576b0e7 100644 --- a/device/common/include/traccc/seeding/device/find_triplets.hpp +++ b/device/common/include/traccc/seeding/device/find_triplets.hpp @@ -1,24 +1,24 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" +// Local include(s). #include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.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/definitions/qualifiers.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" -// System include(s). -#include - namespace traccc::device { /// Function finding all of the spacepoint triplets @@ -39,7 +39,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void find_triplets( - std::size_t globalIndex, const seedfinder_config& config, + global_index_t globalIndex, const seedfinder_config& config, const seedfilter_config& filter_config, const sp_grid_const_view& sp_view, const doublet_counter_collection_types::const_view& dc_view, const device_doublet_collection_types::const_view& mid_top_doublet_view, diff --git a/device/common/include/traccc/seeding/device/form_spacepoints.hpp b/device/common/include/traccc/seeding/device/form_spacepoints.hpp index b6ddf8f2d9..2102635053 100644 --- a/device/common/include/traccc/seeding/device/form_spacepoints.hpp +++ b/device/common/include/traccc/seeding/device/form_spacepoints.hpp @@ -1,21 +1,21 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Local include(s). +#include "traccc/device/global_index.hpp" + // Project include(s). #include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint.hpp" #include "traccc/geometry/silicon_detector_description.hpp" -// System include(s). -#include - namespace traccc::device { /// Function for creating 3D spacepoints out of 2D measurements @@ -28,7 +28,7 @@ namespace traccc::device { /// template TRACCC_HOST_DEVICE inline void form_spacepoints( - std::size_t globalIndex, typename detector_t::view_type det_view, + global_index_t globalIndex, typename detector_t::view_type det_view, const measurement_collection_types::const_view& measurements_view, unsigned int measurement_count, spacepoint_collection_types::view spacepoints_view); diff --git a/device/common/include/traccc/seeding/device/impl/count_doublets.ipp b/device/common/include/traccc/seeding/device/impl/count_doublets.ipp index cac4991c14..406f2bc260 100644 --- a/device/common/include/traccc/seeding/device/impl/count_doublets.ipp +++ b/device/common/include/traccc/seeding/device/impl/count_doublets.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -20,7 +20,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void count_doublets( - const std::size_t globalIndex, const seedfinder_config& config, + const global_index_t globalIndex, const seedfinder_config& config, const sp_grid_const_view& sp_view, const vecmem::data::vector_view& sp_ps_view, doublet_counter_collection_types::view doublet_view, unsigned int& nMidBot, @@ -33,8 +33,7 @@ inline void count_doublets( } // Get the middle spacepoint that we need to be looking at. - const prefix_sum_element_t middle_sp_idx = - sp_prefix_sum[static_cast(globalIndex)]; + const prefix_sum_element_t middle_sp_idx = sp_prefix_sum.at(globalIndex); // Set up the device containers. const const_sp_grid_device sp_grid(sp_view); diff --git a/device/common/include/traccc/seeding/device/impl/count_grid_capacities.ipp b/device/common/include/traccc/seeding/device/impl/count_grid_capacities.ipp index ae9d684329..9616a0ee54 100644 --- a/device/common/include/traccc/seeding/device/impl/count_grid_capacities.ipp +++ b/device/common/include/traccc/seeding/device/impl/count_grid_capacities.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2022 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -17,7 +17,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void count_grid_capacities( - const std::size_t globalIndex, const seedfinder_config& config, + const global_index_t globalIndex, const seedfinder_config& config, const sp_grid::axis_p0_type& phi_axis, const sp_grid::axis_p1_type& z_axis, const spacepoint_collection_types::const_view& spacepoints_view, vecmem::data::vector_view grid_capacities_view) { @@ -28,8 +28,7 @@ inline void count_grid_capacities( if (globalIndex >= spacepoints.size()) { return; } - const spacepoint sp = - spacepoints.at(static_cast(globalIndex)); + const spacepoint sp = spacepoints.at(globalIndex); /// Check out if the spacepoint can be used for seeding. if (is_valid_sp(config, sp) != detray::detail::invalid_value()) { diff --git a/device/common/include/traccc/seeding/device/impl/count_triplets.ipp b/device/common/include/traccc/seeding/device/impl/count_triplets.ipp index 17f9c2c914..7738085cc1 100644 --- a/device/common/include/traccc/seeding/device/impl/count_triplets.ipp +++ b/device/common/include/traccc/seeding/device/impl/count_triplets.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -17,7 +17,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void count_triplets( - const std::size_t globalIndex, const seedfinder_config& config, + const global_index_t globalIndex, const seedfinder_config& config, const sp_grid_const_view& sp_view, const doublet_counter_collection_types::const_view& dc_view, const device_doublet_collection_types::const_view& mid_bot_doublet_view, @@ -34,8 +34,7 @@ inline void count_triplets( } // Get current mid bottom doublet - const device_doublet mid_bot = - mid_bot_doublet_device.at(static_cast(globalIndex)); + const device_doublet mid_bot = mid_bot_doublet_device.at(globalIndex); // Create device copy of input parameters const device_doublet_collection_types::const_device mid_top_doublet_device( diff --git a/device/common/include/traccc/seeding/device/impl/estimate_track_params.ipp b/device/common/include/traccc/seeding/device/impl/estimate_track_params.ipp index 28538a0611..1c8d15337e 100644 --- a/device/common/include/traccc/seeding/device/impl/estimate_track_params.ipp +++ b/device/common/include/traccc/seeding/device/impl/estimate_track_params.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2022 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -15,7 +15,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void estimate_track_params( - const std::size_t globalIndex, + const global_index_t globalIndex, const spacepoint_collection_types::const_view& spacepoints_view, const seed_collection_types::const_view& seeds_view, const vector3& bfield, const std::array& stddev, @@ -32,8 +32,7 @@ inline void estimate_track_params( bound_track_parameters_collection_types::device params_device(params_view); - const seed& this_seed = - seeds_device.at(static_cast(globalIndex)); + const seed& this_seed = seeds_device.at(globalIndex); // Get bound track parameter bound_track_parameters track_params; @@ -52,7 +51,7 @@ inline void estimate_track_params( track_params.set_surface_link(spB.meas.surface_link); // Save the object into global memory. - params_device[static_cast(globalIndex)] = track_params; + params_device.at(globalIndex) = track_params; } } // namespace traccc::device diff --git a/device/common/include/traccc/seeding/device/impl/find_doublets.ipp b/device/common/include/traccc/seeding/device/impl/find_doublets.ipp index 4c3ca5bc2d..ce8e348c02 100644 --- a/device/common/include/traccc/seeding/device/impl/find_doublets.ipp +++ b/device/common/include/traccc/seeding/device/impl/find_doublets.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -20,7 +20,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void find_doublets( - const std::size_t globalIndex, const seedfinder_config& config, + const global_index_t globalIndex, const seedfinder_config& config, const sp_grid_const_view& sp_view, const doublet_counter_collection_types::const_view& dc_view, device_doublet_collection_types::view mb_doublets_view, @@ -34,8 +34,7 @@ inline void find_doublets( } // Get the middle spacepoint that we need to be looking at. - const doublet_counter middle_sp_counter = - doublet_counts.at(static_cast(globalIndex)); + const doublet_counter middle_sp_counter = doublet_counts.at(globalIndex); // Set up the device containers. const const_sp_grid_device sp_grid(sp_view); @@ -113,9 +112,8 @@ inline void find_doublets( // Add it as a candidate to the middle-bottom container. const unsigned int pos = mid_bot_start_idx + mid_bot_idx++; assert(pos < mb_doublets.size()); - mb_doublets.at(pos) = { - {other_bin_idx, other_sp_idx}, - static_cast(globalIndex)}; + mb_doublets.at(pos) = {{other_bin_idx, other_sp_idx}, + globalIndex}; } // Check if this spacepoint is a compatible "top" spacepoint to // the thread's "middle" spacepoint. @@ -126,9 +124,8 @@ inline void find_doublets( // Add it as a candidate to the middle-top container. const unsigned int pos = mid_top_start_idx + mid_top_idx++; assert(pos < mt_doublets.size()); - mt_doublets.at(pos) = { - {other_bin_idx, other_sp_idx}, - static_cast(globalIndex)}; + mt_doublets.at(pos) = {{other_bin_idx, other_sp_idx}, + globalIndex}; } } } diff --git a/device/common/include/traccc/seeding/device/impl/find_triplets.ipp b/device/common/include/traccc/seeding/device/impl/find_triplets.ipp index 542dca7a2a..94f9c9cbca 100644 --- a/device/common/include/traccc/seeding/device/impl/find_triplets.ipp +++ b/device/common/include/traccc/seeding/device/impl/find_triplets.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -17,7 +17,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void find_triplets( - const std::size_t globalIndex, const seedfinder_config& config, + const global_index_t globalIndex, const seedfinder_config& config, const seedfilter_config& filter_config, const sp_grid_const_view& sp_view, const doublet_counter_collection_types::const_view& dc_view, const device_doublet_collection_types::const_view& mid_top_doublet_view, @@ -42,8 +42,7 @@ inline void find_triplets( spM_tc_view); // Get the current work item information - const triplet_counter mid_bot_counter = - triplet_counts.at(static_cast(globalIndex)); + const triplet_counter mid_bot_counter = triplet_counts.at(globalIndex); const triplet_counter_spM spM_counter = triplet_counts_spM.at(mid_bot_counter.spM_counter_link); const doublet_counter doublet_count = @@ -107,7 +106,7 @@ inline void find_triplets( // Add triplet to jagged vector triplets.at(posTriplets++) = device_triplet( - {spT_loc, static_cast(globalIndex), curvature, + {spT_loc, globalIndex, curvature, -impact_parameter * filter_config.impactWeightFactor, lb.Zo()}); } diff --git a/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp b/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp index 5e96cfcdd5..3f9307834d 100644 --- a/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp +++ b/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -17,7 +17,7 @@ namespace traccc::device { template TRACCC_HOST_DEVICE inline void form_spacepoints( - std::size_t globalIndex, typename detector_t::view_type det_view, + const global_index_t globalIndex, typename detector_t::view_type det_view, const measurement_collection_types::const_view& measurements_view, unsigned int measurement_count, spacepoint_collection_types::view spacepoints_view) { @@ -36,7 +36,7 @@ TRACCC_HOST_DEVICE inline void form_spacepoints( assert(measurements.size() == measurement_count); spacepoint_collection_types::device spacepoints(spacepoints_view); - const auto& meas = measurements.at(static_cast(globalIndex)); + const auto& meas = measurements.at(globalIndex); // Fill the spacepoint using the common function. if (details::is_valid_measurement(meas)) { diff --git a/device/common/include/traccc/seeding/device/impl/populate_grid.ipp b/device/common/include/traccc/seeding/device/impl/populate_grid.ipp index d1ecd76abe..01da7a2b2f 100644 --- a/device/common/include/traccc/seeding/device/impl/populate_grid.ipp +++ b/device/common/include/traccc/seeding/device/impl/populate_grid.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2022 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -14,7 +14,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void populate_grid( - unsigned int globalIndex, const seedfinder_config& config, + const global_index_t globalIndex, const seedfinder_config& config, const spacepoint_collection_types::const_view& spacepoints_view, sp_grid_view grid_view) { diff --git a/device/common/include/traccc/seeding/device/impl/reduce_triplet_counts.ipp b/device/common/include/traccc/seeding/device/impl/reduce_triplet_counts.ipp index 7d911935f1..7f2f060201 100644 --- a/device/common/include/traccc/seeding/device/impl/reduce_triplet_counts.ipp +++ b/device/common/include/traccc/seeding/device/impl/reduce_triplet_counts.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -14,7 +14,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void reduce_triplet_counts( - const std::size_t globalIndex, + const global_index_t globalIndex, const doublet_counter_collection_types::const_view& dc_view, triplet_counter_spM_collection_types::view spM_tc_view, unsigned int& num_triplets) { @@ -33,8 +33,7 @@ inline void reduce_triplet_counts( assert(doublet_counts.size() == spM_counts.size()); // Get triplet counter for this middle spacepoint - triplet_counter_spM& this_spM_counter = - spM_counts.at(static_cast(globalIndex)); + triplet_counter_spM& this_spM_counter = spM_counts.at(globalIndex); // Check if anything needs to be done. if (this_spM_counter.m_nTriplets == 0) { @@ -42,8 +41,7 @@ inline void reduce_triplet_counts( } // Fill the middle spacepoint information of the spM triplet counter - this_spM_counter.spM = - doublet_counts.at(static_cast(globalIndex)).m_spM; + this_spM_counter.spM = doublet_counts.at(globalIndex).m_spM; // Increment total number of triplets and claim position for this middle // spacepoint's triplets diff --git a/device/common/include/traccc/seeding/device/impl/select_seeds.ipp b/device/common/include/traccc/seeding/device/impl/select_seeds.ipp index 0bd5ce2ea0..09d2b5fa47 100644 --- a/device/common/include/traccc/seeding/device/impl/select_seeds.ipp +++ b/device/common/include/traccc/seeding/device/impl/select_seeds.ipp @@ -1,18 +1,19 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// System include(s) -#include - // Project include(s). #include "traccc/seeding/seed_selecting_helper.hpp" +// System include(s) +#include +#include + namespace traccc::device { namespace details { @@ -58,7 +59,7 @@ TRACCC_HOST_DEVICE void insertionSort(triplet* arr, // Select seeds kernel TRACCC_HOST_DEVICE inline void select_seeds( - const std::size_t globalIndex, const seedfilter_config& filter_config, + const global_index_t globalIndex, const seedfilter_config& filter_config, const spacepoint_collection_types::const_view& spacepoints_view, const sp_grid_const_view& internal_sp_view, const triplet_counter_spM_collection_types::const_view& spM_tc_view, @@ -84,8 +85,7 @@ inline void select_seeds( seed_collection_types::device seeds_device(seed_view); // Current work item = middle spacepoint - const triplet_counter_spM spM_counter = - triplet_counts_spM.at(static_cast(globalIndex)); + const triplet_counter_spM spM_counter = triplet_counts_spM.at(globalIndex); const sp_location spM_loc = spM_counter.spM; const internal_spacepoint spM = internal_sp_device.bin(spM_loc.bin_idx)[spM_loc.sp_idx]; diff --git a/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp b/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp index 56be55dfb5..6a381a70c7 100644 --- a/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp +++ b/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -17,7 +17,7 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void update_triplet_weights( - const std::size_t globalIndex, const seedfilter_config& filter_config, + const global_index_t globalIndex, const seedfilter_config& filter_config, const sp_grid_const_view& sp_view, const triplet_counter_spM_collection_types::const_view& spM_tc_view, const triplet_counter_collection_types::const_view& tc_view, scalar* data, @@ -37,8 +37,7 @@ inline void update_triplet_weights( tc_view); // Current work item - device_triplet this_triplet = - triplets[static_cast(globalIndex)]; + device_triplet this_triplet = triplets.at(globalIndex); const sp_location& spT_idx = this_triplet.spT; @@ -128,8 +127,7 @@ inline void update_triplet_weights( } } - triplets[static_cast(globalIndex)].weight = - this_triplet.weight; + triplets.at(globalIndex).weight = this_triplet.weight; } } // namespace traccc::device diff --git a/device/common/include/traccc/seeding/device/populate_grid.hpp b/device/common/include/traccc/seeding/device/populate_grid.hpp index 6d2fe5f488..2370e6f617 100644 --- a/device/common/include/traccc/seeding/device/populate_grid.hpp +++ b/device/common/include/traccc/seeding/device/populate_grid.hpp @@ -1,25 +1,22 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2022 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Local include(s). +#include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" + // Project include(s). #include "traccc/definitions/qualifiers.hpp" -#include "traccc/device/fill_prefix_sum.hpp" #include "traccc/edm/spacepoint.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" -// VecMem include(s). -#include - -// System include(s). -#include - namespace traccc::device { /// Function populating the spacepoint grid @@ -31,7 +28,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void populate_grid( - unsigned int globalIndex, const seedfinder_config& config, + global_index_t globalIndex, const seedfinder_config& config, const spacepoint_collection_types::const_view& spacepoints, sp_grid_view grid); diff --git a/device/common/include/traccc/seeding/device/reduce_triplet_counts.hpp b/device/common/include/traccc/seeding/device/reduce_triplet_counts.hpp index a2b46f3248..6d862a87be 100644 --- a/device/common/include/traccc/seeding/device/reduce_triplet_counts.hpp +++ b/device/common/include/traccc/seeding/device/reduce_triplet_counts.hpp @@ -1,15 +1,18 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Local include(s). +#include "traccc/device/global_index.hpp" +#include "traccc/edm/device/triplet_counter.hpp" + // Project include(s). #include "traccc/definitions/qualifiers.hpp" -#include "traccc/edm/device/triplet_counter.hpp" namespace traccc::device { @@ -27,7 +30,7 @@ namespace traccc::device { /// @param[out] num_triplets The total number of triplets TRACCC_HOST_DEVICE inline void reduce_triplet_counts( - std::size_t globalIndex, + global_index_t globalIndex, const doublet_counter_collection_types::const_view& dc_view, triplet_counter_spM_collection_types::view spM_tc_view, unsigned int& num_triplets); diff --git a/device/common/include/traccc/seeding/device/select_seeds.hpp b/device/common/include/traccc/seeding/device/select_seeds.hpp index 9db06c9890..6d9d35efc5 100644 --- a/device/common/include/traccc/seeding/device/select_seeds.hpp +++ b/device/common/include/traccc/seeding/device/select_seeds.hpp @@ -1,24 +1,23 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" +// Local include(s). #include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" #include "traccc/edm/device/device_triplet.hpp" #include "traccc/edm/device/triplet_counter.hpp" + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/seed.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" #include "traccc/seeding/detail/triplet.hpp" -// System include(s). -#include -#include - namespace traccc::device { /// Function used for selecting good triplets to be recorded into seed @@ -35,7 +34,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void select_seeds( - std::size_t globalIndex, const seedfilter_config& filter_config, + global_index_t globalIndex, const seedfilter_config& filter_config, const spacepoint_collection_types::const_view& spacepoints_view, const sp_grid_const_view& internal_sp_view, const triplet_counter_spM_collection_types::const_view& spM_tc_view, @@ -46,4 +45,4 @@ inline void select_seeds( } // namespace traccc::device // Include the implementation. -#include "traccc/seeding/device/impl/select_seeds.ipp" \ No newline at end of file +#include "traccc/seeding/device/impl/select_seeds.ipp" diff --git a/device/common/include/traccc/seeding/device/update_triplet_weights.hpp b/device/common/include/traccc/seeding/device/update_triplet_weights.hpp index fceac38bff..6278875c86 100644 --- a/device/common/include/traccc/seeding/device/update_triplet_weights.hpp +++ b/device/common/include/traccc/seeding/device/update_triplet_weights.hpp @@ -1,23 +1,23 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2023 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" +// Local include(s). #include "traccc/device/fill_prefix_sum.hpp" +#include "traccc/device/global_index.hpp" #include "traccc/edm/device/device_triplet.hpp" #include "traccc/edm/device/triplet_counter.hpp" + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/seeding/detail/seeding_config.hpp" #include "traccc/seeding/detail/spacepoint_grid.hpp" -// System include(s) -#include - namespace traccc::device { /// Function used for updating the triplets' weights @@ -33,7 +33,7 @@ namespace traccc::device { /// TRACCC_HOST_DEVICE inline void update_triplet_weights( - std::size_t globalIndex, const seedfilter_config& filter_config, + global_index_t globalIndex, const seedfilter_config& filter_config, const sp_grid_const_view& sp_view, const triplet_counter_spM_collection_types::const_view& spM_tc_view, const triplet_counter_collection_types::const_view& tc_view, scalar* data, @@ -42,4 +42,4 @@ inline void update_triplet_weights( } // namespace traccc::device // Include the implementation. -#include "traccc/seeding/device/impl/update_triplet_weights.ipp" \ No newline at end of file +#include "traccc/seeding/device/impl/update_triplet_weights.ipp" From 2aa77314ccae20b96ca863db9a057ea726f6ca85 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 8 Jan 2025 23:25:25 +0100 Subject: [PATCH 3/5] Call device functions consistently from CUDA. --- .../src/clusterization/kernels/ccl_kernel.cu | 2 +- .../src/finding/kernels/apply_interaction.cuh | 7 ++-- .../cuda/src/finding/kernels/build_tracks.cu | 10 +++--- .../cuda/src/finding/kernels/build_tracks.cuh | 7 ++-- .../src/finding/kernels/fill_sort_keys.cu | 10 ++++-- .../src/finding/kernels/fill_sort_keys.cuh | 5 +-- .../finding/kernels/make_barcode_sequence.cu | 12 ++++--- .../finding/kernels/make_barcode_sequence.cuh | 5 +-- .../kernels/propagate_to_next_surface.cuh | 6 ++-- .../cuda/src/finding/kernels/prune_tracks.cu | 11 +++--- .../cuda/src/finding/kernels/prune_tracks.cuh | 5 +-- .../apply_interaction_default_detector.cu | 6 +++- .../specializations/apply_interaction_src.cuh | 15 ++++---- ...pagate_to_next_surface_default_detector.cu | 5 +-- .../propagate_to_next_surface_src.cuh | 9 +++-- device/cuda/src/fitting/fitting_algorithm.cu | 14 ++++---- device/cuda/src/sanity/contiguous_on.cuh | 9 ++--- device/cuda/src/sanity/ordered_on.cuh | 6 ++-- device/cuda/src/seeding/seed_finding.cu | 34 +++++++++---------- device/cuda/src/seeding/spacepoint_binning.cu | 11 +++--- .../seeding/spacepoint_formation_algorithm.cu | 9 ++--- .../src/seeding/track_params_estimation.cu | 8 ++--- device/cuda/src/utils/global_index.hpp | 21 ++++++++++++ device/cuda/src/utils/thread_id.hpp | 19 +++++------ 24 files changed, 145 insertions(+), 101 deletions(-) create mode 100644 device/cuda/src/utils/global_index.hpp diff --git a/device/cuda/src/clusterization/kernels/ccl_kernel.cu b/device/cuda/src/clusterization/kernels/ccl_kernel.cu index ac1313f45b..78054fc119 100644 --- a/device/cuda/src/clusterization/kernels/ccl_kernel.cu +++ b/device/cuda/src/clusterization/kernels/ccl_kernel.cu @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ diff --git a/device/cuda/src/finding/kernels/apply_interaction.cuh b/device/cuda/src/finding/kernels/apply_interaction.cuh index 04ceefd305..01ce47d998 100644 --- a/device/cuda/src/finding/kernels/apply_interaction.cuh +++ b/device/cuda/src/finding/kernels/apply_interaction.cuh @@ -1,15 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "traccc/edm/track_parameters.hpp" +// Project include(s). #include "traccc/finding/device/apply_interaction.hpp" -#include "traccc/geometry/detector.hpp" +#include "traccc/finding/finding_config.hpp" namespace traccc::cuda::kernels { @@ -17,4 +17,5 @@ template __global__ void apply_interaction( const finding_config cfg, device::apply_interaction_payload payload); + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/build_tracks.cu b/device/cuda/src/finding/kernels/build_tracks.cu index 801bf118d6..ded6c57a0b 100644 --- a/device/cuda/src/finding/kernels/build_tracks.cu +++ b/device/cuda/src/finding/kernels/build_tracks.cu @@ -1,11 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ +// Local include(s). +#include "../../utils/global_index.hpp" #include "build_tracks.cuh" + +// Project include(s). #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate.hpp" #include "traccc/edm/track_parameters.hpp" @@ -18,8 +22,6 @@ namespace traccc::cuda::kernels { __global__ void build_tracks(const finding_config cfg, device::build_tracks_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; - - device::build_tracks(gid, cfg, payload); + device::build_tracks(details::global_index1(), cfg, payload); } } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/build_tracks.cuh b/device/cuda/src/finding/kernels/build_tracks.cuh index 0cf59cd628..8fc338fe7e 100644 --- a/device/cuda/src/finding/kernels/build_tracks.cuh +++ b/device/cuda/src/finding/kernels/build_tracks.cuh @@ -1,16 +1,13 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "traccc/edm/measurement.hpp" -#include "traccc/edm/track_candidate.hpp" -#include "traccc/edm/track_parameters.hpp" -#include "traccc/finding/candidate_link.hpp" +// Project include(s). #include "traccc/finding/device/build_tracks.hpp" #include "traccc/finding/finding_config.hpp" diff --git a/device/cuda/src/finding/kernels/fill_sort_keys.cu b/device/cuda/src/finding/kernels/fill_sort_keys.cu index 4115841c0a..f8f5b21e56 100644 --- a/device/cuda/src/finding/kernels/fill_sort_keys.cu +++ b/device/cuda/src/finding/kernels/fill_sort_keys.cu @@ -1,18 +1,22 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ +// Local include(s). +#include "../../utils/global_index.hpp" #include "fill_sort_keys.cuh" -#include "traccc/edm/track_parameters.hpp" + +// Project include(s). #include "traccc/finding/device/fill_sort_keys.hpp" namespace traccc::cuda::kernels { __global__ void fill_sort_keys(device::fill_sort_keys_payload payload) { - device::fill_sort_keys(threadIdx.x + blockIdx.x * blockDim.x, payload); + device::fill_sort_keys(details::global_index1(), payload); } + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/fill_sort_keys.cuh b/device/cuda/src/finding/kernels/fill_sort_keys.cuh index 5f9aedb22c..94797fe820 100644 --- a/device/cuda/src/finding/kernels/fill_sort_keys.cuh +++ b/device/cuda/src/finding/kernels/fill_sort_keys.cuh @@ -1,16 +1,17 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "traccc/edm/track_parameters.hpp" +// Project include(s). #include "traccc/finding/device/fill_sort_keys.hpp" namespace traccc::cuda::kernels { __global__ void fill_sort_keys(device::fill_sort_keys_payload payload); + } diff --git a/device/cuda/src/finding/kernels/make_barcode_sequence.cu b/device/cuda/src/finding/kernels/make_barcode_sequence.cu index e6587b553d..a0d0f72e6d 100644 --- a/device/cuda/src/finding/kernels/make_barcode_sequence.cu +++ b/device/cuda/src/finding/kernels/make_barcode_sequence.cu @@ -1,12 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ +// Local include(s). +#include "../../utils/global_index.hpp" #include "make_barcode_sequence.cuh" -#include "traccc/edm/measurement.hpp" + +// Project include(s). #include "traccc/finding/device/make_barcode_sequence.hpp" namespace traccc::cuda::kernels { @@ -14,8 +17,7 @@ namespace traccc::cuda::kernels { __global__ void make_barcode_sequence( device::make_barcode_sequence_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; - - device::make_barcode_sequence(gid, payload); + device::make_barcode_sequence(details::global_index1(), payload); } + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/make_barcode_sequence.cuh b/device/cuda/src/finding/kernels/make_barcode_sequence.cuh index 13f147a047..a8d1a40468 100644 --- a/device/cuda/src/finding/kernels/make_barcode_sequence.cuh +++ b/device/cuda/src/finding/kernels/make_barcode_sequence.cuh @@ -1,17 +1,18 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "traccc/edm/measurement.hpp" +// Project include(s). #include "traccc/finding/device/make_barcode_sequence.hpp" namespace traccc::cuda::kernels { __global__ void make_barcode_sequence( device::make_barcode_sequence_payload payload); + } diff --git a/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh b/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh index c5df625145..49ab46c16d 100644 --- a/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh +++ b/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh @@ -1,16 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "./specializations/types.hpp" +// Project include(s). #include "traccc/finding/device/propagate_to_next_surface.hpp" #include "traccc/finding/finding_config.hpp" -#include "traccc/geometry/detector.hpp" namespace traccc::cuda::kernels { @@ -18,4 +17,5 @@ template __global__ void propagate_to_next_surface( const finding_config cfg, device::propagate_to_next_surface_payload payload); + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/prune_tracks.cu b/device/cuda/src/finding/kernels/prune_tracks.cu index f431676a4d..a7f339707b 100644 --- a/device/cuda/src/finding/kernels/prune_tracks.cu +++ b/device/cuda/src/finding/kernels/prune_tracks.cu @@ -1,19 +1,22 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ +// Local include(s). +#include "../../utils/global_index.hpp" #include "prune_tracks.cuh" + +// Project include(s). #include "traccc/finding/device/prune_tracks.hpp" namespace traccc::cuda::kernels { __global__ void prune_tracks(device::prune_tracks_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; - - device::prune_tracks(gid, payload); + device::prune_tracks(details::global_index1(), payload); } + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/prune_tracks.cuh b/device/cuda/src/finding/kernels/prune_tracks.cuh index 2aaa23b6fa..4dcad7c9e9 100644 --- a/device/cuda/src/finding/kernels/prune_tracks.cuh +++ b/device/cuda/src/finding/kernels/prune_tracks.cuh @@ -1,16 +1,17 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "traccc/edm/track_parameters.hpp" +// Project include(s). #include "traccc/finding/device/prune_tracks.hpp" namespace traccc::cuda::kernels { __global__ void prune_tracks(device::prune_tracks_payload payload); + } diff --git a/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu b/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu index 5fa6f073b2..9e0840f836 100644 --- a/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu +++ b/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu @@ -1,14 +1,18 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ +// Local include(s). #include "apply_interaction_src.cuh" +#include "types.hpp" namespace traccc::cuda::kernels { + template __global__ void apply_interaction( const finding_config, device::apply_interaction_payload); + } diff --git a/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh b/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh index 53b6b2e2a9..7623af52e7 100644 --- a/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh +++ b/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh @@ -1,15 +1,18 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once -#include "traccc/edm/track_parameters.hpp" +// Local include(s). +#include "../../../utils/global_index.hpp" +#include "../apply_interaction.cuh" + +// Project include(s). #include "traccc/finding/device/apply_interaction.hpp" -#include "traccc/geometry/detector.hpp" namespace traccc::cuda::kernels { @@ -18,8 +21,8 @@ __global__ void apply_interaction( const finding_config cfg, device::apply_interaction_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; - - device::apply_interaction(gid, cfg, payload); + device::apply_interaction(details::global_index1(), cfg, + payload); } + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu index c992a67e1d..be14dff104 100644 --- a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu +++ b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu @@ -1,12 +1,12 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ -#include "./types.hpp" #include "propagate_to_next_surface_src.cuh" +#include "types.hpp" namespace traccc::cuda::kernels { @@ -16,4 +16,5 @@ propagate_to_next_surface); + } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh index b4286c2d4a..505d1d580f 100644 --- a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh +++ b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh @@ -7,9 +7,12 @@ #pragma once +// Local include(s). +#include "../../../utils/global_index.hpp" +#include "../propagate_to_next_surface.cuh" + +// Project include(s). #include "traccc/finding/device/propagate_to_next_surface.hpp" -#include "traccc/finding/finding_config.hpp" -#include "traccc/geometry/detector.hpp" namespace traccc::cuda::kernels { @@ -19,7 +22,7 @@ __global__ void propagate_to_next_surface( device::propagate_to_next_surface_payload payload) { device::propagate_to_next_surface( - threadIdx.x + blockIdx.x * blockDim.x, cfg, payload); + details::global_index1(), cfg, payload); } } // namespace traccc::cuda::kernels diff --git a/device/cuda/src/fitting/fitting_algorithm.cu b/device/cuda/src/fitting/fitting_algorithm.cu index 1c1f747761..a521ec3059 100644 --- a/device/cuda/src/fitting/fitting_algorithm.cu +++ b/device/cuda/src/fitting/fitting_algorithm.cu @@ -1,12 +1,13 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ // Project include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/fitting/fitting_algorithm.hpp" #include "traccc/fitting/device/fill_sort_keys.hpp" @@ -33,8 +34,8 @@ __global__ void fill_sort_keys( vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view) { - device::fill_sort_keys(threadIdx.x + blockIdx.x * blockDim.x, - track_candidates_view, keys_view, ids_view); + device::fill_sort_keys(details::global_index1(), track_candidates_view, + keys_view, ids_view); } template @@ -45,10 +46,9 @@ __global__ void fit( vecmem::data::vector_view param_ids_view, track_state_container_types::view track_states_view) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; - - device::fit(gid, det_data, field_data, cfg, track_candidates_view, - param_ids_view, track_states_view); + device::fit(details::global_index1(), det_data, field_data, cfg, + track_candidates_view, param_ids_view, + track_states_view); } } // namespace kernels diff --git a/device/cuda/src/sanity/contiguous_on.cuh b/device/cuda/src/sanity/contiguous_on.cuh index 6e5d198338..668a611806 100644 --- a/device/cuda/src/sanity/contiguous_on.cuh +++ b/device/cuda/src/sanity/contiguous_on.cuh @@ -1,7 +1,7 @@ /** * traccc library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,6 +10,7 @@ // Project include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/utils/stream.hpp" @@ -38,7 +39,7 @@ requires std::regular_invocable().at(0))> __global__ void is_contiguous_on_compress_adjacent( P projection, VIEW _in, vecmem::data::vector_view out_view) { - int tid = threadIdx.x + blockIdx.x * blockDim.x; + const device::global_index_t tid = details::global_index1(); const CONTAINER in(_in); vecmem::device_vector out(out_view); @@ -60,8 +61,8 @@ template __global__ void is_contiguous_on_all_unique( vecmem::data::vector_view in_view, bool* out) { - int tid_x = threadIdx.x + blockIdx.x * blockDim.x; - int tid_y = threadIdx.y + blockIdx.y * blockDim.y; + const unsigned int tid_x = threadIdx.x + blockIdx.x * blockDim.x; + const unsigned int tid_y = threadIdx.y + blockIdx.y * blockDim.y; const vecmem::device_vector in(in_view); diff --git a/device/cuda/src/sanity/ordered_on.cuh b/device/cuda/src/sanity/ordered_on.cuh index fe360d2293..11dc5e71cd 100644 --- a/device/cuda/src/sanity/ordered_on.cuh +++ b/device/cuda/src/sanity/ordered_on.cuh @@ -1,7 +1,7 @@ /** * traccc library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,6 +10,7 @@ // Project include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/utils/stream.hpp" @@ -31,7 +32,8 @@ template requires std::regular_invocable().at(0)), decltype(std::declval().at(0))> __global__ void is_ordered_on_kernel(R relation, VIEW _in, bool* out) { - int tid = threadIdx.x + blockIdx.x * blockDim.x; + + const device::global_index_t tid = details::global_index1(); const CONTAINER in(_in); diff --git a/device/cuda/src/seeding/seed_finding.cu b/device/cuda/src/seeding/seed_finding.cu index e2ab6fd558..81dfc8f883 100644 --- a/device/cuda/src/seeding/seed_finding.cu +++ b/device/cuda/src/seeding/seed_finding.cu @@ -1,12 +1,13 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ // Local include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/seeding/seed_finding.hpp" @@ -44,9 +45,8 @@ __global__ void count_doublets( device::doublet_counter_collection_types::view doublet_counter, unsigned int& nMidBot, unsigned int& nMidTop) { - device::count_doublets(threadIdx.x + blockIdx.x * blockDim.x, config, - sp_grid, sp_prefix_sum, doublet_counter, nMidBot, - nMidTop); + device::count_doublets(details::global_index1(), config, sp_grid, + sp_prefix_sum, doublet_counter, nMidBot, nMidTop); } /// CUDA kernel for running @c traccc::device::find_doublets @@ -56,8 +56,8 @@ __global__ void find_doublets( device::device_doublet_collection_types::view mb_doublets, device::device_doublet_collection_types::view mt_doublets) { - device::find_doublets(threadIdx.x + blockIdx.x * blockDim.x, config, - sp_grid, doublet_counter, mb_doublets, mt_doublets); + device::find_doublets(details::global_index1(), config, sp_grid, + doublet_counter, mb_doublets, mt_doublets); } /// CUDA kernel for running @c traccc::device::count_triplets @@ -69,8 +69,8 @@ __global__ void count_triplets( device::triplet_counter_spM_collection_types::view spM_counter, device::triplet_counter_collection_types::view midBot_counter) { - device::count_triplets(threadIdx.x + blockIdx.x * blockDim.x, config, - sp_grid, doublet_counter, mb_doublets, mt_doublets, + device::count_triplets(details::global_index1(), config, sp_grid, + doublet_counter, mb_doublets, mt_doublets, spM_counter, midBot_counter); } @@ -80,8 +80,8 @@ __global__ void reduce_triplet_counts( device::triplet_counter_spM_collection_types::view spM_counter, unsigned int& num_triplets) { - device::reduce_triplet_counts(threadIdx.x + blockIdx.x * blockDim.x, - doublet_counter, spM_counter, num_triplets); + device::reduce_triplet_counts(details::global_index1(), doublet_counter, + spM_counter, num_triplets); } /// CUDA kernel for running @c traccc::device::find_triplets @@ -94,9 +94,9 @@ __global__ void find_triplets( device::triplet_counter_collection_types::const_view midBot_tc, device::device_triplet_collection_types::view triplet_view) { - device::find_triplets(threadIdx.x + blockIdx.x * blockDim.x, config, - filter_config, sp_grid, doublet_counter, mt_doublets, - spM_tc, midBot_tc, triplet_view); + device::find_triplets(details::global_index1(), config, filter_config, + sp_grid, doublet_counter, mt_doublets, spM_tc, + midBot_tc, triplet_view); } /// CUDA kernel for running @c traccc::device::update_triplet_weights @@ -112,9 +112,9 @@ __global__ void update_triplet_weights( // Each thread uses compatSeedLimit elements of the array scalar* dataPos = &data[threadIdx.x * filter_config.compatSeedLimit]; - device::update_triplet_weights(threadIdx.x + blockIdx.x * blockDim.x, - filter_config, sp_grid, spM_tc, midBot_tc, - dataPos, triplet_view); + device::update_triplet_weights(details::global_index1(), filter_config, + sp_grid, spM_tc, midBot_tc, dataPos, + triplet_view); } /// CUDA kernel for running @c traccc::device::select_seeds @@ -133,7 +133,7 @@ __global__ void select_seeds( // Each thread uses max_triplets_per_spM elements of the array triplet* dataPos = &data2[threadIdx.x * filter_config.max_triplets_per_spM]; - device::select_seeds(threadIdx.x + blockIdx.x * blockDim.x, filter_config, + device::select_seeds(details::global_index1(), filter_config, spacepoints_view, internal_sp_view, spM_tc, midBot_tc, triplet_view, dataPos, seed_view); } diff --git a/device/cuda/src/seeding/spacepoint_binning.cu b/device/cuda/src/seeding/spacepoint_binning.cu index 1c1297d68f..f767bc76b2 100644 --- a/device/cuda/src/seeding/spacepoint_binning.cu +++ b/device/cuda/src/seeding/spacepoint_binning.cu @@ -1,12 +1,13 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ // Local include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/seeding/spacepoint_binning.hpp" @@ -28,9 +29,8 @@ __global__ void count_grid_capacities( spacepoint_collection_types::const_view spacepoints, vecmem::data::vector_view grid_capacities) { - device::count_grid_capacities(threadIdx.x + blockIdx.x * blockDim.x, config, - phi_axis, z_axis, spacepoints, - grid_capacities); + device::count_grid_capacities(details::global_index1(), config, phi_axis, + z_axis, spacepoints, grid_capacities); } /// CUDA kernel for running @c traccc::device::populate_grid @@ -38,8 +38,7 @@ __global__ void populate_grid( seedfinder_config config, spacepoint_collection_types::const_view spacepoints, sp_grid_view grid) { - device::populate_grid(threadIdx.x + blockIdx.x * blockDim.x, config, - spacepoints, grid); + device::populate_grid(details::global_index1(), config, spacepoints, grid); } } // namespace kernels diff --git a/device/cuda/src/seeding/spacepoint_formation_algorithm.cu b/device/cuda/src/seeding/spacepoint_formation_algorithm.cu index c2f2751e4a..a2ae5b384e 100644 --- a/device/cuda/src/seeding/spacepoint_formation_algorithm.cu +++ b/device/cuda/src/seeding/spacepoint_formation_algorithm.cu @@ -1,12 +1,13 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ // Local include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/seeding/spacepoint_formation_algorithm.hpp" @@ -24,9 +25,9 @@ __global__ void __launch_bounds__(1024, 1) const unsigned int measurement_count, spacepoint_collection_types::view spacepoints_view) { - device::form_spacepoints(threadIdx.x + blockIdx.x * blockDim.x, - det_view, measurements_view, - measurement_count, spacepoints_view); + device::form_spacepoints(details::global_index1(), det_view, + measurements_view, measurement_count, + spacepoints_view); } } // namespace kernels diff --git a/device/cuda/src/seeding/track_params_estimation.cu b/device/cuda/src/seeding/track_params_estimation.cu index fd602e3461..aa17b81b74 100644 --- a/device/cuda/src/seeding/track_params_estimation.cu +++ b/device/cuda/src/seeding/track_params_estimation.cu @@ -1,12 +1,13 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ // Local include(s). #include "../utils/cuda_error_handling.hpp" +#include "../utils/global_index.hpp" #include "../utils/utils.hpp" #include "traccc/cuda/seeding/track_params_estimation.hpp" @@ -26,9 +27,8 @@ __global__ void estimate_track_params( const std::array stddev, bound_track_parameters_collection_types::view params_view) { - device::estimate_track_params(threadIdx.x + blockIdx.x * blockDim.x, - spacepoints_view, seed_view, bfield, stddev, - params_view); + device::estimate_track_params(details::global_index1(), spacepoints_view, + seed_view, bfield, stddev, params_view); } } // namespace kernels diff --git a/device/cuda/src/utils/global_index.hpp b/device/cuda/src/utils/global_index.hpp new file mode 100644 index 0000000000..fce8d392b8 --- /dev/null +++ b/device/cuda/src/utils/global_index.hpp @@ -0,0 +1,21 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/device/global_index.hpp" + +namespace traccc::cuda::details { + +/// Function creating a global index in a 1D CUDA kernel +__device__ inline device::global_index_t global_index1() { + + return blockIdx.x * blockDim.x + threadIdx.x; +} + +} // namespace traccc::cuda::details diff --git a/device/cuda/src/utils/thread_id.hpp b/device/cuda/src/utils/thread_id.hpp index c913631103..c9ab9f2201 100644 --- a/device/cuda/src/utils/thread_id.hpp +++ b/device/cuda/src/utils/thread_id.hpp @@ -9,38 +9,35 @@ #pragma once // Project include(s). -#include "traccc/definitions/qualifiers.hpp" #include "traccc/device/concepts/thread_id.hpp" namespace traccc::cuda::details { /// A CUDA thread identifier type struct thread_id1 { - TRACCC_DEVICE thread_id1() {} + __device__ thread_id1() {} - inline unsigned int TRACCC_DEVICE getLocalThreadId() const { + inline unsigned int __device__ getLocalThreadId() const { return threadIdx.x; } - inline unsigned int TRACCC_DEVICE getLocalThreadIdX() const { + inline unsigned int __device__ getLocalThreadIdX() const { return threadIdx.x; } - inline unsigned int TRACCC_DEVICE getGlobalThreadId() const { + inline unsigned int __device__ getGlobalThreadId() const { return threadIdx.x + blockIdx.x * blockDim.x; } - inline unsigned int TRACCC_DEVICE getGlobalThreadIdX() const { + inline unsigned int __device__ getGlobalThreadIdX() const { return threadIdx.x + blockIdx.x * blockDim.x; } - inline unsigned int TRACCC_DEVICE getBlockIdX() const { return blockIdx.x; } + inline unsigned int __device__ getBlockIdX() const { return blockIdx.x; } - inline unsigned int TRACCC_DEVICE getBlockDimX() const { - return blockDim.x; - } + inline unsigned int __device__ getBlockDimX() const { return blockDim.x; } - inline unsigned int TRACCC_DEVICE getGridDimX() const { return gridDim.x; } + inline unsigned int __device__ getGridDimX() const { return gridDim.x; } }; // struct thread_id1 From 74badb90811471b12bad9d913d877c45b2151308 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 8 Jan 2025 23:26:01 +0100 Subject: [PATCH 4/5] Call device functions consistently from SYCL. --- device/sycl/src/fitting/fit_tracks.hpp | 7 ++-- device/sycl/src/seeding/seed_finding.sycl | 29 ++++++++------- .../silicon_pixel_spacepoint_formation.hpp | 5 +-- .../sycl/src/seeding/spacepoint_binning.sycl | 14 ++++---- .../src/seeding/track_params_estimation.sycl | 7 ++-- device/sycl/src/utils/global_index.hpp | 36 +++++++++++++++++++ 6 files changed, 67 insertions(+), 31 deletions(-) create mode 100644 device/sycl/src/utils/global_index.hpp diff --git a/device/sycl/src/fitting/fit_tracks.hpp b/device/sycl/src/fitting/fit_tracks.hpp index 41140dccb8..f8ef1c3ce0 100644 --- a/device/sycl/src/fitting/fit_tracks.hpp +++ b/device/sycl/src/fitting/fit_tracks.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -9,6 +9,7 @@ // Local include(s). #include "../utils/calculate1DimNdRange.hpp" +#include "../utils/global_index.hpp" // Project include(s). #include "traccc/edm/device/sort_key.hpp" @@ -96,7 +97,7 @@ track_state_container_types::buffer fit_tracks( [track_candidates_view, keys_view = vecmem::get_data(keys_buffer), param_ids_view = vecmem::get_data(param_ids_buffer)](::sycl::nd_item<1> item) { - device::fill_sort_keys(item.get_global_linear_id(), + device::fill_sort_keys(details::global_index(item), track_candidates_view, keys_view, param_ids_view); }); @@ -120,7 +121,7 @@ track_state_container_types::buffer fit_tracks( range, [det_view, field_view, config, track_candidates_view, param_ids_view = vecmem::get_data(param_ids_buffer), track_states_view](::sycl::nd_item<1> item) { - device::fit(item.get_global_linear_id(), det_view, + device::fit(details::global_index(item), det_view, field_view, config, track_candidates_view, param_ids_view, track_states_view); diff --git a/device/sycl/src/seeding/seed_finding.sycl b/device/sycl/src/seeding/seed_finding.sycl index 9ecb5d4f4a..09f4bc8e47 100644 --- a/device/sycl/src/seeding/seed_finding.sycl +++ b/device/sycl/src/seeding/seed_finding.sycl @@ -1,19 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ -// System include(s). -#include - -// SYCL library include(s). -#include "traccc/sycl/seeding/seed_finding.hpp" - -// SYCL library include(s). +// Local include(s). #include "../utils/calculate1DimNdRange.hpp" #include "../utils/get_queue.hpp" +#include "../utils/global_index.hpp" +#include "traccc/sycl/seeding/seed_finding.hpp" #include "traccc/sycl/utils/make_prefix_sum_buff.hpp" // Project include(s). @@ -34,6 +30,9 @@ // VecMem include(s). #include +// System include(s). +#include + namespace traccc::sycl { namespace kernels { @@ -123,7 +122,7 @@ seed_finding::output_type seed_finding::operator()( [config = m_seedfinder_config, g2_view, sp_grid_prefix_sum_view, doublet_counter_view, aux_globalCounter](::sycl::nd_item<1> item) { - device::count_doublets(item.get_global_linear_id(), config, + device::count_doublets(details::global_index(item), config, g2_view, sp_grid_prefix_sum_view, doublet_counter_view, (*aux_globalCounter).m_nMidBot, @@ -171,7 +170,7 @@ seed_finding::output_type seed_finding::operator()( doubletFindRange, [config = m_seedfinder_config, g2_view, doublet_counter_view, mb_view, mt_view](::sycl::nd_item<1> item) { - device::find_doublets(item.get_global_linear_id(), config, + device::find_doublets(details::global_index(item), config, g2_view, doublet_counter_view, mb_view, mt_view); }); @@ -210,7 +209,7 @@ seed_finding::output_type seed_finding::operator()( mb_view, mt_view, triplet_counter_spM_view, triplet_counter_midBot_view](::sycl::nd_item<1> item) { device::count_triplets( - item.get_global_linear_id(), config, g2_view, + details::global_index(item), config, g2_view, doublet_counter_view, mb_view, mt_view, triplet_counter_spM_view, triplet_counter_midBot_view); }); @@ -232,7 +231,7 @@ seed_finding::output_type seed_finding::operator()( [doublet_counter_view, triplet_counter_spM_view, aux_globalCounter](::sycl::nd_item<1> item) { device::reduce_triplet_counts( - item.get_global_linear_id(), doublet_counter_view, + details::global_index(item), doublet_counter_view, triplet_counter_spM_view, (*aux_globalCounter).m_nTriplets); }); @@ -270,7 +269,7 @@ seed_finding::output_type seed_finding::operator()( triplet_counter_midBot_view, triplet_view](::sycl::nd_item<1> item) { device::find_triplets( - item.get_global_linear_id(), config, filter_config, + details::global_index(item), config, filter_config, g2_view, doublet_counter_view, mt_view, triplet_counter_spM_view, triplet_counter_midBot_view, triplet_view); @@ -311,7 +310,7 @@ seed_finding::output_type seed_finding::operator()( filter_config.compatSeedLimit]; device::update_triplet_weights( - item.get_global_linear_id(), filter_config, g2_view, + details::global_index(item), filter_config, g2_view, triplet_counter_spM_view, triplet_counter_midBot_view, dataPos, triplet_view); }); @@ -359,7 +358,7 @@ seed_finding::output_type seed_finding::operator()( &local_mem[item.get_local_id() * filter_config.max_triplets_per_spM]; - device::select_seeds(item.get_global_linear_id(), + device::select_seeds(details::global_index(item), filter_config, spacepoints_view, g2_view, triplet_counter_spM_view, triplet_counter_midBot_view, diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp index ecd76b04a5..a1ef106cb3 100644 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023-2024 CERN for the benefit of the ACTS project + * (c) 2023-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,6 +10,7 @@ // Local include(s). #include "../utils/calculate1DimNdRange.hpp" #include "../utils/get_queue.hpp" +#include "../utils/global_index.hpp" // Project include(s). #include "traccc/edm/measurement.hpp" @@ -69,7 +70,7 @@ spacepoint_collection_types::buffer silicon_pixel_spacepoint_formation( spacepoints_view = vecmem::get_data(result)]( ::sycl::nd_item<1> item) { device::form_spacepoints( - item.get_global_linear_id(), det_view, + details::global_index(item), det_view, measurements_view, n_measurements, spacepoints_view); }); }) diff --git a/device/sycl/src/seeding/spacepoint_binning.sycl b/device/sycl/src/seeding/spacepoint_binning.sycl index f3974667c8..5545c2b465 100644 --- a/device/sycl/src/seeding/spacepoint_binning.sycl +++ b/device/sycl/src/seeding/spacepoint_binning.sycl @@ -1,16 +1,15 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ // Local include(s). #include "../utils/calculate1DimNdRange.hpp" -#include "traccc/sycl/seeding/spacepoint_binning.hpp" - -// Local include(s). #include "../utils/get_queue.hpp" +#include "../utils/global_index.hpp" +#include "traccc/sycl/seeding/spacepoint_binning.hpp" // Project include(s). #include "traccc/seeding/device/count_grid_capacities.hpp" @@ -72,7 +71,7 @@ sp_grid_buffer spacepoint_binning::operator()( z_axis = m_axes.second, spacepoints = spacepoints_view, grid_capacities = grid_capacities_view](::sycl::nd_item<1> item) { - device::count_grid_capacities(item.get_global_linear_id(), + device::count_grid_capacities(details::global_index(item), config, phi_axis, z_axis, spacepoints, grid_capacities); }); @@ -99,9 +98,8 @@ sp_grid_buffer spacepoint_binning::operator()( h.parallel_for( range, [config = m_config, spacepoints = spacepoints_view, grid = grid_view](::sycl::nd_item<1> item) { - device::populate_grid( - static_cast(item.get_global_linear_id()), - config, spacepoints, grid); + device::populate_grid(details::global_index(item), config, + spacepoints, grid); }); }) .wait_and_throw(); diff --git a/device/sycl/src/seeding/track_params_estimation.sycl b/device/sycl/src/seeding/track_params_estimation.sycl index d919015184..3ab9fd172b 100644 --- a/device/sycl/src/seeding/track_params_estimation.sycl +++ b/device/sycl/src/seeding/track_params_estimation.sycl @@ -1,13 +1,14 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2021-2024 CERN for the benefit of the ACTS project + * (c) 2021-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ -// SYCL library include(s). +// Local include(s). #include "../utils/calculate1DimNdRange.hpp" #include "../utils/get_queue.hpp" +#include "../utils/global_index.hpp" #include "traccc/sycl/seeding/track_params_estimation.hpp" // Project include(s). @@ -63,7 +64,7 @@ track_params_estimation::output_type track_params_estimation::operator()( trackParamsNdRange, [spacepoints_view, seeds_view, bfield, stddev, params_view](::sycl::nd_item<1> item) { - device::estimate_track_params(item.get_global_linear_id(), + device::estimate_track_params(details::global_index(item), spacepoints_view, seeds_view, bfield, stddev, params_view); }); diff --git a/device/sycl/src/utils/global_index.hpp b/device/sycl/src/utils/global_index.hpp new file mode 100644 index 0000000000..d455469156 --- /dev/null +++ b/device/sycl/src/utils/global_index.hpp @@ -0,0 +1,36 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/device/global_index.hpp" + +// SYCL include(s). +#include + +namespace traccc::sycl::details { + +/// Function creating a global index in a 1D SYCL kernel +inline device::global_index_t global_index(const ::sycl::nd_item<1>& item) { + + return static_cast(item.get_global_linear_id()); +} + +/// Function creating a global index in a 2D SYCL kernel +inline device::global_index_t global_index(const ::sycl::nd_item<2>& item) { + + return static_cast(item.get_global_linear_id()); +} + +/// Function creating a global index in a 3D SYCL kernel +inline device::global_index_t global_index(const ::sycl::nd_item<3>& item) { + + return static_cast(item.get_global_linear_id()); +} + +} // namespace traccc::sycl::details From b0ffe4c9ce636fe0a9852f7e798662921f9c8e8c Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Thu, 9 Jan 2025 11:06:14 +0100 Subject: [PATCH 5/5] Implement (some of) Stephen's suggestions. --- device/alpaka/src/utils/thread_id.hpp | 29 +++++++++++++----------- device/cuda/src/sanity/contiguous_on.cuh | 4 ++-- 2 files changed, 18 insertions(+), 15 deletions(-) diff --git a/device/alpaka/src/utils/thread_id.hpp b/device/alpaka/src/utils/thread_id.hpp index 832ca5a432..cbca42093f 100644 --- a/device/alpaka/src/utils/thread_id.hpp +++ b/device/alpaka/src/utils/thread_id.hpp @@ -25,34 +25,37 @@ template struct thread_id1 { TRACCC_HOST_DEVICE explicit thread_id1(const Acc& acc) : m_acc(acc) {} - auto inline TRACCC_HOST_DEVICE getLocalThreadId() const { - return ::alpaka::getIdx<::alpaka::Block, ::alpaka::Threads>(m_acc)[0u]; + unsigned int inline TRACCC_HOST_DEVICE getLocalThreadId() const { + return static_cast( + ::alpaka::getIdx<::alpaka::Block, ::alpaka::Threads>(m_acc)[0u]); } - auto inline TRACCC_HOST_DEVICE getLocalThreadIdX() const { + unsigned int inline TRACCC_HOST_DEVICE getLocalThreadIdX() const { return getLocalThreadId(); } - auto inline TRACCC_HOST_DEVICE getGlobalThreadId() const { + unsigned int inline TRACCC_HOST_DEVICE getGlobalThreadId() const { return getLocalThreadId() + getBlockIdX() * getBlockDimX(); } - auto inline TRACCC_HOST_DEVICE getGlobalThreadIdX() const { + unsigned int inline TRACCC_HOST_DEVICE getGlobalThreadIdX() const { return getLocalThreadId() + getBlockIdX() * getBlockDimX(); } - auto inline TRACCC_HOST_DEVICE getBlockIdX() const { - return ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Blocks>(m_acc)[0u]; + unsigned int inline TRACCC_HOST_DEVICE getBlockIdX() const { + return static_cast( + ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Blocks>(m_acc)[0u]); } - auto inline TRACCC_HOST_DEVICE getBlockDimX() const { - return ::alpaka::getWorkDiv<::alpaka::Block, ::alpaka::Threads>( - m_acc)[0u]; + unsigned int inline TRACCC_HOST_DEVICE getBlockDimX() const { + return static_cast( + ::alpaka::getWorkDiv<::alpaka::Block, ::alpaka::Threads>( + m_acc)[0u]); } - auto inline TRACCC_HOST_DEVICE getGridDimX() const { - return ::alpaka::getWorkDiv<::alpaka::Grid, ::alpaka::Blocks>( - m_acc)[0u]; + unsigned int inline TRACCC_HOST_DEVICE getGridDimX() const { + return static_cast( + ::alpaka::getWorkDiv<::alpaka::Grid, ::alpaka::Blocks>(m_acc)[0u]); } private: diff --git a/device/cuda/src/sanity/contiguous_on.cuh b/device/cuda/src/sanity/contiguous_on.cuh index 668a611806..81963fcd65 100644 --- a/device/cuda/src/sanity/contiguous_on.cuh +++ b/device/cuda/src/sanity/contiguous_on.cuh @@ -61,8 +61,8 @@ template __global__ void is_contiguous_on_all_unique( vecmem::data::vector_view in_view, bool* out) { - const unsigned int tid_x = threadIdx.x + blockIdx.x * blockDim.x; - const unsigned int tid_y = threadIdx.y + blockIdx.y * blockDim.y; + const device::global_index_t tid_x = threadIdx.x + blockIdx.x * blockDim.x; + const device::global_index_t tid_y = threadIdx.y + blockIdx.y * blockDim.y; const vecmem::device_vector in(in_view);