diff options
| author | Artem Belevich <tra@google.com> | 2017-12-06 17:40:35 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2017-12-06 17:40:35 +0000 |
| commit | 4631ef1e43703a70686f817960fdc0116c117e1e (patch) | |
| tree | 286cedea55742e869b582e2ca4f10316b9c2bda0 /clang/lib | |
| parent | ffaed72089f359320d92e07f9ed00de070264969 (diff) | |
| download | bcm5719-llvm-4631ef1e43703a70686f817960fdc0116c117e1e.tar.gz bcm5719-llvm-4631ef1e43703a70686f817960fdc0116c117e1e.zip | |
[CUDA] Added overloads for '[unsigned] long' variants of shfl builtins.
Differential Revision: https://reviews.llvm.org/D40871
llvm-svn: 319908
Diffstat (limited to 'clang/lib')
| -rw-r--r-- | clang/lib/Headers/__clang_cuda_intrinsics.h | 18 |
1 files changed, 18 insertions, 0 deletions
diff --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h index bc5b876577f..3f14b4f2dde 100644 --- a/clang/lib/Headers/__clang_cuda_intrinsics.h +++ b/clang/lib/Headers/__clang_cuda_intrinsics.h @@ -135,6 +135,24 @@ __MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f); return static_cast<unsigned long long>(::__FnName( \ __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, \ int __offset, int __width = warpSize) { \ long long __tmp; \ |
