[libc-commits] [libc] c830c12 - [libc] Remove leftover target dependent intrinsic

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2024-02-13T12:44:26-06:00
New Revision: c830c1205dc164b645edb9c40cccbe768d5b337c

URL: https://github.com/llvm/llvm-project/commit/c830c1205dc164b645edb9c40cccbe768d5b337c
DIFF: https://github.com/llvm/llvm-project/commit/c830c1205dc164b645edb9c40cccbe768d5b337c.diff

LOG: [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.

Added: 
    

Modified: 
    libc/src/__support/GPU/nvptx/utils.h

Removed: 
    


################################################################################
diff  --git a/libc/src/__support/GPU/nvptx/utils.h b/libc/src/__support/GPU/nvptx/utils.h
index 22a46e87cfc05d..a92c8847b6ecdf 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(); }