Changed the inference to treat all implicit functions uniformly. Extended test
to include copy constructor.
http://reviews.llvm.org/D6565
Files:
include/clang/Basic/Attr.td
include/clang/Basic/DiagnosticSemaKinds.td
include/clang/Sema/Sema.h
lib/Sema/SemaCUDA.cpp
test/SemaCUDA/implicit-copy.cu
test/SemaCUDA/implicit-member-target.cu
Index: include/clang/Basic/Attr.td
===================================================================
--- include/clang/Basic/Attr.td
+++ include/clang/Basic/Attr.td
@@ -563,6 +563,13 @@
let Documentation = [Undocumented];
}
+def CUDAImplicitTarget : InheritableAttr {
+ let Spellings = [];
+ let Subjects = SubjectList<[Function]>;
+ let LangOpts = [CUDA];
+ let Documentation = [Undocumented];
+}
+
def CUDAInvalidTarget : InheritableAttr {
let Spellings = [];
let Subjects = SubjectList<[Function]>;
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -3068,8 +3068,8 @@
"function (the implicit copy assignment operator)|"
"function (the implicit move assignment operator)|"
"constructor (inherited)}0 not viable: call to "
- "%select{__device__|__global__|__host__|__host__ __device__|invalid}1 function from"
- " %select{__device__|__global__|__host__|__host__ __device__|invalid}2 function">;
+ "%select{__device__|__global__|__host__|__host__ __device__|implicit|invalid}1 function from"
+ " %select{__device__|__global__|__host__|__host__ __device__|implicit|invalid}2 function">;
def note_implicit_member_target_infer_collision : Note<
"implicit %select{"
"default constructor|"
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8295,6 +8295,7 @@
CFT_Global,
CFT_Host,
CFT_HostDevice,
+ CFT_ImplicitTarget,
CFT_InvalidTarget
};
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -52,9 +52,7 @@
} else if (D->hasAttr<CUDAHostAttr>()) {
return CFT_Host;
} else if (D->isImplicit()) {
- // Some implicit declarations (like intrinsic functions) are not marked.
- // Set the most lenient target on them for maximal flexibility.
- return CFT_HostDevice;
+ return CFT_ImplicitTarget;
}
return CFT_Host;
@@ -73,6 +71,10 @@
if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
return true;
+ // If one of the targets is implicit, the check always passes.
+ if (CallerTarget == CFT_ImplicitTarget || CalleeTarget == CFT_ImplicitTarget)
+ return false;
+
// CUDA B.1.1 "The __device__ qualifier declares a function that is [...]
// Callable from the device only."
if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
Index: test/SemaCUDA/implicit-copy.cu
===================================================================
--- test/SemaCUDA/implicit-copy.cu
+++ test/SemaCUDA/implicit-copy.cu
@@ -0,0 +1,46 @@
+// 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 { // expected-note 4 {{candidate function}}
+ CopyableH b;
+};
+struct SimpleD { // expected-note 4 {{candidate function}}
+ CopyableD b;
+};
+
+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;
+}
Index: test/SemaCUDA/implicit-member-target.cu
===================================================================
--- test/SemaCUDA/implicit-member-target.cu
+++ test/SemaCUDA/implicit-member-target.cu
@@ -146,11 +146,12 @@
struct B7_with_copy_assign : A7_with_copy_assign {
};
-// expected-note@-3 {{copy assignment operator of 'B7_with_copy_assign' is implicitly deleted}}
+// 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 {{object of type 'B7_with_copy_assign' cannot be assigned because its copy assignment operator is implicitly deleted}}
+ b1 = b2; // expected-error {{no viable overloaded '='}}
}
//------------------------------------------------------------------------------
@@ -176,9 +177,10 @@
struct B8_with_move_assign : A8_with_move_assign {
};
-// expected-note@-3 {{copy assignment operator of 'B8_with_move_assign' is implicitly deleted because base class 'A8_with_move_assign' has no copy assignment operator}}
+// 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 {{object of type 'B8_with_move_assign' cannot be assigned because its copy assignment operator is implicitly deleted}}
+ b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
