From 2a9beb2f57348499a2f08d604592b060c3f9ad00 Mon Sep 17 00:00:00 2001
From: Justin Lebar <jlebar@google.com>
Date: Tue, 12 Jul 2016 23:23:01 +0000
Subject: [PATCH] [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
---
 lib/Sema/SemaOverload.cpp                   |  3 +--
 test/SemaCUDA/call-overloaded-destructor.cu | 17 +++++++++++++++++
 2 files changed, 18 insertions(+), 2 deletions(-)
 create mode 100644 test/SemaCUDA/call-overloaded-destructor.cu

diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp
index 038be79309..b025a397ed 100644
--- 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
index 0000000000..24b0e7d330
--- /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();
+}
-- 
2.40.0