return getRecord(ID).Type;
}
+ /// \brief Return true if this function is a target-specific builtin
+ bool isTSBuiltin(unsigned ID) const {
+ return ID >= Builtin::FirstTSBuiltin;
+ }
+
/// \brief Return true if this function has no side effects and doesn't
/// read memory.
bool isConst(unsigned ID) const {
// Since the target specific builtins for each arch overlap, only check those
// of the arch we are compiling for.
- if (BuiltinID >= Builtin::FirstTSBuiltin) {
+ if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
switch (Context.getTargetInfo().getTriple().getArch()) {
case llvm::Triple::arm:
case llvm::Triple::armeb:
FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation()));
if (Context.BuiltinInfo.isConst(BuiltinID) && !FD->hasAttr<ConstAttr>())
FD->addAttr(ConstAttr::CreateImplicit(Context, FD->getLocation()));
+ if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads &&
+ Context.BuiltinInfo.isTSBuiltin(BuiltinID) &&
+ !FD->hasAttr<CUDADeviceAttr>() && !FD->hasAttr<CUDAHostAttr>()) {
+ // Target-specific builtins are assumed to be intended for use
+ // in this particular CUDA compilation mode and should have
+ // appropriate attribute set so we can enforce CUDA function
+ // call restrictions.
+ if (getLangOpts().CUDAIsDevice)
+ FD->addAttr(CUDADeviceAttr::CreateImplicit(Context, FD->getLocation()));
+ else
+ FD->addAttr(CUDAHostAttr::CreateImplicit(Context, FD->getLocation()));
+ }
}
IdentifierInfo *Name = FD->getIdentifier();
--- /dev/null
+// Tests that target-specific builtins have appropriate host/device
+// attributes and that CUDA call restrictions are enforced. Also
+// verify that non-target builtins can be used from both host and
+// device functions.
+//
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown \
+// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
+// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
+
+
+#ifdef __CUDA_ARCH__
+// Device-side builtins are not allowed to be called from host functions.
+void hf() {
+ int x = __builtin_ptx_read_tid_x(); // expected-note {{'__builtin_ptx_read_tid_x' declared here}}
+ // expected-error@-1 {{reference to __device__ function '__builtin_ptx_read_tid_x' in __host__ function}}
+ x = __builtin_abs(1);
+}
+__attribute__((device)) void df() {
+ int x = __builtin_ptx_read_tid_x();
+ x = __builtin_abs(1);
+}
+#else
+// Host-side builtins are not allowed to be called from device functions.
+__attribute__((device)) void df() {
+ int x = __builtin_ia32_rdtsc(); // expected-note {{'__builtin_ia32_rdtsc' declared here}}
+ // expected-error@-1 {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
+ x = __builtin_abs(1);
+}
+void hf() {
+ int x = __builtin_ia32_rdtsc();
+ x = __builtin_abs(1);
+}
+#endif
-// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
+// RUN: -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
+// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
#include "Inputs/cuda.h"
// expected-no-diagnostics
__device__ void __threadfence_system() {
- // This shouldn't produce an error, since __nvvm_membar_sys is inferred to
- // be __host__ __device__ and thus callable from device code.
+ // This shouldn't produce an error, since __nvvm_membar_sys should be
+ // __device__ and thus callable from device code.
__nvvm_membar_sys();
}