[CUDA] Disallow overloading destructors.

[jlebar]

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.

[rnk]

lgtm

clang/lib/Sema/SemaOverload.cpp

1131

I feel like we should exit early on destructors here, before we do any target checks. The assert also feels kind of trivial because we only come into this overload machinery if looking up New's DeclarationName found Old.

[jlebar]

clang/lib/Sema/SemaOverload.cpp

1131

Will do when I check this in. Thank you for the review!

clang/lib/Sema/SemaOverload.cpp

if (NewI == NewE || OldI == OldE)
       return true;
     llvm::FoldingSetNodeID NewID, OldID;
     NewI->getCond()->Profile(NewID, Context, true);
     OldI->getCond()->Profile(OldID, Context, true);
     if (NewID != OldID)
       return true;
   }
 
   if (getLangOpts().CUDA && ConsiderCudaAttrs) {
     CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New),
                        OldTarget = IdentifyCUDATarget(Old);
     if (NewTarget == CFT_InvalidTarget || NewTarget == CFT_Global)
       return false;
 
     assert((OldTarget != CFT_InvalidTarget) && "Unexpected invalid target.");
 
     // Don't allow mixing of HD with other kinds. This guarantees that
     // we have only one viable function with this signature on any
     // side of CUDA compilation .
     // __global__ functions can't be overloaded based on attribute
     // difference because, like HD, they also exist on both sides.
     if ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice) ||
         (NewTarget == CFT_Global) || (OldTarget == CFT_Global))
       return false;
 
-    // Allow overloading of functions with same signature, but
-    // different CUDA target attributes.
+    // Don't allow overloading of destructors.  (In theory we could, but it
+    // would be a giant change to clang.)
+    if (isa<CXXDestructorDecl>(New)) {
+      assert(isa<CXXDestructorDecl>(Old) &&
+             "If New is a destructor, Old must be as well.");
+      return false;
+    }
+
+    // Allow overloading of functions with same signature and different CUDA
+    // target attributes.
     return NewTarget != OldTarget;
   }
 
   // The signatures match; this is not an overload.
   return false;
 }
 
 /// \brief Checks availability of the function depending on the current

clang/test/CodeGenCUDA/function-overload.cu

 
 #include "Inputs/cuda.h"
 
 // Check constructors/destructors for D/H functions
 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 {
   __host__ __device__ s_cd_hd() { x = 31; }
   __host__ __device__ ~s_cd_hd() { x = 32; }
 };
 
 // CHECK-BOTH: define void @_Z7wrapperv
 #if defined(__CUDA_ARCH__)
 __device__
 #else
 __host__
 #endif
 void wrapper() {
   s_cd_dh scddh;
   // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
   s_cd_hd scdhd;
   // 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
 
 // Now it's time to check what's been generated for the methods we used.
 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
 // CHECK-HOST:   store i32 11,
 // CHECK-DEVICE: store i32 12,
 // CHECK-BOTH: ret void
 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
 // CHECK-BOTH:   store i32 31,
 // CHECK-BOTH: ret void
 
 // 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

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

clang/test/SemaCUDA/function-overload.cu

 
 // Test overloading of destructors
 // Can't mix H and unattributed destructors
 struct d_h {
   ~d_h() {} // expected-note {{previous declaration is here}}
   __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}}
   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
 };
 
 // D/H overloading is OK
 struct m_dh {

clang/test/SemaCUDA/no-destructor-overload.cu

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