aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorJoseph Huber <huberjn@outlook.com>2024-02-13 12:43:44 -0600
committerJoseph Huber <huberjn@outlook.com>2024-02-13 12:44:26 -0600
commitc830c1205dc164b645edb9c40cccbe768d5b337c (patch)
tree5f74b75c012d1188ea045736a69875b82198aaf1
parent16140ff219b68f61fedf92df13019d89a4990a47 (diff)
downloadllvm-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.h8
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(); }