[CUDA] Don't assume that destructors can't be overloaded.

Summary:
You can overload a destructor in CUDA, and SemaOverload needs to be
tweaked not to crash when it sees an explicit call to an overloaded
destructor.

Reviewers: rsmith

Subscribers: cfe-commits, tra

Differential Revision: http://reviews.llvm.org/D21912

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@275231 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp
index 038be79309f6f4b7efdb6e8a915c8cebb01a3940..b025a397edca077e4ee1242b2b05e8f4d9c6f139 100644 (file)
--- a/lib/Sema/SemaOverload.cpp
+++ b/lib/Sema/SemaOverload.cpp
@@ -12407,8 +12407,7 @@ Sema::BuildCallToMemberFunction(Scope *S, Expr *MemExprE,
   if (CXXDestructorDecl *DD =
           dyn_cast<CXXDestructorDecl>(TheCall->getMethodDecl())) {
     // a->A::f() doesn't go through the vtable, except in AppleKext mode.
-    bool CallCanBeVirtual = !cast<MemberExpr>(NakedMemExpr)->hasQualifier() ||
-                            getLangOpts().AppleKext;
+    bool CallCanBeVirtual = !MemExpr->hasQualifier() || getLangOpts().AppleKext;
     CheckVirtualDtorCall(DD, MemExpr->getLocStart(), /*IsDelete=*/false,
                          CallCanBeVirtual, /*WarnOnNonAbstractTypes=*/true,
                          MemExpr->getMemberLoc());

diff --git a/test/SemaCUDA/call-overloaded-destructor.cu b/test/SemaCUDA/call-overloaded-destructor.cu
new file mode 100644 (file)
index 0000000..24b0e7d
--- /dev/null
+++ b/test/SemaCUDA/call-overloaded-destructor.cu
@@ -0,0 +1,17 @@
+// expected-no-diagnostics
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+struct S {
+  __host__ ~S() {}
+  __device__ ~S() {}
+};
+
+__host__ __device__ void test() {
+  S s;
+  // This should not crash clang.
+  s.~S();
+}