diff options
author | Joseph Huber <huberjn@outlook.com> | 2024-02-13 12:43:44 -0600 |
---|---|---|
committer | Joseph Huber <huberjn@outlook.com> | 2024-02-13 12:44:26 -0600 |
commit | c830c1205dc164b645edb9c40cccbe768d5b337c (patch) | |
tree | 5f74b75c012d1188ea045736a69875b82198aaf1 | |
parent | 16140ff219b68f61fedf92df13019d89a4990a47 (diff) | |
download | llvm-c830c1205dc164b645edb9c40cccbe768d5b337c.zip llvm-c830c1205dc164b645edb9c40cccbe768d5b337c.tar.gz llvm-c830c1205dc164b645edb9c40cccbe768d5b337c.tar.bz2 |
[libc] Remove leftover target dependent intrinsic
Summary:
I forgot to remove these because I thought I did it already. This caused
the build to fail when actually linked.
-rw-r--r-- | libc/src/__support/GPU/nvptx/utils.h | 8 |
1 files changed, 0 insertions, 8 deletions
diff --git a/libc/src/__support/GPU/nvptx/utils.h b/libc/src/__support/GPU/nvptx/utils.h index 22a46e8..a92c884 100644 --- a/libc/src/__support/GPU/nvptx/utils.h +++ b/libc/src/__support/GPU/nvptx/utils.h @@ -110,21 +110,13 @@ LIBC_INLINE uint32_t get_lane_size() { return 32; } uint32_t x) { uint32_t mask = static_cast<uint32_t>(lane_mask); uint32_t id = __builtin_ffs(mask) - 1; -#if __CUDA_ARCH__ >= 600 return __nvvm_shfl_sync_idx_i32(mask, x, id, get_lane_size() - 1); -#else - return __nvvm_shfl_idx_i32(x, id, get_lane_size() - 1); -#endif } /// Returns a bitmask of threads in the current lane for which \p x is true. [[clang::convergent]] LIBC_INLINE uint64_t ballot(uint64_t lane_mask, bool x) { uint32_t mask = static_cast<uint32_t>(lane_mask); -#if __CUDA_ARCH__ >= 600 return __nvvm_vote_ballot_sync(mask, x); -#else - return 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(); } |