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
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
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)
// 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;
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();
--- /dev/null
+// 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;
+}
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 '='}}
}
//------------------------------------------------------------------------------
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 '='}}
}