This is an archive of the discontinued LLVM Phabricator instance.

Differential D40871

[CUDA] Added overloads for '[unsigned] long' variants of shfl builtins.
ClosedPublic

Authored by tra on Dec 5 2017, 4:19 PM.

Details

Reviewers
jlebar
Commits
rG4631ef1e4370: [CUDA] Added overloads for '[unsigned] long' variants of shfl builtins.
rC319908: [CUDA] Added overloads for '[unsigned] long' variants of shfl builtins.
rL319908: [CUDA] Added overloads for '[unsigned] long' variants of shfl builtins.

Diff Detail

Event Timeline

tra created this revision.Dec 5 2017, 4:19 PM
Herald added a subscriber: sanjoy. · View Herald TranscriptDec 5 2017, 4:19 PM
jlebar accepted this revision.Dec 5 2017, 4:28 PM
This revision is now accepted and ready to land.Dec 5 2017, 4:28 PM
Closed by commit rL319908: [CUDA] Added overloads for '[unsigned] long' variants of shfl builtins. (authored by tra). · Explain WhyDec 6 2017, 9:41 AM
Closed by commit rC319908: [CUDA] Added overloads for '[unsigned] long' variants of shfl builtins. (authored by tra). · Explain Why
This revision was automatically updated to reflect the committed changes.
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
clang/
lib/
Headers/
18 lines

Diff 125648

clang/lib/Headers/__clang_cuda_intrinsics.h

Show First 20 Lines • Show All 129 Lines • ▼ Show 20 Lines inline __device__ long long __FnName(unsigned int __mask, long long __val, \
return __ret; \ return __ret; \
} \ } \
inline __device__ unsigned long long __FnName( \ inline __device__ unsigned long long __FnName( \
unsigned int __mask, unsigned long long __val, int __offset, \ unsigned int __mask, unsigned long long __val, int __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return static_cast<unsigned long long>(::__FnName( \ return static_cast<unsigned long long>(::__FnName( \
__mask, static_cast<unsigned long long>(__val), __offset, __width)); \ __mask, static_cast<unsigned long long>(__val), __offset, __width)); \
} \ } \
inline __device__ long __FnName(unsigned int __mask, long __val, \
int __offset, int __width = warpSize) { \
_Static_assert(sizeof(long) == sizeof(long long) || \
sizeof(long) == sizeof(int)); \
if (sizeof(long) == sizeof(long long)) { \
return static_cast<long>(::__FnName( \
__mask, static_cast<long long>(__val), __offset, __width)); \
} else if (sizeof(long) == sizeof(int)) { \
return static_cast<long>( \
::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \
} \
} \
inline __device__ unsigned long __FnName(unsigned int __mask, \
unsigned long __val, int __offset, \
int __width = warpSize) { \
return static_cast<unsigned long>( \
::__FnName(__mask, static_cast<long>(__val), __offset, __width)); \
} \
inline __device__ double __FnName(unsigned int __mask, double __val, \ inline __device__ double __FnName(unsigned int __mask, double __val, \
int __offset, int __width = warpSize) { \ int __offset, int __width = warpSize) { \
long long __tmp; \ long long __tmp; \
_Static_assert(sizeof(__tmp) == sizeof(__val)); \ _Static_assert(sizeof(__tmp) == sizeof(__val)); \
memcpy(&__tmp, &__val, sizeof(__val)); \ memcpy(&__tmp, &__val, sizeof(__val)); \
__tmp = ::__FnName(__mask, __tmp, __offset, __width); \ __tmp = ::__FnName(__mask, __tmp, __offset, __width); \
double __ret; \ double __ret; \
memcpy(&__ret, &__tmp, sizeof(__ret)); \ memcpy(&__ret, &__tmp, sizeof(__ret)); \
▲ Show 20 Lines • Show All 301 Lines • Show Last 20 Lines