[CUDA] Fixed order of words in the names of shfl builtins.

[jlebar]

Naturally they're different orders in the PTX and CUDA. :)

clang/lib/Headers/__clang_cuda_intrinsics.h

inline __device__ double __FnName(unsigned int __mask, double __val,         \
     double __ret;                                                              \
     memcpy(&__ret, &__tmp, sizeof(__ret));                                     \
     return __ret;                                                              \
   }
 __MAKE_SYNC_SHUFFLES(__shfl_sync, __nvvm_shfl_sync_idx_i32,
                      __nvvm_shfl_sync_idx_f32, 0x1f);
 // We use 0 rather than 31 as our mask, because shfl.up applies to lanes >=
 // maxLane.
-__MAKE_SYNC_SHUFFLES(__shfl_sync_up, __nvvm_shfl_sync_up_i32,
+__MAKE_SYNC_SHUFFLES(__shfl_up_sync, __nvvm_shfl_sync_up_i32,
                      __nvvm_shfl_sync_up_f32, 0);
-__MAKE_SYNC_SHUFFLES(__shfl_sync_down, __nvvm_shfl_sync_down_i32,
+__MAKE_SYNC_SHUFFLES(__shfl_down_sync, __nvvm_shfl_sync_down_i32,
                      __nvvm_shfl_sync_down_f32, 0x1f);
-__MAKE_SYNC_SHUFFLES(__shfl_sync_xor, __nvvm_shfl_sync_bfly_i32,
+__MAKE_SYNC_SHUFFLES(__shfl_xor_sync, __nvvm_shfl_sync_bfly_i32,
                      __nvvm_shfl_sync_bfly_f32, 0x1f);
-
 #pragma pop_macro("__MAKE_SYNC_SHUFFLES")
 
 inline __device__ void __syncwarp(unsigned int mask = 0xffffffff) {
   return __nvvm_bar_warp_sync(mask);
 }
 
 inline __device__ void __barrier_sync(unsigned int id) {
   __nvvm_barrier_sync(id);