[CUDA] Fix names of __nvvm_vote* intrinsics.

Also fixed a syntax error in activemask().

Differential Revision: https://reviews.llvm.org/D38188

llvm-svn: 314129

diff --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index 15a5cc4..0e2141a 100644 (file)
--- a/clang/lib/Headers/__clang_cuda_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -170,22 +170,22 @@ inline __device__ void __barrier_sync_count(unsigned int id,
 }
 
 inline __device__ int __all_sync(unsigned int mask, int pred) {
-  return __nvvm_vote_sync_all(mask, pred);
+  return __nvvm_vote_all_sync(mask, pred);
 }
 
 inline __device__ int __any_sync(unsigned int mask, int pred) {
-  return __nvvm_vote_sync_any(mask, pred);
+  return __nvvm_vote_any_sync(mask, pred);
 }
 
 inline __device__ int __uni_sync(unsigned int mask, int pred) {
-  return __nvvm_vote_sync_uni(mask, pred);
+  return __nvvm_vote_uni_sync(mask, pred);
 }
 
 inline __device__ unsigned int __ballot_sync(unsigned int mask, int pred) {
-  return __nvvm_vote_sync_ballot(mask, pred);
+  return __nvvm_vote_ballot_sync(mask, pred);
 }
 
-inline __device__ activemask() { return __nvvm_vote.ballot(1); }
+inline __device__ unsigned int activemask() { return __nvvm_vote_ballot(1); }
 
 #endif // __CUDA_VERSION >= 9000 && (!defined(__CUDA_ARCH__) ||
        // __CUDA_ARCH__ >= 300)