yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
Clang infers defaulted ctor as `__device__ __host__` and virtual dtor as
`__host__`.
It diagnose the following code in device compilation as B() references ~A()
implicitly.
struct A { virtual ~A(); };
struct B: public A { B(); };
B::B() = default;
Clang should not diagnose the above code since B() may be used in host function
only.
Clang should diagnose only if B() is used in device function or kernel.
This patch fixes that by assuming defaulted class member as not always emitted.
Therefore
a definition of defaulted special member function does not trigger a
diagnostic, unless
the defaulted special member function is really called.
https://reviews.llvm.org/D67509
Files:
lib/Sema/SemaCUDA.cpp
test/SemaCUDA/default-ctor.cu
Index: test/SemaCUDA/default-ctor.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/default-ctor.cu
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only
-fcuda-is-device -verify %s
+
+#define __host__ __attribute__((host))
+#define __device__ __attribute__((device))
+
+struct Base {
+ virtual ~Base(); // expected-note {{'~Base' declared here}}
+};
+
+struct Unused : public Base {
+ Unused();
+};
+
+// Unused defaulted constructor should not cause diagnostics.
+Unused::Unused() = default;
+
+struct Used : public Base // expected-note {{'~Used' declared here}}
+{
+ Used();
+};
+
+Used::Used() = default; // expected-error {{reference to __host__ function
'~Base' in __host__ __device__ function}}
+
+__device__ void f() {
+ Used x; // expected-error {{reference to __host__ function '~Used' in
__device__ function}}
+ // expected-note@-1 {{called by 'f'}}
+}
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -611,8 +611,12 @@
// emitted, because (say) the definition could include "inline".
FunctionDecl *Def = FD->getDefinition();
+ CXXMethodDecl *MD = dyn_cast_or_null<CXXMethodDecl>(FD);
+
if (Def &&
-
!isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(Def)))
+ !isDiscardableGVALinkage(
+ S.getASTContext().GetGVALinkageForFunction(Def)) &&
+ !(MD && MD->isDefaulted()))
return true;
// Otherwise, the function is known-emitted if it's in our set of
Index: test/SemaCUDA/default-ctor.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/default-ctor.cu
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
+
+#define __host__ __attribute__((host))
+#define __device__ __attribute__((device))
+
+struct Base {
+ virtual ~Base(); // expected-note {{'~Base' declared here}}
+};
+
+struct Unused : public Base {
+ Unused();
+};
+
+// Unused defaulted constructor should not cause diagnostics.
+Unused::Unused() = default;
+
+struct Used : public Base // expected-note {{'~Used' declared here}}
+{
+ Used();
+};
+
+Used::Used() = default; // expected-error {{reference to __host__ function '~Base' in __host__ __device__ function}}
+
+__device__ void f() {
+ Used x; // expected-error {{reference to __host__ function '~Used' in __device__ function}}
+ // expected-note@-1 {{called by 'f'}}
+}
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -611,8 +611,12 @@
// emitted, because (say) the definition could include "inline".
FunctionDecl *Def = FD->getDefinition();
+ CXXMethodDecl *MD = dyn_cast_or_null<CXXMethodDecl>(FD);
+
if (Def &&
- !isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(Def)))
+ !isDiscardableGVALinkage(
+ S.getASTContext().GetGVALinkageForFunction(Def)) &&
+ !(MD && MD->isDefaulted()))
return true;
// Otherwise, the function is known-emitted if it's in our set of
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
