9 files changed, 532 insertions, 8 deletions
diff --git a/test/SemaCUDA/amdgpu-num-gpr-attr.cu b/test/SemaCUDA/amdgpu-num-gpr-attr.cu
new file mode 100644
index 0000000..83acbc5
--- /dev/null
+++ b/test/SemaCUDA/amdgpu-num-gpr-attr.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+__attribute__((amdgpu_num_vgpr(64)))
+__global__ void test_num_vgpr() { } // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+
+__attribute__((amdgpu_num_sgpr(32)))
+__global__ void test_num_sgpr() { } // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+
+// fixme-expected-error@+3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+// expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
+__global__ void test_num_vgpr_num_sgpr() { }
diff --git a/test/SemaCUDA/function-target.cu b/test/SemaCUDA/function-target.cu
index 51bc8c9..ca56030 100644
--- a/test/SemaCUDA/function-target.cu
+++ b/test/SemaCUDA/function-target.cu
@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
 
 #include "Inputs/cuda.h"
 
@@ -31,14 +32,40 @@ __device__ void d1(void) {
   d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
 }
 
-__host__ void hd1h(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-__device__ void hd1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+// Expected 0-1 as in one of host/device side compilation it is an error, while
+// not in the other
+__host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+__device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+__host__ void hd1hg(void);
+__device__ void hd1dg(void);
+#ifdef __CUDA_ARCH__
+__host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#else
+__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#endif
 __host__ __device__ void hd1hd(void);
 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
 
 __host__ __device__ void hd1(void) {
-  hd1h(); // expected-error {{no matching function}}
-  hd1d(); // expected-error {{no matching function}}
+  // Expected 0-1 as in one of host/device side compilation it is an error,
+  // while not in the other
+  hd1d(); // expected-error 0-1 {{no matching function}}
+  hd1h(); // expected-error 0-1 {{no matching function}}
+
+  // No errors as guarded
+#ifdef __CUDA_ARCH__
+  hd1d();
+#else
+  hd1h();
+#endif
+
+  // Errors as incorrectly guarded
+#ifndef __CUDA_ARCH__
+  hd1dig(); // expected-error {{no matching function}}
+#else
+  hd1hig(); // expected-error {{no matching function}}
+#endif
+
   hd1hd();
   hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
 }
diff --git a/test/SemaCUDA/implicit-copy.cu b/test/SemaCUDA/implicit-copy.cu
new file mode 100644
index 0000000..b1b4887
--- /dev/null
+++ b/test/SemaCUDA/implicit-copy.cu
@@ -0,0 +1,51 @@
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fcuda-is-device -fsyntax-only -verify %s
+
+struct CopyableH {
+  const CopyableH& operator=(const CopyableH& x) { return *this; }
+};
+struct CopyableD {
+  __attribute__((device)) const CopyableD& operator=(const CopyableD x) { return *this; }
+};
+
+struct SimpleH {
+  CopyableH b;
+};
+// expected-note@-3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __host__ function from __device__ function}}
+// expected-note@-4 2 {{candidate function (the implicit move assignment operator) not viable: call to __host__ function from __device__ function}}
+
+struct SimpleD {
+  CopyableD b;
+};
+// expected-note@-3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-4 2 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void foo1hh() {
+  SimpleH a, b;
+  a = b;
+}
+__attribute__((device)) void foo1hd() {
+  SimpleH a, b;
+  a = b; // expected-error {{no viable overloaded}}
+}
+void foo1dh() {
+  SimpleD a, b;
+  a = b; // expected-error {{no viable overloaded}}
+}
+__attribute__((device)) void foo1dd() {
+  SimpleD a, b;
+  a = b;
+}
+
+void foo2hh(SimpleH &a, SimpleH &b) {
+  a = b;
+}
+__attribute__((device)) void foo2hd(SimpleH &a, SimpleH &b) {
+  a = b; // expected-error {{no viable overloaded}}
+}
+void foo2dh(SimpleD &a, SimpleD &b) {
+  a = b; // expected-error {{no viable overloaded}}
+}
+__attribute__((device)) void foo2dd(SimpleD &a, SimpleD &b) {
+  a = b;
+}
diff --git a/test/SemaCUDA/implicit-intrinsic.cu b/test/SemaCUDA/implicit-intrinsic.cu
new file mode 100644
index 0000000..3d24aa7
--- /dev/null
+++ b/test/SemaCUDA/implicit-intrinsic.cu
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -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.
+  __nvvm_membar_sys();
+}
diff --git a/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
new file mode 100644
index 0000000..f038c37
--- /dev/null
+++ b/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
@@ -0,0 +1,111 @@
+// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: collision between two bases
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+
+struct B1_with_device_ctor {
+  __device__ B1_with_device_ctor() {}
+};
+
+struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit default constructor) not viable}}
+// expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo1() {
+  C1_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+  A1_with_host_ctor aa;
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-8 {{candidate constructor (the implicit move constructor}} not viable
+
+void hostfoo2() {
+  C2_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-7 {{candidate constructor (the implicit move constructor}} not viable
+
+void hostfoo3() {
+  C3_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 4: collision on resolving a copy ctor
+
+struct A4_with_host_copy_ctor {
+  A4_with_host_copy_ctor() {}
+  A4_with_host_copy_ctor(const A4_with_host_copy_ctor&) {}
+};
+
+struct B4_with_device_copy_ctor {
+  B4_with_device_copy_ctor() {}
+  __device__ B4_with_device_copy_ctor(const B4_with_device_copy_ctor&) {}
+};
+
+struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-4 {{implicit copy constructor inferred target collision}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo4() {
+  C4_with_collision c;
+  C4_with_collision c2 = c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: collision on resolving a move ctor
+
+struct A5_with_host_move_ctor {
+  A5_with_host_move_ctor() {}
+  A5_with_host_move_ctor(A5_with_host_move_ctor&&) {}
+// expected-note@-1 {{copy constructor is implicitly deleted because 'A5_with_host_move_ctor' has a user-declared move constructor}}
+};
+
+struct B5_with_device_move_ctor {
+  B5_with_device_move_ctor() {}
+  __device__ B5_with_device_move_ctor(B5_with_device_move_ctor&&) {}
+};
+
+struct C5_with_collision : A5_with_host_move_ctor, B5_with_device_move_ctor {
+};
+// expected-note@-2 {{deleted}}
+
+void hostfoo5() {
+  C5_with_collision c;
+  // What happens here:
+  // This tries to find the move ctor. Since the move ctor is deleted due to
+  // collision, it then looks for a copy ctor. But copy ctors are implicitly
+  // deleted when move ctors are declared explicitly.
+  C5_with_collision c2(static_cast<C5_with_collision&&>(c)); // expected-error {{call to implicitly-deleted}}
+}
diff --git a/test/SemaCUDA/implicit-member-target-collision.cu b/test/SemaCUDA/implicit-member-target-collision.cu
new file mode 100644
index 0000000..25cdccb
--- /dev/null
+++ b/test/SemaCUDA/implicit-member-target-collision.cu
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: collision between two bases
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+
+struct B1_with_device_ctor {
+  __device__ B1_with_device_ctor() {}
+};
+
+struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo1() {
+  C1_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+  A1_with_host_ctor aa;
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo2() {
+  C2_with_collision c; // expected-error {{no matching constructor}}
+
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo3() {
+  C3_with_collision c; // expected-error {{no matching constructor}}
+}
diff --git a/test/SemaCUDA/implicit-member-target.cu b/test/SemaCUDA/implicit-member-target.cu
new file mode 100644
index 0000000..6064560
--- /dev/null
+++ b/test/SemaCUDA/implicit-member-target.cu
@@ -0,0 +1,186 @@
+// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: infer default ctor to be host.
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+
+// The implicit default constructor is inferred to be host because it only needs
+// to invoke a single host constructor (A1_with_host_ctor's). So we'll encounter
+// an error when calling it from a __device__ function, but not from a __host__
+// function.
+struct B1_with_implicit_default_ctor : A1_with_host_ctor {
+};
+
+// expected-note@-3 {{call to __host__ function from __device__}}
+// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo() {
+  B1_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo() {
+  B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: infer default ctor to be device.
+
+struct A2_with_device_ctor {
+  __device__ A2_with_device_ctor() {}
+};
+
+struct B2_with_implicit_default_ctor : A2_with_device_ctor {
+};
+
+// expected-note@-3 {{call to __device__ function from __host__}}
+// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo2() {
+  B2_with_implicit_default_ctor b;  // expected-error {{no matching constructor}}
+}
+
+__device__ void devicefoo2() {
+  B2_with_implicit_default_ctor b;
+}
+
+//------------------------------------------------------------------------------
+// Test 3: infer copy ctor
+
+struct A3_with_device_ctors {
+  __host__ A3_with_device_ctors() {}
+  __device__ A3_with_device_ctors(const A3_with_device_ctors&) {}
+};
+
+struct B3_with_implicit_ctors : A3_with_device_ctors {
+};
+
+// expected-note@-3 {{copy constructor of 'B3_with_implicit_ctors' is implicitly deleted}}
+
+void hostfoo3() {
+  B3_with_implicit_ctors b;  // this is OK because the inferred default ctor
+                             // here is __host__
+  B3_with_implicit_ctors b2 = b; // expected-error {{call to implicitly-deleted copy constructor}}
+
+}
+
+//------------------------------------------------------------------------------
+// Test 4: infer default ctor from a field, not a base
+
+struct A4_with_host_ctor {
+  A4_with_host_ctor() {}
+};
+
+struct B4_with_implicit_default_ctor {
+  A4_with_host_ctor field;
+};
+
+// expected-note@-4 {{call to __host__ function from __device__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo4() {
+  B4_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo4() {
+  B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: copy ctor with non-const param
+
+struct A5_copy_ctor_constness {
+  __host__ A5_copy_ctor_constness() {}
+  __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}
+};
+
+struct B5_copy_ctor_constness : A5_copy_ctor_constness {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
+// expected-note@-4 {{candidate constructor (the implicit default constructor) not viable}}
+
+void hostfoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg;
+}
+
+__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: explicitly defaulted ctor: since they are spelled out, they have
+// a host/device designation explicitly so no inference needs to be done.
+
+struct A6_with_device_ctor {
+  __device__ A6_with_device_ctor() {}
+};
+
+struct B6_with_defaulted_ctor : A6_with_device_ctor {
+  __host__ B6_with_defaulted_ctor() = default;
+};
+
+// expected-note@-3 {{candidate constructor not viable: call to __host__ function from __device__ function}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+__device__ void devicefoo6() {
+  B6_with_defaulted_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 7: copy assignment operator
+
+struct A7_with_copy_assign {
+  A7_with_copy_assign() {}
+  __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}
+};
+
+struct B7_with_copy_assign : A7_with_copy_assign {
+};
+
+// expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void hostfoo7() {
+  B7_with_copy_assign b1, b2;
+  b1 = b2; // expected-error {{no viable overloaded '='}}
+}
+
+//------------------------------------------------------------------------------
+// Test 8: move assignment operator
+
+// definitions for std::move
+namespace std {
+inline namespace foo {
+template <class T> struct remove_reference { typedef T type; };
+template <class T> struct remove_reference<T&> { typedef T type; };
+template <class T> struct remove_reference<T&&> { typedef T type; };
+
+template <class T> typename remove_reference<T>::type&& move(T&& t);
+}
+}
+
+struct A8_with_move_assign {
+  A8_with_move_assign() {}
+  __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
+  __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {}
+};
+
+struct B8_with_move_assign : A8_with_move_assign {
+};
+
+// expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void hostfoo8() {
+  B8_with_move_assign b1, b2;
+  b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
+}
diff --git a/test/SemaCUDA/launch_bounds.cu b/test/SemaCUDA/launch_bounds.cu
index bed7658..8edc41b 100644
--- a/test/SemaCUDA/launch_bounds.cu
+++ b/test/SemaCUDA/launch_bounds.cu
@@ -6,9 +6,6 @@ __launch_bounds__(128, 7) void Test1(void);
 __launch_bounds__(128) void Test2(void);
 
 __launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
-
-// FIXME: the error should read that the attribute takes exactly one or two arguments, but there
-// is no support for such a diagnostic currently.
-__launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
+__launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
 
 int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
diff --git a/test/SemaCUDA/method-target.cu b/test/SemaCUDA/method-target.cu
new file mode 100644
index 0000000..4fa2907
--- /dev/null
+++ b/test/SemaCUDA/method-target.cu
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: host method called from device function
+
+struct S1 {
+  void method() {}
+};
+
+__device__ void foo1(S1& s) {
+  s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: host method called from device function, for overloaded method
+
+struct S2 {
+  void method(int) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
+  void method(float) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
+};
+
+__device__ void foo2(S2& s, int i, float f) {
+  s.method(f); // expected-error {{no matching member function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 3: device method called from host function
+
+struct S3 {
+  __device__ void method() {}
+};
+
+void foo3(S3& s) {
+  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 4: device method called from host&device function
+
+struct S4 {
+  __device__ void method() {}
+};
+
+__host__ __device__ void foo4(S4& s) {
+  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: overloaded operators
+
+struct S5 {
+  S5() {}
+  S5& operator=(const S5&) {return *this;} // expected-note {{candidate function not viable}}
+};
+
+__device__ void foo5(S5& s, S5& t) {
+  s = t; // expected-error {{no viable overloaded '='}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: call method through pointer
+
+struct S6 {
+  void method() {}
+};
+
+__device__ void foo6(S6* s) {
+  s->method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+}