kernels/specializations/find_tracks_src.cuh

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2023-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 #pragma once
0009
0010 // Project include(s).
0011 #include "traccc/finding/device/find_tracks.hpp"
0012
0013 // Local include(s).
0014 #include "../../../utils/barrier.hpp"
0015 #include "../../../utils/hints.hpp"
0016 #include "../../../utils/thread_id.hpp"
0017
0018 // System include(s).
0019 #include <utility>
0020
0021 namespace traccc::cuda {
0022 namespace kernels {
0023
0024 template <typename detector_t>
0025 __global__ void find_tracks(
0026     const __grid_constant__ finding_config cfg,
0027     const __grid_constant__ typename detector_t::const_view_type det,
0028     const __grid_constant__ device::find_tracks_payload payload) {
0029   __shared__ unsigned int shared_num_out_params;
0030   __shared__ unsigned int shared_candidates_size;
0031   extern __shared__ unsigned long long int s[];
0032   unsigned long long int* shared_insertion_mutex = s;
0033   std::pair<unsigned int, unsigned int>* shared_candidates =
0034       reinterpret_cast<std::pair<unsigned int, unsigned int>*>(
0035           &shared_insertion_mutex[blockDim.x]);
0036
0037   cuda::barrier barrier;
0038   details::thread_id1 thread_id;
0039
0040   device::find_tracks<detector_t>(
0041       thread_id, barrier, cfg, det, payload,
0042       device::find_tracks_shared_payload{
0043           .shared_num_out_params = shared_num_out_params,
0044           .shared_insertion_mutex = shared_insertion_mutex,
0045           .shared_candidates = shared_candidates,
0046           .shared_candidates_size = shared_candidates_size});
0047 }
0048
0049 }  // namespace kernels
0050
0051 template <typename detector_t>
0052 void find_tracks(const dim3& grid_size, const dim3& block_size,
0053                  std::size_t shared_mem_size, const cudaStream_t& stream,
0054                  const finding_config& cfg,
0055                  const typename detector_t::const_view_type& det,
0056                  const device::find_tracks_payload& payload) {
0057   kernels::find_tracks<detector_t>
0058       <<<grid_size, block_size, shared_mem_size, stream>>>(cfg, det, payload);
0059 }
0060 }  // namespace traccc::cuda