Warning, /acts/Traccc/device/hip/src/gbts_seeding/gbts_seeding_algorithm.hip is written in an unsupported language. File is not indexed.
0001 // This file is part of the ACTS project.
0002 //
0003 // Copyright (C) 2016 CERN for the benefit of the ACTS project
0004 //
0005 // This Source Code Form is subject to the terms of the Mozilla Public
0006 // License, v. 2.0. If a copy of the MPL was not distributed with this
0007 // file, You can obtain one at https://mozilla.org/MPL/2.0/.
0008
0009 // Local include(s).
0010 #include "../utils/barrier.hpp"
0011 #include "../utils/hip_error_handling.hpp"
0012 #include "../utils/thread_id.hpp"
0013 #include "../utils/utils.hpp"
0014 #include "traccc/hip/gbts_seeding/gbts_seeding_algorithm.hpp"
0015
0016 // Project include(s).
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_build_edge_work_list.hpp"
0020 #include "traccc/gbts_seeding/device/gbts_compress_graph.hpp"
0021 #include "traccc/gbts_seeding/device/gbts_convert_seeds.hpp"
0022 #include "traccc/gbts_seeding/device/gbts_count_graph_edges.hpp"
0023 #include "traccc/gbts_seeding/device/gbts_count_paths.hpp"
0024 #include "traccc/gbts_seeding/device/gbts_fill_graph_edges.hpp"
0025 #include "traccc/gbts_seeding/device/gbts_fill_path_store.hpp"
0026 #include "traccc/gbts_seeding/device/gbts_find_minmax_radius.hpp"
0027 #include "traccc/gbts_seeding/device/gbts_finish_cca.hpp"
0028 #include "traccc/gbts_seeding/device/gbts_match_graph_edges.hpp"
0029 #include "traccc/gbts_seeding/device/gbts_run_cca_iteration.hpp"
0030 #include "traccc/gbts_seeding/device/gbts_sort_nodes.hpp"
0031 #include "traccc/gbts_seeding/gbts_types.hpp"
0032
0033 // VecMem include(s).
0034 #include <vecmem/containers/data/vector_view.hpp>
0035
0036 // (Roc)Thrust include(s).
0037 #include <thrust/execution_policy.h>
0038 #include <thrust/scan.h>
0039 #include <thrust/sort.h>
0040
0041 // System include(s).
0042 #include <algorithm>
0043 #include <memory_resource>
0044
0045 namespace traccc::hip {
0046
0047 namespace kernels {
0048
0049 using float4 = traccc::float4;
0050 using uint2 = traccc::uint2;
0051 using int2 = traccc::int2;
0052
0053 // ---------------------------------------------------------------------------
0054 // Stage 1 — nodes-making kernels
0055 // ---------------------------------------------------------------------------
0056
0057 /// HIP kernel for running @c traccc::device::gbts_bin_spacepoints
0058 __global__ void gbts_bin_spacepoints(
0059 const device::gbts_bin_spacepoints_payload payload) {
0060 device::gbts_bin_spacepoints(details::thread_id1{}, payload);
0061 }
0062
0063 /// HIP kernel for running @c traccc::device::gbts_sort_nodes
0064 __global__ void gbts_sort_nodes(const device::gbts_sort_nodes_payload payload) {
0065 device::gbts_sort_nodes(details::thread_id1{}, payload);
0066 }
0067
0068 /// HIP kernel for running @c traccc::device::gbts_find_minmax_radius
0069 __global__ void gbts_find_minmax_radius(
0070 const device::gbts_find_minmax_radius_payload payload) {
0071 device::gbts_find_minmax_radius(details::thread_id1{}, payload);
0072 }
0073
0074 // ---------------------------------------------------------------------------
0075 // Stage 2 — graph-making kernels
0076 // ---------------------------------------------------------------------------
0077
0078 /// HIP kernel for running @c traccc::device::gbts_build_edge_work_list
0079 __global__ void gbts_build_edge_work_list(
0080 const device::gbts_build_edge_work_list_payload payload) {
0081 __shared__ unsigned int scratch[device::gbts_build_edge_work_list_block_size];
0082 const traccc::hip::barrier barrier;
0083 device::gbts_build_edge_work_list(
0084 details::thread_id1{}, barrier, payload,
0085 {vecmem::data::vector_view<unsigned int>(
0086 device::gbts_build_edge_work_list_block_size, scratch)});
0087 }
0088
0089 /// HIP kernel for running @c traccc::device::gbts_count_graph_edges
0090 __global__ void gbts_count_graph_edges(
0091 const device::gbts_count_graph_edges_payload payload) {
0092 device::gbts_count_graph_edges(details::thread_id1{}, payload);
0093 }
0094
0095 /// HIP kernel for running @c traccc::device::gbts_fill_graph_edges
0096 __global__ void gbts_fill_graph_edges(
0097 const device::gbts_fill_graph_edges_payload payload) {
0098 device::gbts_fill_graph_edges(details::thread_id1{}, payload);
0099 }
0100
0101 /// HIP kernel for running @c traccc::device::gbts_match_graph_edges
0102 __global__ void gbts_match_graph_edges(
0103 const device::gbts_match_graph_edges_payload payload) {
0104 device::gbts_match_graph_edges(details::thread_id1{}, payload);
0105 }
0106
0107 /// HIP kernel for running @c traccc::device::gbts_compress_graph
0108 __global__ void gbts_compress_graph(
0109 const device::gbts_compress_graph_payload payload) {
0110 device::gbts_compress_graph(details::thread_id1{}, payload);
0111 }
0112
0113 // ---------------------------------------------------------------------------
0114 // Stage 3 — graph-processing kernels
0115 // ---------------------------------------------------------------------------
0116
0117 /// HIP kernel for running @c traccc::device::gbts_run_cca_iteration
0118 __global__ void gbts_run_cca_iteration(
0119 const device::gbts_run_cca_iteration_payload payload) {
0120 device::gbts_run_cca_iteration(details::thread_id1{}, payload);
0121 }
0122
0123 /// HIP kernel for running @c traccc::device::gbts_finish_cca
0124 __global__ void gbts_finish_cca(const device::gbts_finish_cca_payload payload) {
0125 device::gbts_finish_cca(details::thread_id1{}, payload);
0126 }
0127
0128 /// HIP kernel for running @c traccc::device::gbts_count_paths
0129 __global__ void gbts_count_paths(
0130 const device::gbts_count_paths_payload payload) {
0131 device::gbts_count_paths(details::thread_id1{}, payload);
0132 }
0133
0134 /// HIP kernel for running @c traccc::device::gbts_fill_path_store
0135 __global__ void gbts_fill_path_store(
0136 const device::gbts_fill_path_store_payload payload) {
0137 device::gbts_fill_path_store(details::thread_id1{}, payload);
0138 }
0139
0140 /// HIP kernel for running @c traccc::device::gbts_bid_seeds_for_hits
0141 __global__ void gbts_bid_seeds_for_hits(
0142 const device::gbts_bid_seeds_for_hits_payload payload) {
0143 device::gbts_bid_seeds_for_hits(details::thread_id1{}, payload);
0144 }
0145
0146 /// HIP kernel for running @c traccc::device::gbts_convert_seeds
0147 __global__ void gbts_convert_seeds(
0148 const device::gbts_convert_seeds_payload payload) {
0149 device::gbts_convert_seeds(details::thread_id1{}, payload);
0150 }
0151
0152 } // namespace kernels
0153
0154 // ===========================================================================
0155 // gbts_seeding_algorithm: kernel launchers
0156 // ===========================================================================
0157
0158 gbts_seeding_algorithm::gbts_seeding_algorithm(
0159 const gbts_seedfinder_config& cfg, const memory_resource& mr,
0160 const vecmem::copy& copy, const stream_wrapper& str,
0161 std::unique_ptr<const Logger> logger, await_function_type await_func)
0162 : device::gbts_seeding_algorithm(cfg, mr, copy, std::move(logger)),
0163 hip::algorithm_base{str, std::move(await_func)} {}
0164
0165 void gbts_seeding_algorithm::gbts_bin_spacepoints_kernel(
0166 const device::gbts_bin_spacepoints_payload& payload) const {
0167 const unsigned int n_threads = 128;
0168 const unsigned int n_blocks = 1 + (payload.nSp - 1) / n_threads;
0169 hipLaunchKernelGGL(kernels::gbts_bin_spacepoints, n_blocks, n_threads, 0,
0170 details::get_stream(stream()), payload);
0171 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0172 }
0173
0174 void gbts_seeding_algorithm::gbts_sort_nodes_kernel(
0175 const device::gbts_sort_nodes_payload& payload) const {
0176 // Order the nodes by their (eta bin, phi, spacepoint index bits) keys,
0177 // carrying the full spacepoint index along as the value.
0178 vecmem::device_vector<unsigned long long int> d_sort_keys(payload.sort_keys);
0179 vecmem::device_vector<unsigned int> d_sort_values(payload.sort_values);
0180 // Set up the Thrust execution policy. Unfortunately this has to be done in
0181 // a backend specific way. :-(
0182 #if defined(__HIP_PLATFORM_AMD__)
0183 auto policy =
0184 thrust::hip::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0185 .on(details::get_stream(stream()));
0186 #elif defined(__HIP_PLATFORM_NVIDIA__)
0187 auto policy =
0188 thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0189 .on(details::get_stream(stream()));
0190 #else
0191 #error "Unrecognised HIP platform."
0192 #endif
0193 thrust::sort_by_key(policy, d_sort_keys.begin(),
0194 d_sort_keys.begin() + static_cast<int>(payload.nSp),
0195 d_sort_values.begin());
0196
0197 const unsigned int n_threads = 256;
0198 const unsigned int n_blocks = 1 + (payload.nSp - 1) / n_threads;
0199 hipLaunchKernelGGL(kernels::gbts_sort_nodes, n_blocks, n_threads, 0,
0200 details::get_stream(stream()), payload);
0201 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0202 }
0203
0204 void gbts_seeding_algorithm::gbts_find_minmax_radius_kernel(
0205 const device::gbts_find_minmax_radius_payload& payload) const {
0206 const unsigned int n_threads = 128;
0207 const unsigned int n_blocks = 1 + (payload.nEtaBins - 1) / n_threads;
0208 hipLaunchKernelGGL(kernels::gbts_find_minmax_radius, n_blocks, n_threads, 0,
0209 details::get_stream(stream()), payload);
0210 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0211 }
0212
0213 void gbts_seeding_algorithm::gbts_build_edge_work_list_kernel(
0214 const device::gbts_build_edge_work_list_payload& payload) const {
0215 hipLaunchKernelGGL(kernels::gbts_build_edge_work_list, 1,
0216 device::gbts_build_edge_work_list_block_size, 0,
0217 details::get_stream(stream()), payload);
0218 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0219 }
0220
0221 void gbts_seeding_algorithm::gbts_count_graph_edges_kernel(
0222 const device::gbts_count_graph_edges_payload& payload) const {
0223 // One thread per inner node of a chunk. The blocks stride over the work
0224 // items.
0225 const unsigned int n_threads = device::gbts_consts::edge_chunk_size;
0226 const unsigned int n_blocks =
0227 std::min(payload.nWorkMax, device::gbts_count_graph_edges_max_blocks);
0228 hipLaunchKernelGGL(kernels::gbts_count_graph_edges, n_blocks, n_threads, 0,
0229 details::get_stream(stream()), payload);
0230 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0231
0232 // Turn the per-node counts into the edge buckets.
0233 vecmem::device_vector<unsigned int> d_num_outgoing_edges(
0234 payload.num_outgoing_edges);
0235 // Set up the Thrust execution policy. Unfortunately this has to be done in
0236 // a backend specific way. :-(
0237 #if defined(__HIP_PLATFORM_AMD__)
0238 auto policy =
0239 thrust::hip::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0240 .on(details::get_stream(stream()));
0241 #elif defined(__HIP_PLATFORM_NVIDIA__)
0242 auto policy =
0243 thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0244 .on(details::get_stream(stream()));
0245 #else
0246 #error "Unrecognised HIP platform."
0247 #endif
0248 thrust::inclusive_scan(policy, d_num_outgoing_edges.begin(),
0249 d_num_outgoing_edges.end(),
0250 d_num_outgoing_edges.begin());
0251 }
0252
0253 void gbts_seeding_algorithm::gbts_fill_graph_edges_kernel(
0254 const device::gbts_fill_graph_edges_payload& payload) const {
0255 const unsigned int n_threads = device::gbts_consts::edge_chunk_size;
0256 const unsigned int n_blocks =
0257 std::min(payload.nWorkMax, device::gbts_fill_graph_edges_max_blocks);
0258 hipLaunchKernelGGL(kernels::gbts_fill_graph_edges, n_blocks, n_threads, 0,
0259 details::get_stream(stream()), payload);
0260 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0261 }
0262
0263 void gbts_seeding_algorithm::gbts_match_graph_edges_kernel(
0264 const device::gbts_match_graph_edges_payload& payload) const {
0265 const unsigned int n_threads = 256;
0266 const unsigned int n_blocks = 1u + (payload.nEdgesMax - 1u) / n_threads;
0267 hipLaunchKernelGGL(kernels::gbts_match_graph_edges, n_blocks, n_threads, 0,
0268 details::get_stream(stream()), payload);
0269 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0270
0271 // Compact the kept edges with a prefix sum over their 0/1 flags.
0272 // Set up the Thrust execution policy. Unfortunately this has to be done in
0273 // a backend specific way. :-(
0274 #if defined(__HIP_PLATFORM_AMD__)
0275 auto policy =
0276 thrust::hip::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0277 .on(details::get_stream(stream()));
0278 #elif defined(__HIP_PLATFORM_NVIDIA__)
0279 auto policy =
0280 thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0281 .on(details::get_stream(stream()));
0282 #else
0283 #error "Unrecognised HIP platform."
0284 #endif
0285 thrust::inclusive_scan(policy, payload.kept.ptr(),
0286 payload.kept.ptr() + payload.nEdgesMax,
0287 payload.reIndexer.ptr());
0288 }
0289
0290 void gbts_seeding_algorithm::gbts_compress_graph_kernel(
0291 const device::gbts_compress_graph_payload& payload) const {
0292 const unsigned int n_threads = 256;
0293 const unsigned int n_blocks = 1u + (payload.nEdgesMax - 1u) / n_threads;
0294 hipLaunchKernelGGL(kernels::gbts_compress_graph, n_blocks, n_threads, 0,
0295 details::get_stream(stream()), payload);
0296 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0297 }
0298
0299 void gbts_seeding_algorithm::gbts_run_cca_iteration_kernel(
0300 const device::gbts_run_cca_iteration_payload& payload) const {
0301 const unsigned int n_threads = 256;
0302 const unsigned int n_blocks = 1 + (payload.nConnectedEdges - 1) / n_threads;
0303 hipLaunchKernelGGL(kernels::gbts_run_cca_iteration, n_blocks, n_threads, 0,
0304 details::get_stream(stream()), payload);
0305 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0306 }
0307
0308 void gbts_seeding_algorithm::gbts_finish_cca_kernel(
0309 const device::gbts_finish_cca_payload& payload) const {
0310 const unsigned int n_threads = 256;
0311 const unsigned int n_blocks = 1 + (payload.nConnectedEdges - 1) / n_threads;
0312 hipLaunchKernelGGL(kernels::gbts_finish_cca, n_blocks, n_threads, 0,
0313 details::get_stream(stream()), payload);
0314 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0315 }
0316
0317 void gbts_seeding_algorithm::gbts_count_paths_kernel(
0318 const device::gbts_count_paths_payload& payload) const {
0319 const unsigned int n_threads = 128;
0320 const unsigned int n_blocks = 1 + (payload.nConnectedEdges - 1) / n_threads;
0321 hipLaunchKernelGGL(kernels::gbts_count_paths, n_blocks, n_threads, 0,
0322 details::get_stream(stream()), payload);
0323 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0324
0325 // Path offsets of the path store.
0326 vecmem::device_vector<unsigned int> d_path_counts(payload.path_counts);
0327 // Set up the Thrust execution policy. Unfortunately this has to be done in
0328 // a backend specific way. :-(
0329 #if defined(__HIP_PLATFORM_AMD__)
0330 auto policy =
0331 thrust::hip::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0332 .on(details::get_stream(stream()));
0333 #elif defined(__HIP_PLATFORM_NVIDIA__)
0334 auto policy =
0335 thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&(mr().main)))
0336 .on(details::get_stream(stream()));
0337 #else
0338 #error "Unrecognised HIP platform."
0339 #endif
0340 thrust::inclusive_scan(policy, d_path_counts.begin(),
0341 d_path_counts.begin() + payload.nConnectedEdges,
0342 d_path_counts.begin());
0343 }
0344
0345 void gbts_seeding_algorithm::gbts_fill_path_store_kernel(
0346 const device::gbts_fill_path_store_payload& payload) const {
0347 const unsigned int n_threads = 128;
0348 const unsigned int n_blocks = 1 + (payload.nPathsGrid - 1) / n_threads;
0349 hipLaunchKernelGGL(kernels::gbts_fill_path_store, n_blocks, n_threads, 0,
0350 details::get_stream(stream()), payload);
0351 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0352 }
0353
0354 void gbts_seeding_algorithm::gbts_bid_seeds_for_hits_kernel(
0355 const device::gbts_bid_seeds_for_hits_payload& payload) const {
0356 const unsigned int n_threads = 128;
0357 const unsigned int n_blocks = 1 + (payload.nPathsGrid - 1) / n_threads;
0358 hipLaunchKernelGGL(kernels::gbts_bid_seeds_for_hits, n_blocks, n_threads, 0,
0359 details::get_stream(stream()), payload);
0360 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0361 }
0362
0363 void gbts_seeding_algorithm::gbts_convert_seeds_kernel(
0364 const device::gbts_convert_seeds_payload& payload) const {
0365 const unsigned int n_threads = 128;
0366 const unsigned int n_blocks = 1 + (payload.nPathsGrid - 1) / n_threads;
0367 hipLaunchKernelGGL(kernels::gbts_convert_seeds, n_blocks, n_threads, 0,
0368 details::get_stream(stream()), payload);
0369 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0370 }
0371
0372 void gbts_seeding_algorithm::synchronize() const {
0373 stream().synchronize();
0374 }
0375
0376 } // namespace traccc::hip
