[CUDA] Fix names of __nvvm_vote* intrinsics.

[tra]

Summary

Also fixed a syntax error in activemask().

[jlebar]

Should we add tests to the test-suite? Or, are these already caught by the existing tests we have?

[tra]

In D38188#880318, @jlebar wrote:

Should we add tests to the test-suite? Or, are these already caught by the existing tests we have?

That's the plan. Once clang can compile CUDA headers, I'll add CUDA-9 specific tests to the testsuite and update the buildbot to compile/run tests with CUDA-9.

clang/lib/Headers/__clang_cuda_intrinsics.h

 }
 
 inline __device__ void __barrier_sync_count(unsigned int id,
                                             unsigned int count) {
   __nvvm_barrier_sync_cnt(id, count);
 }
 
 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)
 
 // sm_32 intrinsics: __ldg and __funnelshift_{l,lc,r,rc}.
 
 // Prevent the vanilla sm_32 intrinsics header from being included.
 #define __SM_32_INTRINSICS_H__