Index: cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h +++ cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h @@ -135,6 +135,24 @@ return static_cast(::__FnName( \ __mask, static_cast(__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(::__FnName( \ + __mask, static_cast(__val), __offset, __width)); \ + } else if (sizeof(long) == sizeof(int)) { \ + return static_cast( \ + ::__FnName(__mask, static_cast(__val), __offset, __width)); \ + } \ + } \ + inline __device__ unsigned long __FnName(unsigned int __mask, \ + unsigned long __val, int __offset, \ + int __width = warpSize) { \ + return static_cast( \ + ::__FnName(__mask, static_cast(__val), __offset, __width)); \ + } \ inline __device__ double __FnName(unsigned int __mask, double __val, \ int __offset, int __width = warpSize) { \ long long __tmp; \