This is an archive of the discontinued LLVM Phabricator instance.

Differential D28325

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

Authored by jlebar on Jan 4 2017, 2:02 PM.

Details

Reviewers
tra
Commits
rGb8f7a3b8b17b: [CUDA] Rename keywords used in macro so they don't conflict with MSVC.
rC291138: [CUDA] Rename keywords used in macro so they don't conflict with MSVC.
rL291138: [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.

Diff Detail

Event Timeline

jlebar updated this revision to Diff 83133.Jan 4 2017, 2:02 PM
jlebar retitled this revision from to [CUDA] Rename keywords used in macro so they don't conflict with MSVC..
jlebar updated this object.
jlebar added a reviewer: tra.
jlebar added a subscriber: cfe-commits.
jlebar added a child revision: D28326: [Docs] Update docs to indicate that CUDA compilation is supported on Windows..Jan 4 2017, 2:05 PM
tra accepted this revision.Jan 4 2017, 5:13 PM
tra edited edge metadata.
This revision is now accepted and ready to land.Jan 4 2017, 5:13 PM
Closed by commit rL291138: [CUDA] Rename keywords used in macro so they don't conflict with MSVC. (authored by jlebar). · Explain WhyJan 5 2017, 9:05 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
clang/
lib/
Headers/
42 lines

Diff 83133

clang/lib/Headers/__clang_cuda_intrinsics.h

Show All 29 Lines
#define __SM_30_INTRINSICS_H__ #define __SM_30_INTRINSICS_H__
#define __SM_30_INTRINSICS_HPP__ #define __SM_30_INTRINSICS_HPP__
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300
#pragma push_macro("__MAKE_SHUFFLES") #pragma push_macro("__MAKE_SHUFFLES")
#define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask) \ #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) { \ int __width = warpSize) { \
return __IntIntrinsic(__in, __offset, \ return __IntIntrinsic(__val, __offset, \
((warpSize - __width) << 8) | (__Mask)); \ ((warpSize - __width) << 8) | (__Mask)); \
} \ } \
inline __device__ float __FnName(float __in, int __offset, \ inline __device__ float __FnName(float __val, int __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return __FloatIntrinsic(__in, __offset, \ return __FloatIntrinsic(__val, __offset, \
((warpSize - __width) << 8) | (__Mask)); \ ((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) { \ int __width = warpSize) { \
return static_cast<unsigned int>( \ 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) { \ int __width = warpSize) { \
struct __Bits { \ struct __Bits { \
int __a, __b; \ int __a, __b; \
}; \ }; \
_Static_assert(sizeof(__in) == sizeof(__Bits)); \ _Static_assert(sizeof(__val) == sizeof(__Bits)); \
_Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \ _Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \
__Bits __tmp; \ __Bits __tmp; \
memcpy(&__in, &__tmp, sizeof(__in)); \ memcpy(&__val, &__tmp, sizeof(__val)); \
__tmp.__a = ::__FnName(__tmp.__a, __offset, __width); \ __tmp.__a = ::__FnName(__tmp.__a, __offset, __width); \
__tmp.__b = ::__FnName(__tmp.__b, __offset, __width); \ __tmp.__b = ::__FnName(__tmp.__b, __offset, __width); \
long long __out; \ long long __ret; \
memcpy(&__out, &__tmp, sizeof(__tmp)); \ memcpy(&__ret, &__tmp, sizeof(__tmp)); \
return __out; \ return __ret; \
} \ } \
inline __device__ unsigned long long __FnName( \ inline __device__ unsigned long long __FnName( \
unsigned long long __in, int __offset, int __width = warpSize) { \ unsigned long long __val, int __offset, int __width = warpSize) { \
return static_cast<unsigned long long>( \ return static_cast<unsigned long long>(::__FnName( \
::__FnName(static_cast<unsigned long long>(__in), __offset, __width)); \ 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) { \ int __width = warpSize) { \
long long __tmp; \ long long __tmp; \
_Static_assert(sizeof(__tmp) == sizeof(__in)); \ _Static_assert(sizeof(__tmp) == sizeof(__val)); \
memcpy(&__tmp, &__in, sizeof(__in)); \ memcpy(&__tmp, &__val, sizeof(__val)); \
__tmp = ::__FnName(__tmp, __offset, __width); \ __tmp = ::__FnName(__tmp, __offset, __width); \
double __out; \ double __ret; \
memcpy(&__out, &__tmp, sizeof(__out)); \ memcpy(&__ret, &__tmp, sizeof(__ret)); \
return __out; \ return __ret; \
} }
__MAKE_SHUFFLES(__shfl, __nvvm_shfl_idx_i32, __nvvm_shfl_idx_f32, 0x1f); __MAKE_SHUFFLES(__shfl, __nvvm_shfl_idx_i32, __nvvm_shfl_idx_f32, 0x1f);
// We use 0 rather than 31 as our mask, because shfl.up applies to lanes >= // We use 0 rather than 31 as our mask, because shfl.up applies to lanes >=
// maxLane. // maxLane.
__MAKE_SHUFFLES(__shfl_up, __nvvm_shfl_up_i32, __nvvm_shfl_up_f32, 0); __MAKE_SHUFFLES(__shfl_up, __nvvm_shfl_up_i32, __nvvm_shfl_up_f32, 0);
__MAKE_SHUFFLES(__shfl_down, __nvvm_shfl_down_i32, __nvvm_shfl_down_f32, 0x1f); __MAKE_SHUFFLES(__shfl_down, __nvvm_shfl_down_i32, __nvvm_shfl_down_f32, 0x1f);
__MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f); __MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f);
▲ Show 20 Lines • Show All 233 Lines • Show Last 20 Lines