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

[Artem Belevich via Phabricator via cfe-commits]

tra created this revision.
Herald added a subscriber: sanjoy.

https://reviews.llvm.org/D40871

Files:
  clang/lib/Headers/__clang_cuda_intrinsics.h


Index: clang/lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- clang/lib/Headers/__clang_cuda_intrinsics.h
+++ clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -135,6 +135,24 @@
     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;                                                           \


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D40871.125648.patch
Type: text/x-patch
Size: 2204 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20171206/97043e9f/attachment.bin>