src/utils/barrier.hpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2023-2024 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 #pragma once
0009
0010 namespace traccc::cuda {
0011
0012 struct barrier {
0013   __device__ inline void blockBarrier() const { __syncthreads(); }
0014
0015   __device__ inline bool blockAnd(bool predicate) const {
0016     return __syncthreads_and(predicate);
0017   }
0018
0019   __device__ inline bool blockOr(bool predicate) const {
0020     return __syncthreads_or(predicate);
0021   }
0022
0023   __device__ inline int blockCount(bool predicate) const {
0024     return __syncthreads_count(predicate);
0025   }
0026 };
0027
0028 }  // namespace traccc::cuda