6 files changed, 417 insertions, 9 deletions
diff --git a/test/SemaCUDA/asm-constraints-mixed.cu b/test/SemaCUDA/asm-constraints-mixed.cu
index a4ac9c65c99f..a3b1e1a08c51 100644
--- a/test/SemaCUDA/asm-constraints-mixed.cu
+++ b/test/SemaCUDA/asm-constraints-mixed.cu
@@ -1,15 +1,39 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify %s
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
-// expected-no-diagnostics
+
+__attribute__((device)) register long global_dev_reg asm("r0");
+__attribute__((device)) register long
+    global_dev_hreg asm("rsp"); // device-side error
+
+register long global_host_reg asm("rsp");
+register long global_host_dreg asm("r0"); // host-side error
 
 __attribute__((device)) void df() {
+  register long local_dev_reg asm("r0");
+  register long local_host_reg asm("rsp"); // device-side error
   short h;
   // asm with PTX constraints. Some of them are PTX-specific.
-  __asm__("dont care" : "=h"(h): "f"(0.0), "d"(0.0), "h"(0), "r"(0), "l"(0));
+  __asm__("dont care" : "=h"(h) : "f"(0.0), "d"(0.0), "h"(0), "r"(0), "l"(0));
 }
 
 void hf() {
+  register long local_dev_reg asm("r0"); // host-side error
+  register long local_host_reg asm("rsp");
   int a;
-  // Asm with x86 constraints that are not supported by PTX.
-  __asm__("dont care" : "=a"(a): "a"(0), "b"(0), "c"(0));
+  // Asm with x86 constraints and registers that are not supported by PTX.
+  __asm__("dont care" : "=a"(a) : "a"(0), "b"(0), "c"(0) : "flags");
 }
+
+// Check errors in named register variables.
+// We should only see errors relevant to current compilation mode.
+#if defined(__CUDA_ARCH__)
+// Device-side compilation:
+// expected-error@8 {{unknown register name 'rsp' in asm}}
+// expected-error@15 {{unknown register name 'rsp' in asm}}
+#else
+// Host-side compilation:
+// expected-error@11 {{unknown register name 'r0' in asm}}
+// expected-error@22 {{unknown register name 'r0' in asm}}
+#endif
diff --git a/test/SemaCUDA/attributes.cu b/test/SemaCUDA/attributes.cu
new file mode 100644
index 000000000000..ce4dc925a3f3
--- /dev/null
+++ b/test/SemaCUDA/attributes.cu
@@ -0,0 +1,33 @@
+// Tests handling of CUDA attributes.
+//
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// Now pretend that we're compiling a C file. There should be warnings.
+// RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s
+
+#if defined(EXPECT_WARNINGS)
+// expected-warning@+12 {{'device' attribute ignored}}
+// expected-warning@+12 {{'global' attribute ignored}}
+// expected-warning@+12 {{'constant' attribute ignored}}
+// expected-warning@+12 {{'shared' attribute ignored}}
+// expected-warning@+12 {{'host' attribute ignored}}
+//
+// NOTE: IgnoredAttr in clang which is used for the rest of
+// attributes ignores LangOpts, so there are no warnings.
+#else
+// expected-no-diagnostics
+#endif
+
+__attribute__((device)) void f_device();
+__attribute__((global)) void f_global();
+__attribute__((constant)) int* g_constant;
+__attribute__((shared)) float *g_shared;
+__attribute__((host)) void f_host();
+__attribute__((device_builtin)) void f_device_builtin();
+typedef __attribute__((device_builtin)) const void *t_device_builtin;
+enum __attribute__((device_builtin)) e_device_builtin {E};
+__attribute__((device_builtin)) int v_device_builtin;
+__attribute__((cudart_builtin)) void f_cudart_builtin();
+__attribute__((nv_weak)) void f_nv_weak();
+__attribute__((device_builtin_surface_type)) unsigned long long surface_var;
+__attribute__((device_builtin_texture_type)) unsigned long long texture_var;
diff --git a/test/SemaCUDA/builtins.cu b/test/SemaCUDA/builtins.cu
new file mode 100644
index 000000000000..32b575862cfe
--- /dev/null
+++ b/test/SemaCUDA/builtins.cu
@@ -0,0 +1,31 @@
+// Tests that host and target builtins can be used in the same TU,
+// 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:     -aux-triple nvptx64-unknown-cuda \
+// RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
+// RUN:     -aux-triple x86_64-unknown-unknown \
+// RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
+
+#if !(defined(__amd64__) && defined(__PTX__))
+#error "Expected to see preprocessor macros from both sides of compilation."
+#endif
+
+void hf() {
+  int x = __builtin_ia32_rdtsc();
+  int y = __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();
+  int y = __builtin_ia32_rdtsc(); // expected-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
+                                  // expected-note@20 {{'__builtin_ia32_rdtsc' declared here}}
+  x = __builtin_abs(1);
+}
diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu
new file mode 100644
index 000000000000..bd3fb508bfab
--- /dev/null
+++ b/test/SemaCUDA/function-overload.cu
@@ -0,0 +1,317 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we handle target overloads correctly.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN:    -fsyntax-only -fcuda-target-overloads -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
+// RUN:    -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s
+
+// Check target overloads handling with disabled call target checks.
+// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s
+// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
+// RUN:    -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+typedef void (*gp_t)(void);
+
+// Host and unattributed functions can't be overloaded
+__host__ int hh(void) { return 1; } // expected-note {{previous definition is here}}
+int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}}
+
+// H/D overloading is OK
+__host__ int dh(void) { return 2; }
+__device__ int dh(void) { return 2; }
+
+// H/HD and D/HD are not allowed
+__host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}}
+__host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}}
+
+__host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}}
+__host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+__host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}}
+__device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}}
+
+__device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}}
+__host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+// Same tests for extern "C" functions
+extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}}
+extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}}
+
+// H/D overloading is OK
+extern "C" __device__ int cdh(void) {return 10;}
+extern "C" __host__ int cdh(void) {return 11;}
+
+// H/HD and D/HD overloading is not allowed.
+extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}}
+extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}}
+
+extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}}
+extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+// Helper functions to verify calling restrictions.
+__device__ int d(void) { return 8; }
+__host__ int h(void) { return 9; }
+__global__ void g(void) {}
+extern "C" __device__ int cd(void) {return 10;}
+extern "C" __host__ int ch(void) {return 11;}
+
+__host__ void hostf(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __device__ function 'd' in __host__ function}}
+  // expected-note@65 {{'d' declared here}}
+  // expected-error@-4 {{reference to __device__ function 'cd' in __host__ function}}
+  // expected-note@68 {{'cd' declared here}}
+#endif
+  fp_t hp = h;
+  fp_t chp = ch;
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g;
+
+  d();
+  cd();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'd'}}
+  // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}}
+  // expected-error@-4 {{no matching function for call to 'cd'}}
+  // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}}
+#endif
+  h();
+  ch();
+  dh();
+  cdh();
+  g(); // expected-error {{call to global function g not configured}}
+  g<<<0,0>>>();
+}
+
+
+__device__ void devicef(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __host__ function 'h' in __device__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-4 {{reference to __host__ function 'ch' in __device__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
+               // expected-note@67 {{'g' declared here}}
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}}
+  // expected-error@-4 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}}
+#endif
+  dh();
+  cdh();
+  g(); // expected-error {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}}
+  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+}
+
+__global__ void globalf(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __host__ function 'h' in __global__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-4 {{reference to __host__ function 'ch' in __global__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
+               // expected-note@67 {{'g' declared here}}
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}}
+  // expected-error@-4 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}}
+#endif
+  dh();
+  cdh();
+  g(); // expected-error {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}}
+  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
+  // expected-note@67 {{'g' declared here}}
+}
+
+__host__ __device__ void hostdevicef(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}}
+  // expected-note@65 {{'d' declared here}}
+  // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+  // expected-note@68 {{'cd' declared here}}
+#else
+  // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g;
+#if defined(__CUDA_ARCH__)
+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+#endif
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-6 {{no matching function for call to 'd'}}
+  // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+  // expected-error@-7 {{no matching function for call to 'cd'}}
+  // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#else
+  // expected-error@-9 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+  // expected-error@-10 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+#endif
+
+  dh();
+  cdh();
+  g();
+  g<<<0,0>>>();
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-3 {{call to global function g not configured}}
+#else
+  // expected-error@-5 {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
+  // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+#endif  // __CUDA_ARCH__
+}
+
+// Test for address of overloaded function resolution in the global context.
+fp_t hp = h;
+fp_t chp = ch;
+fp_t dhp = dh;
+fp_t cdhp = cdh;
+gp_t gp = g;
+
+
+// Test overloading of destructors
+// Can't mix H and unattributed destructors
+struct d_h {
+  ~d_h() {} // expected-note {{previous declaration is here}}
+  __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+// H/D overloading is OK
+struct d_dh {
+  __device__ ~d_dh() {}
+  __host__ ~d_dh() {}
+};
+
+// HD is OK
+struct d_hd {
+  __host__ __device__ ~d_hd() {}
+};
+
+// Mixing H/D and HD is not allowed.
+struct d_dhhd {
+  __device__ ~d_dhhd() {}
+  __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hhd {
+  __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hdh {
+  __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
+  __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_dhd {
+  __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hdd {
+  __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
+  __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+// Test overloading of member functions
+struct m_h {
+  void operator delete(void *ptr); // expected-note {{previous declaration is here}}
+  __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
+};
+
+// D/H overloading is OK
+struct m_dh {
+  __device__ void operator delete(void *ptr);
+  __host__ void operator delete(void *ptr);
+};
+
+// HD by itself is OK
+struct m_hd {
+  __device__ __host__ void operator delete(void *ptr);
+};
+
+struct m_hhd {
+  __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
+
+struct m_hdh {
+  __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
+
+struct m_dhd {
+  __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
+
+struct m_hdd {
+  __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
diff --git a/test/SemaCUDA/function-target-hd.cu b/test/SemaCUDA/function-target-hd.cu
index 25fcc6e9188f..685f4f9cda62 100644
--- a/test/SemaCUDA/function-target-hd.cu
+++ b/test/SemaCUDA/function-target-hd.cu
@@ -8,9 +8,9 @@
 // host device functions are not allowed to call device functions.
 
 // RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
 
 #include "Inputs/cuda.h"
 
diff --git a/test/SemaCUDA/implicit-intrinsic.cu b/test/SemaCUDA/implicit-intrinsic.cu
index 3d24aa719e57..0793d64b1017 100644
--- a/test/SemaCUDA/implicit-intrinsic.cu
+++ b/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();
 }