diff options
Diffstat (limited to 'libc/src/__support/GPU/nvptx/utils.h')
| -rw-r--r-- | libc/src/__support/GPU/nvptx/utils.h | 7 |
1 files changed, 7 insertions, 0 deletions
diff --git a/libc/src/__support/GPU/nvptx/utils.h b/libc/src/__support/GPU/nvptx/utils.h index 443b8c72fc85..9f20edf16c62 100644 --- a/libc/src/__support/GPU/nvptx/utils.h +++ b/libc/src/__support/GPU/nvptx/utils.h @@ -118,6 +118,13 @@ LIBC_INLINE uint32_t get_lane_size() { return LANE_SIZE; } #endif } +[[clang::convergent]] LIBC_INLINE uint64_t ballot(uint64_t lane_mask, bool x) { +#if __CUDA_ARCH__ >= 600 + return __nvvm_vote_ballot_sync(lane_mask, x); +#else + return lane_mask & __nvvm_vote_ballot(x); +#endif +} /// Waits for all the threads in the block to converge and issues a fence. [[clang::convergent]] LIBC_INLINE void sync_threads() { __syncthreads(); } |