From: Artem Belevich <tra@google.com>
Date: Wed, 6 Dec 2017 17:40:35 +0000 (+0000)
Subject: [CUDA] Added overloads for '[unsigned] long' variants of shfl builtins.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=7540bcca41e7b1c551b6ac2a2bace497760e0983;p=clang

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

Differential Revision: https://reviews.llvm.org/D40871

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@319908 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/Headers/__clang_cuda_intrinsics.h b/lib/Headers/__clang_cuda_intrinsics.h
index bc5b876577..3f14b4f2dd 100644
--- a/lib/Headers/__clang_cuda_intrinsics.h
+++ b/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;                                                           \