From e605c161735ecf1ef4c825b77c4df1dcdf02fceb Mon Sep 17 00:00:00 2001 From: Jacques Pienaar Date: Tue, 16 Dec 2014 20:12:38 +0000 Subject: [PATCH] Consider calls from implict host device functions as valid in SemaCUDA. In SemaCUDA all implicit functions were considered host device, this led to errors such as the following code snippet failing to compile: struct Copyable { const Copyable& operator=(const Copyable& x) { return *this; } }; struct Simple { Copyable b; }; void foo() { Simple a, b; a = b; } Above the implicit copy assignment operator was inferred as host device but there was only a host assignment copy defined which is an error in device compilation mode. Differential Revision: http://reviews.llvm.org/D6565 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@224358 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Sema/Sema.h | 3 -- lib/Sema/SemaCUDA.cpp | 12 +++--- lib/Sema/SemaExpr.cpp | 7 ++-- test/SemaCUDA/implicit-copy.cu | 51 +++++++++++++++++++++++++ test/SemaCUDA/implicit-member-target.cu | 10 +++-- 5 files changed, 66 insertions(+), 17 deletions(-) create mode 100644 test/SemaCUDA/implicit-copy.cu diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h index 10d8b03fcc..17154b636d 100644 --- a/include/clang/Sema/Sema.h +++ b/include/clang/Sema/Sema.h @@ -8325,9 +8325,6 @@ public: 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 diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp index 46a83173fc..64222fbf8a 100644 --- a/lib/Sema/SemaCUDA.cpp +++ b/lib/Sema/SemaCUDA.cpp @@ -62,12 +62,9 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { 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 @@ bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget, // 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; diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp index 04497f3efd..22cc925c25 100644 --- a/lib/Sema/SemaExpr.cpp +++ b/lib/Sema/SemaExpr.cpp @@ -1590,11 +1590,10 @@ Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK, if (getLangOpts().CUDA) if (const FunctionDecl *Caller = dyn_cast(CurContext)) if (const FunctionDecl *Callee = dyn_cast(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(); diff --git a/test/SemaCUDA/implicit-copy.cu b/test/SemaCUDA/implicit-copy.cu new file mode 100644 index 0000000000..b1b4887f56 --- /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-member-target.cu b/test/SemaCUDA/implicit-member-target.cu index 87b05e61a6..6064560f2c 100644 --- a/test/SemaCUDA/implicit-member-target.cu +++ b/test/SemaCUDA/implicit-member-target.cu @@ -146,11 +146,12 @@ struct A7_with_copy_assign { 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 A8_with_move_assign { 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 '='}} } -- 2.40.0