[CUDA] Rename keywords used in macro so they don't conflict with MSVC.

Summary:
MSVC seems to use "__in" and "__out" for its own purposes, so we have to
pick different names in this macro.

Reviewers: tra

Subscribers: cfe-commits

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@291138 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Headers/__clang_cuda_intrinsics.h b/lib/Headers/__clang_cuda_intrinsics.h
index 3df41fa290d3d27e8fed3bcea4dad9adb4884ecd..b43ce21d0bb32bd6d7dfcd46a2420570433400d4 100644 (file)
--- a/lib/Headers/__clang_cuda_intrinsics.h
+++ b/lib/Headers/__clang_cuda_intrinsics.h
@@ -35,50 +35,50 @@
 
 #pragma push_macro("__MAKE_SHUFFLES")
 #define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask)    \
-  inline __device__ int __FnName(int __in, int __offset,                       \
+  inline __device__ int __FnName(int __val, int __offset,                      \
                                  int __width = warpSize) {                     \
-    return __IntIntrinsic(__in, __offset,                                      \
+    return __IntIntrinsic(__val, __offset,                                     \
                           ((warpSize - __width) << 8) | (__Mask));             \
   }                                                                            \
-  inline __device__ float __FnName(float __in, int __offset,                   \
+  inline __device__ float __FnName(float __val, int __offset,                  \
                                    int __width = warpSize) {                   \
-    return __FloatIntrinsic(__in, __offset,                                    \
+    return __FloatIntrinsic(__val, __offset,                                   \
                             ((warpSize - __width) << 8) | (__Mask));           \
   }                                                                            \
-  inline __device__ unsigned int __FnName(unsigned int __in, int __offset,     \
+  inline __device__ unsigned int __FnName(unsigned int __val, int __offset,    \
                                           int __width = warpSize) {            \
     return static_cast<unsigned int>(                                          \
-        ::__FnName(static_cast<int>(__in), __offset, __width));                \
+        ::__FnName(static_cast<int>(__val), __offset, __width));               \
   }                                                                            \
-  inline __device__ long long __FnName(long long __in, int __offset,           \
+  inline __device__ long long __FnName(long long __val, int __offset,          \
                                        int __width = warpSize) {               \
     struct __Bits {                                                            \
       int __a, __b;                                                            \
     };                                                                         \
-    _Static_assert(sizeof(__in) == sizeof(__Bits));                            \
+    _Static_assert(sizeof(__val) == sizeof(__Bits));                           \
     _Static_assert(sizeof(__Bits) == 2 * sizeof(int));                         \
     __Bits __tmp;                                                              \
-    memcpy(&__in, &__tmp, sizeof(__in));                                       \
+    memcpy(&__val, &__tmp, sizeof(__val));                                     \
     __tmp.__a = ::__FnName(__tmp.__a, __offset, __width);                      \
     __tmp.__b = ::__FnName(__tmp.__b, __offset, __width);                      \
-    long long __out;                                                           \
-    memcpy(&__out, &__tmp, sizeof(__tmp));                                     \
-    return __out;                                                              \
+    long long __ret;                                                           \
+    memcpy(&__ret, &__tmp, sizeof(__tmp));                                     \
+    return __ret;                                                              \
   }                                                                            \
   inline __device__ unsigned long long __FnName(                               \
-      unsigned long long __in, int __offset, int __width = warpSize) {         \
-    return static_cast<unsigned long long>(                                    \
-        ::__FnName(static_cast<unsigned long long>(__in), __offset, __width)); \
+      unsigned long long __val, int __offset, int __width = warpSize) {        \
+    return static_cast<unsigned long long>(::__FnName(                         \
+        static_cast<unsigned long long>(__val), __offset, __width));           \
   }                                                                            \
-  inline __device__ double __FnName(double __in, int __offset,                 \
+  inline __device__ double __FnName(double __val, int __offset,                \
                                     int __width = warpSize) {                  \
     long long __tmp;                                                           \
-    _Static_assert(sizeof(__tmp) == sizeof(__in));                             \
-    memcpy(&__tmp, &__in, sizeof(__in));                                       \
+    _Static_assert(sizeof(__tmp) == sizeof(__val));                            \
+    memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp = ::__FnName(__tmp, __offset, __width);                              \
-    double __out;                                                              \
-    memcpy(&__out, &__tmp, sizeof(__out));                                     \
-    return __out;                                                              \
+    double __ret;                                                              \
+    memcpy(&__ret, &__tmp, sizeof(__ret));                                     \
+    return __ret;                                                              \
   }
 
 __MAKE_SHUFFLES(__shfl, __nvvm_shfl_idx_i32, __nvvm_shfl_idx_f32, 0x1f);