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

[jlebar]

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.

[jlebar]

Friendly ping

[rnk]

lgtm

cfe/trunk/lib/Sema/SemaOverload.cpp

if (isa<CXXThisExpr>(MemExpr->getBase()->IgnoreParenCasts()) &&
              << MD->getParent()->getDeclName()
              << MD->getDeclName();
     }
   }
 
   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());
   }
 
   return MaybeBindToTemporary(TheCall);
 }
 

cfe/trunk/test/SemaCUDA/call-overloaded-destructor.cu

+// 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();
+}