From df0c467c1717d405a598100fb9de4c210d4e9c0c Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Wed, 6 Dec 2017 17:50:05 +0000 Subject: [PATCH] [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 --- include/clang/Basic/BuiltinsNVPTX.def | 3 +++ lib/Headers/__clang_cuda_intrinsics.h | 4 ++++ 2 files changed, 7 insertions(+) diff --git a/include/clang/Basic/BuiltinsNVPTX.def b/include/clang/Basic/BuiltinsNVPTX.def index b596793c9c..7bab73a3b1 100644 --- 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 3f14b4f2dd..02d68a2e61 100644 --- 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. -- 2.50.1