aboutsummaryrefslogtreecommitdiffstats
path: root/lib/Headers/__clang_cuda_intrinsics.h
diff options
context:
space:
mode:
Diffstat (limited to 'lib/Headers/__clang_cuda_intrinsics.h')
-rw-r--r--lib/Headers/__clang_cuda_intrinsics.h10
1 files changed, 9 insertions, 1 deletions
diff --git a/lib/Headers/__clang_cuda_intrinsics.h b/lib/Headers/__clang_cuda_intrinsics.h
index 2970d17f89ee..b67461a146fc 100644
--- a/lib/Headers/__clang_cuda_intrinsics.h
+++ b/lib/Headers/__clang_cuda_intrinsics.h
@@ -211,7 +211,15 @@ inline __device__ unsigned int __ballot_sync(unsigned int mask, int pred) {
return __nvvm_vote_ballot_sync(mask, pred);
}
-inline __device__ unsigned int __activemask() { return __nvvm_vote_ballot(1); }
+inline __device__ unsigned int __activemask() {
+#if CUDA_VERSION < 9020
+ return __nvvm_vote_ballot(1);
+#else
+ unsigned int mask;
+ asm volatile("activemask.b32 %0;" : "=r"(mask));
+ return mask;
+#endif
+}
inline __device__ unsigned int __fns(unsigned mask, unsigned base, int offset) {
return __nvvm_fns(mask, base, offset);