src/gbts_seeding/gbts_seeding_algorithm.cu

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2025-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 // Local include(s).
0009 #include "../utils/barrier.hpp"
0010 #include "../utils/cuda_error_handling.hpp"
0011 #include "../utils/thread_id.hpp"
0012 #include "../utils/utils.hpp"
0013 #include "traccc/cuda/gbts_seeding/gbts_seeding_algorithm.hpp"
0014
0015 // Project include(s).
0016 #include "traccc/gbts_seeding/device/gbts_add_terminus_to_path_store.hpp"
0017 #include "traccc/gbts_seeding/device/gbts_bid_seeds_for_hits.hpp"
0018 #include "traccc/gbts_seeding/device/gbts_bin_spacepoints.hpp"
0019 #include "traccc/gbts_seeding/device/gbts_compress_graph.hpp"
0020 #include "traccc/gbts_seeding/device/gbts_convert_seeds.hpp"
0021 #include "traccc/gbts_seeding/device/gbts_count_eta_phi_bins.hpp"
0022 #include "traccc/gbts_seeding/device/gbts_count_spacepoints_by_layer.hpp"
0023 #include "traccc/gbts_seeding/device/gbts_count_terminus_edges.hpp"
0024 #include "traccc/gbts_seeding/device/gbts_fill_path_store.hpp"
0025 #include "traccc/gbts_seeding/device/gbts_find_minmax_radius.hpp"
0026 #include "traccc/gbts_seeding/device/gbts_fit_segments.hpp"
0027 #include "traccc/gbts_seeding/device/gbts_link_graph_edges.hpp"
0028 #include "traccc/gbts_seeding/device/gbts_make_graph_edges.hpp"
0029 #include "traccc/gbts_seeding/device/gbts_match_graph_edges.hpp"
0030 #include "traccc/gbts_seeding/device/gbts_prefix_sum_eta_phi_bins.hpp"
0031 #include "traccc/gbts_seeding/device/gbts_rebid_seeds_for_edges.hpp"
0032 #include "traccc/gbts_seeding/device/gbts_reindex_edges.hpp"
0033 #include "traccc/gbts_seeding/device/gbts_reset_edge_bids.hpp"
0034 #include "traccc/gbts_seeding/device/gbts_run_cca_iteration.hpp"
0035 #include "traccc/gbts_seeding/device/gbts_sort_nodes.hpp"
0036 #include "traccc/gbts_seeding/gbts_types.hpp"
0037
0038 // VecMem include(s).
0039 #include <vecmem/containers/data/vector_view.hpp>
0040
0041 // System include(s).
0042 #include <algorithm>
0043 #include <memory_resource>
0044
0045 // Thrust include(s).
0046 #include <thrust/execution_policy.h>
0047 #include <thrust/scan.h>
0048
0049 namespace traccc::cuda {
0050
0051 namespace kernels {
0052
0053 using float4 = traccc::float4;
0054 using uint2 = traccc::uint2;
0055 using int2 = traccc::int2;
0056
0057 // ---------------------------------------------------------------------------
0058 // Stage 1 — nodes-making kernels
0059 // ---------------------------------------------------------------------------
0060
0061 /// CUDA kernel for running @c traccc::device::gbts_count_spacepoints_by_layer
0062 __global__ void gbts_count_spacepoints_by_layer(
0063     const device::gbts_count_spacepoints_by_layer_payload payload) {
0064   device::gbts_count_spacepoints_by_layer(details::thread_id1{}, payload);
0065 }
0066
0067 /// CUDA kernel for running @c traccc::device::gbts_bin_spacepoints
0068 __global__ void gbts_bin_spacepoints(
0069     const device::gbts_bin_spacepoints_payload payload) {
0070   device::gbts_bin_spacepoints(details::thread_id1{}, payload);
0071 }
0072
0073 /// CUDA kernel for running @c traccc::device::gbts_count_eta_phi_bins
0074 __global__ void gbts_count_eta_phi_bins(
0075     const device::gbts_count_eta_phi_bins_payload payload) {
0076   device::gbts_count_eta_phi_bins(details::thread_id1{}, payload);
0077 }
0078
0079 /// CUDA kernel for running @c traccc::device::gbts_prefix_sum_eta_phi_bins
0080 __global__ void gbts_prefix_sum_eta_phi_bins(
0081     const device::gbts_prefix_sum_eta_phi_bins_payload payload) {
0082   device::gbts_prefix_sum_eta_phi_bins(details::thread_id1{}, payload);
0083 }
0084
0085 /// CUDA kernel for running @c traccc::device::gbts_sort_nodes
0086 __global__ void gbts_sort_nodes(const device::gbts_sort_nodes_payload payload) {
0087   device::gbts_sort_nodes(details::thread_id1{}, payload);
0088 }
0089
0090 /// CUDA kernel for running @c traccc::device::gbts_find_minmax_radius
0091 __global__ void gbts_find_minmax_radius(
0092     const device::gbts_find_minmax_radius_payload payload) {
0093   device::gbts_find_minmax_radius(details::thread_id1{}, payload);
0094 }
0095
0096 // ---------------------------------------------------------------------------
0097 // Stage 2 — graph-making kernels
0098 // ---------------------------------------------------------------------------
0099
0100 /// CUDA kernel for running @c traccc::device::gbts_make_graph_edges
0101 __global__ void gbts_make_graph_edges(
0102     const device::gbts_make_graph_edges_payload payload) {
0103   __shared__ float phi[traccc::device::gbts_consts::node_buffer_length];
0104   __shared__ float4 node_pack[traccc::device::gbts_consts::node_buffer_length];
0105   const traccc::cuda::barrier barrier;
0106
0107   device::gbts_make_graph_edges(
0108       details::thread_id1{}, barrier, payload,
0109       {vecmem::data::vector_view<float>(
0110            traccc::device::gbts_consts::node_buffer_length, phi),
0111        vecmem::data::vector_view<float4>(
0112            traccc::device::gbts_consts::node_buffer_length, node_pack)});
0113 }
0114
0115 /// CUDA kernel for running @c traccc::device::gbts_link_graph_edges
0116 __global__ void gbts_link_graph_edges(
0117     const device::gbts_link_graph_edges_payload payload) {
0118   device::gbts_link_graph_edges(details::thread_id1{}, payload);
0119 }
0120
0121 /// CUDA kernel for running @c traccc::device::gbts_match_graph_edges
0122 __global__ void gbts_match_graph_edges(
0123     const device::gbts_match_graph_edges_payload payload) {
0124   device::gbts_match_graph_edges(details::thread_id1{}, payload);
0125 }
0126
0127 /// CUDA kernel for running @c traccc::device::gbts_reindex_edges
0128 __global__ void gbts_reindex_edges(
0129     const device::gbts_reindex_edges_payload payload) {
0130   device::gbts_reindex_edges(details::thread_id1{}, payload);
0131 }
0132
0133 /// CUDA kernel for running @c traccc::device::gbts_compress_graph
0134 __global__ void gbts_compress_graph(
0135     const device::gbts_compress_graph_payload payload) {
0136   device::gbts_compress_graph(details::thread_id1{}, payload);
0137 }
0138
0139 // ---------------------------------------------------------------------------
0140 // Stage 3 — graph-processing kernels
0141 // ---------------------------------------------------------------------------
0142
0143 /// CUDA kernel for running @c traccc::device::gbts_run_cca_iteration
0144 __global__ void gbts_run_cca_iteration(
0145     const device::gbts_run_cca_iteration_payload payload) {
0146   device::gbts_run_cca_iteration(details::thread_id1{}, payload);
0147 }
0148
0149 /// CUDA kernel for running @c traccc::device::gbts_count_terminus_edges
0150 __global__ void gbts_count_terminus_edges(
0151     const device::gbts_count_terminus_edges_payload payload) {
0152   device::gbts_count_terminus_edges(details::thread_id1{}, payload);
0153 }
0154
0155 /// CUDA kernel for running @c traccc::device::gbts_add_terminus_to_path_store
0156 __global__ void gbts_add_terminus_to_path_store(
0157     const device::gbts_add_terminus_to_path_store_payload payload) {
0158   device::gbts_add_terminus_to_path_store(details::thread_id1{}, payload);
0159 }
0160
0161 /// CUDA kernel for running @c traccc::device::gbts_fill_path_store
0162 __global__ void gbts_fill_path_store(
0163     const device::gbts_fill_path_store_payload payload) {
0164   __shared__ traccc::uint2
0165       live_paths[traccc::device::gbts_consts::live_path_buffer];
0166   __shared__ int n_live_paths;
0167   const traccc::cuda::barrier barrier;
0168
0169   device::gbts_fill_path_store(
0170       details::thread_id1{}, barrier, payload,
0171       {vecmem::data::vector_view<traccc::uint2>(
0172            traccc::device::gbts_consts::live_path_buffer, live_paths),
0173        n_live_paths});
0174 }
0175
0176 /// CUDA kernel for running @c traccc::device::gbts_fit_segments
0177 __global__ void gbts_fit_segments(
0178     const device::gbts_fit_segments_payload payload) {
0179   device::gbts_fit_segments(details::thread_id1{}, payload);
0180 }
0181
0182 /// CUDA kernel for running @c traccc::device::gbts_reset_edge_bids
0183 __global__ void gbts_reset_edge_bids(
0184     const device::gbts_reset_edge_bids_payload payload) {
0185   device::gbts_reset_edge_bids(details::thread_id1{}, payload);
0186 }
0187
0188 /// CUDA kernel for running @c traccc::device::gbts_rebid_seeds_for_edges
0189 __global__ void gbts_rebid_seeds_for_edges(
0190     const device::gbts_rebid_seeds_for_edges_payload payload) {
0191   device::gbts_rebid_seeds_for_edges(details::thread_id1{}, payload);
0192 }
0193
0194 /// CUDA kernel for running @c traccc::device::gbts_bid_seeds_for_hits
0195 __global__ void gbts_bid_seeds_for_hits(
0196     const device::gbts_bid_seeds_for_hits_payload payload) {
0197   device::gbts_bid_seeds_for_hits(details::thread_id1{}, payload);
0198 }
0199
0200 /// CUDA kernel for running @c traccc::device::gbts_convert_seeds
0201 __global__ void gbts_convert_seeds(
0202     const device::gbts_convert_seeds_payload payload) {
0203   device::gbts_convert_seeds(details::thread_id1{}, payload);
0204 }
0205
0206 }  // namespace kernels
0207
0208 // ===========================================================================
0209 // gbts_seeding_algorithm: kernel launchers
0210 // ===========================================================================
0211
0212 gbts_seeding_algorithm::gbts_seeding_algorithm(
0213     const gbts_seedfinder_config& cfg, const memory_resource& mr,
0214     const vecmem::copy& copy, const stream_wrapper& str,
0215     std::unique_ptr<const Logger> logger)
0216     : device::gbts_seeding_algorithm(cfg, mr, copy, std::move(logger)),
0217       cuda::algorithm_base{str} {}
0218
0219 void gbts_seeding_algorithm::gbts_count_spacepoints_by_layer_kernel(
0220     const device::gbts_count_spacepoints_by_layer_payload& payload) const {
0221   const unsigned int n_threads = 128;
0222   const unsigned int n_blocks = 1 + (payload.nSp - 1) / n_threads;
0223   kernels::gbts_count_spacepoints_by_layer<<<n_blocks, n_threads, 0,
0224                                              details::get_stream(stream())>>>(
0225       payload);
0226   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0227   vecmem::device_vector<unsigned int> d_layer_sums(payload.layerCounts);
0228   thrust::inclusive_scan(
0229       thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0230           .on(details::get_stream(stream())),
0231       d_layer_sums.begin(), d_layer_sums.end(), d_layer_sums.begin());
0232 }
0233
0234 void gbts_seeding_algorithm::gbts_bin_spacepoints_kernel(
0235     const device::gbts_bin_spacepoints_payload& payload) const {
0236   const unsigned int n_threads = 128;
0237   const unsigned int n_blocks = 1 + (payload.nSp - 1) / n_threads;
0238   kernels::gbts_bin_spacepoints<<<n_blocks, n_threads, 0,
0239                                   details::get_stream(stream())>>>(payload);
0240   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0241 }
0242
0243 void gbts_seeding_algorithm::gbts_count_eta_phi_bins_kernel(
0244     const device::gbts_count_eta_phi_bins_payload& payload) const {
0245   const unsigned int n_threads = 128;
0246   const unsigned int n_blocks = 1 + (payload.nEtaBins - 1) / n_threads;
0247   kernels::gbts_count_eta_phi_bins<<<n_blocks, n_threads, 0,
0248                                      details::get_stream(stream())>>>(payload);
0249   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0250   vecmem::device_vector<unsigned int> d_eta_sums(payload.eta_node_counter);
0251   thrust::inclusive_scan(
0252       thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0253           .on(details::get_stream(stream())),
0254       d_eta_sums.begin(), d_eta_sums.end(), d_eta_sums.begin());
0255 }
0256
0257 void gbts_seeding_algorithm::gbts_prefix_sum_eta_phi_bins_kernel(
0258     const device::gbts_prefix_sum_eta_phi_bins_payload& payload) const {
0259   const unsigned int n_threads = 128;
0260   const unsigned int n_blocks = 1 + (payload.nEtaBins - 1) / n_threads;
0261   kernels::gbts_prefix_sum_eta_phi_bins<<<n_blocks, n_threads, 0,
0262                                           details::get_stream(stream())>>>(
0263       payload);
0264   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0265 }
0266
0267 void gbts_seeding_algorithm::gbts_sort_nodes_kernel(
0268     const device::gbts_sort_nodes_payload& payload) const {
0269   const unsigned int n_threads = 256;
0270   const unsigned int n_blocks = 1 + (payload.nNodes - 1) / n_threads;
0271   kernels::gbts_sort_nodes<<<n_blocks, n_threads, 0,
0272                              details::get_stream(stream())>>>(payload);
0273   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0274 }
0275
0276 void gbts_seeding_algorithm::gbts_find_minmax_radius_kernel(
0277     const device::gbts_find_minmax_radius_payload& payload) const {
0278   const unsigned int n_threads = 128;
0279   const unsigned int n_blocks = 1 + (payload.nEtaBins - 1) / n_threads;
0280   kernels::gbts_find_minmax_radius<<<n_blocks, n_threads, 0,
0281                                      details::get_stream(stream())>>>(payload);
0282   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0283 }
0284
0285 void gbts_seeding_algorithm::gbts_make_graph_edges_kernel(
0286     const device::gbts_make_graph_edges_payload& payload) const {
0287   const unsigned int n_threads = 128;
0288   const unsigned int n_blocks = payload.nUsedBinPairs;
0289   kernels::gbts_make_graph_edges<<<n_blocks, n_threads, 0,
0290                                    details::get_stream(stream())>>>(payload);
0291   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0292   vecmem::device_vector<unsigned int> d_num_outgoing_edges(
0293       payload.num_outgoing_edges);
0294   thrust::inclusive_scan(
0295       thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0296           .on(details::get_stream(stream())),
0297       d_num_outgoing_edges.begin(), d_num_outgoing_edges.end(),
0298       d_num_outgoing_edges.begin());
0299 }
0300
0301 void gbts_seeding_algorithm::gbts_link_graph_edges_kernel(
0302     const device::gbts_link_graph_edges_payload& payload) const {
0303   const unsigned int n_threads = 256;
0304   const unsigned int n_blocks = 1 + (payload.nEdges - 1) / n_threads;
0305   kernels::gbts_link_graph_edges<<<n_blocks, n_threads, 0,
0306                                    details::get_stream(stream())>>>(payload);
0307   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0308 }
0309
0310 void gbts_seeding_algorithm::gbts_match_graph_edges_kernel(
0311     const device::gbts_match_graph_edges_payload& payload) const {
0312   const unsigned int n_threads = 256;
0313   const unsigned int n_blocks = 1 + (payload.nEdges - 1) / n_threads;
0314   kernels::gbts_match_graph_edges<<<n_blocks, n_threads, 0,
0315                                     details::get_stream(stream())>>>(payload);
0316   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0317 }
0318
0319 void gbts_seeding_algorithm::gbts_reindex_edges_kernel(
0320     const device::gbts_reindex_edges_payload& payload) const {
0321   const unsigned int n_threads = 256;
0322   const unsigned int n_blocks = 1 + (payload.nEdges - 1) / n_threads;
0323   kernels::gbts_reindex_edges<<<n_blocks, n_threads, 0,
0324                                 details::get_stream(stream())>>>(payload);
0325   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0326 }
0327
0328 void gbts_seeding_algorithm::gbts_compress_graph_kernel(
0329     const device::gbts_compress_graph_payload& payload) const {
0330   const unsigned int n_threads = 256;
0331   const unsigned int n_blocks = 1 + (payload.nEdges - 1) / n_threads;
0332   kernels::gbts_compress_graph<<<n_blocks, n_threads, 0,
0333                                  details::get_stream(stream())>>>(payload);
0334   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0335 }
0336
0337 void gbts_seeding_algorithm::gbts_run_cca_iteration_kernel(
0338     const device::gbts_run_cca_iteration_payload& payload) const {
0339   const unsigned int n_threads = 128;
0340   const unsigned int n_blocks = 1 + (payload.nConnectedEdges - 1) / n_threads;
0341
0342   kernels::gbts_run_cca_iteration<<<n_blocks, n_threads, 0,
0343                                     details::get_stream(stream())>>>(payload);
0344   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0345 }
0346
0347 void gbts_seeding_algorithm::gbts_count_terminus_edges_kernel(
0348     const device::gbts_count_terminus_edges_payload& payload) const {
0349   const unsigned int n_threads = 128;
0350   const unsigned int n_blocks = 1 + (payload.nConnectedEdges - 1) / n_threads;
0351   kernels::gbts_count_terminus_edges<<<n_blocks, n_threads, 0,
0352                                        details::get_stream(stream())>>>(
0353       payload);
0354   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0355 }
0356
0357 void gbts_seeding_algorithm::gbts_add_terminus_to_path_store_kernel(
0358     const device::gbts_add_terminus_to_path_store_payload& payload) const {
0359   const unsigned int n_threads = 128;
0360   const unsigned int n_blocks = 1 + (payload.nConnectedEdges - 1) / n_threads;
0361   kernels::gbts_add_terminus_to_path_store<<<n_blocks, n_threads, 0,
0362                                              details::get_stream(stream())>>>(
0363       payload);
0364   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0365 }
0366
0367 void gbts_seeding_algorithm::gbts_fill_path_store_kernel(
0368     const device::gbts_fill_path_store_payload& payload) const {
0369   const unsigned int n_threads = 128;
0370   const unsigned int pathsPerTerminus =
0371       1 + (payload.nPaths - 1) / payload.nTerminusEdges;
0372   const unsigned int terminusPerBlock = std::min(
0373       n_threads, 1 + (traccc::device::gbts_consts::live_path_buffer - 1) /
0374                          pathsPerTerminus);
0375   const unsigned int n_blocks =
0376       1 + (payload.nTerminusEdges - 1) / terminusPerBlock;
0377   kernels::gbts_fill_path_store<<<n_blocks, n_threads, 0,
0378                                   details::get_stream(stream())>>>(payload);
0379   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0380 }
0381
0382 void gbts_seeding_algorithm::gbts_fit_segments_kernel(
0383     const device::gbts_fit_segments_payload& payload) const {
0384   const unsigned int n_threads = 128;
0385   const unsigned int n_blocks = 1 + (payload.nPaths - 1) / n_threads;
0386   kernels::gbts_fit_segments<<<n_blocks, n_threads, 0,
0387                                details::get_stream(stream())>>>(payload);
0388   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0389 }
0390
0391 void gbts_seeding_algorithm::gbts_reset_edge_bids_kernel(
0392     const device::gbts_reset_edge_bids_payload& payload) const {
0393   const unsigned int n_threads = 128;
0394   const unsigned int n_blocks = 1 + (payload.nProps - 1) / n_threads;
0395   kernels::gbts_reset_edge_bids<<<n_blocks, n_threads, 0,
0396                                   details::get_stream(stream())>>>(payload);
0397   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0398 }
0399
0400 void gbts_seeding_algorithm::gbts_rebid_seeds_for_edges_kernel(
0401     const device::gbts_rebid_seeds_for_edges_payload& payload) const {
0402   const unsigned int n_threads = 128;
0403   const unsigned int n_blocks = 1 + (payload.nProps - 1) / n_threads;
0404   kernels::gbts_rebid_seeds_for_edges<<<n_blocks, n_threads, 0,
0405                                         details::get_stream(stream())>>>(
0406       payload);
0407   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0408 }
0409
0410 void gbts_seeding_algorithm::gbts_bid_seeds_for_hits_kernel(
0411     const device::gbts_bid_seeds_for_hits_payload& payload) const {
0412   const unsigned int n_threads = 128;
0413   const unsigned int n_blocks = 1 + (payload.nProps - 1) / n_threads;
0414   kernels::gbts_bid_seeds_for_hits<<<n_blocks, n_threads, 0,
0415                                      details::get_stream(stream())>>>(payload);
0416   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0417 }
0418
0419 void gbts_seeding_algorithm::gbts_convert_seeds_kernel(
0420     const device::gbts_convert_seeds_payload& payload) const {
0421   const unsigned int n_threads = 128;
0422   const unsigned int n_blocks = 1 + (payload.nProps - 1) / n_threads;
0423   kernels::gbts_convert_seeds<<<n_blocks, n_threads, 0,
0424                                 details::get_stream(stream())>>>(payload);
0425   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());  //
0426 }
0427
0428 }  // namespace traccc::cuda