File tree Expand file tree Collapse file tree 1 file changed +0
-8
lines changed
libc/src/__support/GPU/nvptx Expand file tree Collapse file tree 1 file changed +0
-8
lines changed Original file line number Diff line number Diff line change @@ -110,21 +110,13 @@ LIBC_INLINE uint32_t get_lane_size() { return 32; }
110
110
uint32_t x) {
111
111
uint32_t mask = static_cast <uint32_t >(lane_mask);
112
112
uint32_t id = __builtin_ffs (mask) - 1 ;
113
- #if __CUDA_ARCH__ >= 600
114
113
return __nvvm_shfl_sync_idx_i32 (mask, x, id, get_lane_size () - 1 );
115
- #else
116
- return __nvvm_shfl_idx_i32 (x, id, get_lane_size () - 1 );
117
- #endif
118
114
}
119
115
120
116
// / Returns a bitmask of threads in the current lane for which \p x is true.
121
117
[[clang::convergent]] LIBC_INLINE uint64_t ballot (uint64_t lane_mask, bool x) {
122
118
uint32_t mask = static_cast <uint32_t >(lane_mask);
123
- #if __CUDA_ARCH__ >= 600
124
119
return __nvvm_vote_ballot_sync (mask, x);
125
- #else
126
- return mask & __nvvm_vote_ballot (x);
127
- #endif
128
120
}
129
121
// / Waits for all the threads in the block to converge and issues a fence.
130
122
[[clang::convergent]] LIBC_INLINE void sync_threads () { __syncthreads (); }
You can’t perform that action at this time.
0 commit comments