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

lib/Headers/__clang_cuda_intrinsics.h

inline __device__ long long __FnName(unsigned int __mask, long long __val,   \
     return __ret;                                                              \
   }                                                                            \
   inline __device__ unsigned long long __FnName(                               \
       unsigned int __mask, unsigned long long __val, int __offset,             \
       int __width = warpSize) {                                                \
     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;                                                           \
     _Static_assert(sizeof(__tmp) == sizeof(__val));                            \
     memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp = ::__FnName(__mask, __tmp, __offset, __width);                      \
     double __ret;                                                              \
     memcpy(&__ret, &__tmp, sizeof(__ret));                                     \