[NVPTX,CUDA] Added llvm.nvvm.fns intrinsic and matching __nvvm_fns builtin in clang.

Differential Revision: https://reviews.llvm.org/D40872

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@319909 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsNVPTX.def b/include/clang/Basic/BuiltinsNVPTX.def
index b596793c9c12854e22dcc1aa4a4f51c98c8eec56..7bab73a3b110cdcd081d9e87400e48daa5e18a4e 100644 (file)
--- a/include/clang/Basic/BuiltinsNVPTX.def
+++ b/include/clang/Basic/BuiltinsNVPTX.def
@@ -371,6 +371,9 @@ BUILTIN(__nvvm_bitcast_i2f, "fi", "")
 BUILTIN(__nvvm_bitcast_ll2d, "dLLi", "")
 BUILTIN(__nvvm_bitcast_d2ll, "LLid", "")
 
+// FNS
+TARGET_BUILTIN(__nvvm_fns, "UiUiUii", "n", "ptx60")
+
 // Sync
 
 BUILTIN(__syncthreads, "v", "")

diff --git a/lib/Headers/__clang_cuda_intrinsics.h b/lib/Headers/__clang_cuda_intrinsics.h
index 3f14b4f2dded3d998d1fd328c6fdd6b0e7bc1782..02d68a2e618ed655e1e8cd455263e4b47b0e1868 100644 (file)
--- a/lib/Headers/__clang_cuda_intrinsics.h
+++ b/lib/Headers/__clang_cuda_intrinsics.h
@@ -206,6 +206,10 @@ inline __device__ unsigned int __ballot_sync(unsigned int mask, int pred) {
 
 inline __device__ unsigned int __activemask() { return __nvvm_vote_ballot(1); }
 
+inline __device__ unsigned int __fns(unsigned mask, unsigned base, int offset) {
+  return __nvvm_fns(mask, base, offset);
+}
+
 #endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300
 
 // Define __match* builtins CUDA-9 headers expect to see.