r224358 - Consider calls from implict host device functions as valid in SemaCUDA.
Jacques Pienaar
jpienaar at google.com
Tue Dec 16 12:12:38 PST 2014
Author: jpienaar
Date: Tue Dec 16 14:12:38 2014
New Revision: 224358
URL: http://llvm.org/viewvc/llvm-project?rev=224358&view=rev
Log:
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
Added:
cfe/trunk/test/SemaCUDA/implicit-copy.cu
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/test/SemaCUDA/implicit-member-target.cu
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=224358&r1=224357&r2=224358&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Tue Dec 16 14:12:38 2014
@@ -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
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=224358&r1=224357&r2=224358&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Dec 16 14:12:38 2014
@@ -62,12 +62,9 @@ Sema::CUDAFunctionTarget Sema::IdentifyC
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(CUDAFunctionT
// 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;
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=224358&r1=224357&r2=224358&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Dec 16 14:12:38 2014
@@ -1590,11 +1590,10 @@ Sema::BuildDeclRefExpr(ValueDecl *D, Qua
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();
Added: cfe/trunk/test/SemaCUDA/implicit-copy.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-copy.cu?rev=224358&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-copy.cu (added)
+++ cfe/trunk/test/SemaCUDA/implicit-copy.cu Tue Dec 16 14:12:38 2014
@@ -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 at -3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __host__ function from __device__ function}}
+// expected-note at -4 2 {{candidate function (the implicit move assignment operator) not viable: call to __host__ function from __device__ function}}
+
+struct SimpleD {
+ CopyableD b;
+};
+// expected-note at -3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note at -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;
+}
Modified: cfe/trunk/test/SemaCUDA/implicit-member-target.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-member-target.cu?rev=224358&r1=224357&r2=224358&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-member-target.cu Tue Dec 16 14:12:38 2014
@@ -146,11 +146,12 @@ struct A7_with_copy_assign {
struct B7_with_copy_assign : A7_with_copy_assign {
};
-// expected-note at -3 {{copy assignment operator of 'B7_with_copy_assign' is implicitly deleted}}
+// expected-note at -3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note at -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 at -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 at -3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note at -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 '='}}
}
More information about the cfe-commits
mailing list