Avoided adding an additional attribute.
http://reviews.llvm.org/D6565
Files:
include/clang/Sema/Sema.h
lib/Sema/SemaCUDA.cpp
lib/Sema/SemaExpr.cpp
test/SemaCUDA/implicit-copy.cu
test/SemaCUDA/implicit-member-target.cu
EMAIL PREFERENCES
http://reviews.llvm.org/settings/panel/emailpreferences/
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8300,9 +8300,6 @@
CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
- bool CheckCUDATarget(CUDAFunctionTarget CallerTarget,
- CUDAFunctionTarget CalleeTarget);
-
bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee);
/// Given a implicit special member, infer its CUDA target from the
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -62,12 +62,9 @@
bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
- return CheckCUDATarget(IdentifyCUDATarget(Caller),
- IdentifyCUDATarget(Callee));
-}
+ CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
+ CalleeTarget = IdentifyCUDATarget(Callee);
-bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget,
- CUDAFunctionTarget CalleeTarget) {
// If one of the targets is invalid, the check always fails, no matter what
// the other target is.
if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
@@ -90,8 +87,11 @@
// however, in which case the function is compiled for both the host and the
// device. The __CUDA_ARCH__ macro [...] can be used to differentiate code
// paths between host and device."
- bool InDeviceMode = getLangOpts().CUDAIsDevice;
if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) {
+ // If the caller is implicit then the check always passes.
+ if (Caller->isImplicit()) return false;
+
+ bool InDeviceMode = getLangOpts().CUDAIsDevice;
if ((InDeviceMode && CalleeTarget != CFT_Device) ||
(!InDeviceMode && CalleeTarget != CFT_Host))
return true;
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -1590,11 +1590,10 @@
if (getLangOpts().CUDA)
if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
if (const FunctionDecl *Callee = dyn_cast<FunctionDecl>(D)) {
- CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
- CalleeTarget = IdentifyCUDATarget(Callee);
- if (CheckCUDATarget(CallerTarget, CalleeTarget)) {
+ if (CheckCUDATarget(Caller, Callee)) {
Diag(NameInfo.getLoc(), diag::err_ref_bad_target)
- << CalleeTarget << D->getIdentifier() << CallerTarget;
+ << IdentifyCUDATarget(Callee) << D->getIdentifier()
+ << IdentifyCUDATarget(Caller);
Diag(D->getLocation(), diag::note_previous_decl)
<< D->getIdentifier();
return ExprError();
Index: test/SemaCUDA/implicit-copy.cu
===================================================================
--- test/SemaCUDA/implicit-copy.cu
+++ 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;
+}
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
