[CUDA] Add appropriate host/device attribute to builtins.

The changes are part of attribute-based CUDA function overloading (D12453)
and as such are only enabled when it's in effect (-fcuda-target-overloads).

Differential Revision: http://reviews.llvm.org/D12122

llvm-svn: 248296

2 files changed, 42 insertions, 3 deletions

diff --git a/clang/test/SemaCUDA/builtins.cu b/clang/test/SemaCUDA/builtins.cu
new file mode 100644
index 00000000000..2c619b56339
--- /dev/null
+++ b/clang/test/SemaCUDA/builtins.cu
@@ -0,0 +1,36 @@
+// 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
diff --git a/clang/test/SemaCUDA/implicit-intrinsic.cu b/clang/test/SemaCUDA/implicit-intrinsic.cu
index 3d24aa719e5..0793d64b101 100644
--- a/clang/test/SemaCUDA/implicit-intrinsic.cu
+++ b/clang/test/SemaCUDA/implicit-intrinsic.cu
@@ -1,10 +1,13 @@
-// 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();
 }