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

[Artem Belevich via cfe-commits]

Author: tra
Date: Wed Dec  6 09:50:05 2017
New Revision: 319909

URL: http://llvm.org/viewvc/llvm-project?rev=319909&view=rev
Log:
[NVPTX,CUDA] Added llvm.nvvm.fns intrinsic and matching __nvvm_fns builtin in clang.


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

Modified:
    cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
    cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h

Modified: cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def?rev=319909&r1=319908&r2=319909&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def Wed Dec  6 09:50:05 2017
@@ -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", "")

Modified: cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h?rev=319909&r1=319908&r2=319909&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h Wed Dec  6 09:50:05 2017
@@ -206,6 +206,10 @@ inline __device__ unsigned int __ballot_
 
 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.