From 7540bcca41e7b1c551b6ac2a2bace497760e0983 Mon Sep 17 00:00:00 2001
From: Artem Belevich <tra@google.com>
Date: Wed, 6 Dec 2017 17:40:35 +0000
Subject: [PATCH] [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
---
lib/Headers/__clang_cuda_intrinsics.h | 18 ++++++++++++++++++
1 file changed, 18 insertions(+)
diff --git a/lib/Headers/__clang_cuda_intrinsics.h b/lib/Headers/__clang_cuda_intrinsics.h
index bc5b876577f..3f14b4f2dde 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; \
--
GitLab