diff options
| author | Justin Lebar <jlebar@google.com> | 2017-01-05 16:54:11 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2017-01-05 16:54:11 +0000 |
| commit | a39e49e99a1a828db2e797035d2bfaa552e0d88b (patch) | |
| tree | 4741a91377384a66aff82e0ad44e28fbb6fac0ea /lib/Headers | |
| parent | a2b7c3d9ef33a137d8c253c79033cffcd535bc91 (diff) | |
[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
Diffstat (limited to 'lib/Headers')
| -rw-r--r-- | lib/Headers/__clang_cuda_intrinsics.h | 42 |
1 files changed, 21 insertions, 21 deletions
diff --git a/lib/Headers/__clang_cuda_intrinsics.h b/lib/Headers/__clang_cuda_intrinsics.h index 3df41fa290..b43ce21d0b 100644 --- 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); |