[CUDA] Disallow overloading destructors.

Summary:
We'd attempted to allow this, but turns out we were doing a very bad
job.  :)

Making this work properly would be a giant change in clang.  For
example, we'd need to make CXXRecordDecl::getDestructor()
context-sensitive, because the destructor you end up with depends on
where you're calling it from.

For now (and hopefully for ever), just disallow overloading of
destructors in CUDA.

Reviewers: rsmith

Subscribers: cfe-commits, tra

Differential Revision: https://reviews.llvm.org/D24571

llvm-svn: 283120

5 files changed, 38 insertions, 58 deletions

diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 10aa9977459..ecd5bd2b99b 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -1129,6 +1129,11 @@ bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old,
   }
 
   if (getLangOpts().CUDA && ConsiderCudaAttrs) {
+    // Don't allow overloading of destructors.  (In theory we could, but it
+    // would be a giant change to clang.)
+    if (isa<CXXDestructorDecl>(New))
+      return false;
+
     CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New),
                        OldTarget = IdentifyCUDATarget(Old);
     if (NewTarget == CFT_InvalidTarget || NewTarget == CFT_Global)
diff --git a/clang/test/CodeGenCUDA/function-overload.cu b/clang/test/CodeGenCUDA/function-overload.cu
index 380304af822..c82b2e96f6c 100644
--- a/clang/test/CodeGenCUDA/function-overload.cu
+++ b/clang/test/CodeGenCUDA/function-overload.cu
@@ -16,8 +16,6 @@ int x;
 struct s_cd_dh {
   __host__ s_cd_dh() { x = 11; }
   __device__ s_cd_dh() { x = 12; }
-  __host__ ~s_cd_dh() { x = 21; }
-  __device__ ~s_cd_dh() { x = 22; }
 };
 
 struct s_cd_hd {
@@ -38,7 +36,6 @@ void wrapper() {
   // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
 
   // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
-  // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
 }
 // CHECK-BOTH: ret void
 
@@ -56,8 +53,3 @@ void wrapper() {
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
 // CHECK-BOTH: store i32 32,
 // CHECK-BOTH: ret void
-
-// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
-// CHECK-HOST:   store i32 21,
-// CHECK-DEVICE: store i32 22,
-// CHECK-BOTH: ret void
diff --git a/clang/test/SemaCUDA/call-overloaded-destructor.cu b/clang/test/SemaCUDA/call-overloaded-destructor.cu
deleted file mode 100644
index 24b0e7d330e..00000000000
--- a/clang/test/SemaCUDA/call-overloaded-destructor.cu
+++ /dev/null
@@ -1,17 +0,0 @@
-// 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();
-}
diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu
index ff1db14b494..6071dbb2748 100644
--- a/clang/test/SemaCUDA/function-overload.cu
+++ b/clang/test/SemaCUDA/function-overload.cu
@@ -210,44 +210,11 @@ struct d_h {
   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
 };
 
-// H/D overloading is OK
-struct d_dh {
-  __device__ ~d_dh() {}
-  __host__ ~d_dh() {}
-};
-
 // HD is OK
 struct d_hd {
   __host__ __device__ ~d_hd() {}
 };
 
-// Mixing H/D and HD is not allowed.
-struct d_dhhd {
-  __device__ ~d_dhhd() {}
-  __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
-  __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
-};
-
-struct d_hhd {
-  __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
-  __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
-};
-
-struct d_hdh {
-  __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
-  __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
-};
-
-struct d_dhd {
-  __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
-  __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
-};
-
-struct d_hdd {
-  __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
-  __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
-};
-
 // Test overloading of member functions
 struct m_h {
   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
diff --git a/clang/test/SemaCUDA/no-destructor-overload.cu b/clang/test/SemaCUDA/no-destructor-overload.cu
new file mode 100644
index 00000000000..aa6971ee8ca
--- /dev/null
+++ b/clang/test/SemaCUDA/no-destructor-overload.cu
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// We don't allow destructors to be overloaded.  Making this work would be a
+// giant change to clang, and the use cases seem quite limited.
+
+struct A {
+  ~A() {} // expected-note {{previous declaration is here}}
+  __device__ ~A() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct B {
+  __host__ ~B() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~B() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct C {
+  __host__ __device__ ~C() {} // expected-note {{previous declaration is here}}
+  __host__ ~C() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct D {
+  __device__ ~D() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~D() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct E {
+  __host__ __device__ ~E() {} // expected-note {{previous declaration is here}}
+  __device__ ~E() {} // expected-error {{destructor cannot be redeclared}}
+};
+