src/gbts_seeding/gbts_seeding_algorithm.sycl

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