[PATCH] D67509: [CUDA][HIP] Diagnose defaulted constructor only if it is used

[Yaxun Liu via Phabricator via cfe-commits]

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 at -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


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D67509.219946.patch
Type: text/x-patch
Size: 1629 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190912/dd64195e/attachment.bin>