Index: cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h +++ cfe/trunk/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( \ - ::__FnName(static_cast(__in), __offset, __width)); \ + ::__FnName(static_cast(__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( \ - ::__FnName(static_cast(__in), __offset, __width)); \ + unsigned long long __val, int __offset, int __width = warpSize) { \ + return static_cast(::__FnName( \ + static_cast(__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);