This revision was automatically updated to reflect the committed changes.
Closed by commit rC319908: [CUDA] Added overloads for '[unsigned] long'
variants of shfl builtins. (authored by tra).
Changed prior to commit:
https://reviews.llvm.org/D40871?vs=125648&id=125756#toc
Repository:
rC Clang
https://reviews.llvm.org/D40871
Files:
lib/Headers/__clang_cuda_intrinsics.h
Index: lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- lib/Headers/__clang_cuda_intrinsics.h
+++ 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;
\
Index: lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- lib/Headers/__clang_cuda_intrinsics.h
+++ 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; \
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
