src/utils/barrier.hpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2023-2025 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/device/concepts/barrier.hpp"
0012
0013 // HIP include(s).
0014 #include <hip/hip_runtime.h>
0015
0016 namespace traccc::hip {
0017
0018 /// A HIP barrier implementation
0019 struct barrier {
0020   __device__ inline void blockBarrier() const { __syncthreads(); }
0021
0022   __device__ inline bool blockAnd(bool predicate) const {
0023     return __syncthreads_and(predicate);
0024   }
0025
0026   __device__ inline bool blockOr(bool predicate) const {
0027     return __syncthreads_or(predicate);
0028   }
0029
0030   __device__ inline int blockCount(bool predicate) const {
0031     return __syncthreads_count(predicate);
0032   }
0033
0034 };  // struct barrier
0035
0036 // Ensure that the barrier concept is satisfied
0037 static_assert(traccc::device::concepts::barrier<barrier>);
0038
0039 }  // namespace traccc::hip