Author: jlebar Date: Mon Oct 3 11:48:23 2016 New Revision: 283120 URL: http://llvm.org/viewvc/llvm-project?rev=283120&view=rev Log: [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 Added: cfe/trunk/test/SemaCUDA/no-destructor-overload.cu Removed: cfe/trunk/test/SemaCUDA/call-overloaded-destructor.cu Modified: cfe/trunk/lib/Sema/SemaOverload.cpp cfe/trunk/test/CodeGenCUDA/function-overload.cu cfe/trunk/test/SemaCUDA/function-overload.cu Modified: cfe/trunk/lib/Sema/SemaOverload.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=283120&r1=283119&r2=283120&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaOverload.cpp (original) +++ cfe/trunk/lib/Sema/SemaOverload.cpp Mon Oct 3 11:48:23 2016 @@ -1129,6 +1129,11 @@ bool Sema::IsOverload(FunctionDecl *New, } 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) Modified: cfe/trunk/test/CodeGenCUDA/function-overload.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/function-overload.cu?rev=283120&r1=283119&r2=283120&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/function-overload.cu (original) +++ cfe/trunk/test/CodeGenCUDA/function-overload.cu Mon Oct 3 11:48:23 2016 @@ -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 Removed: cfe/trunk/test/SemaCUDA/call-overloaded-destructor.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/call-overloaded-destructor.cu?rev=283119&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/call-overloaded-destructor.cu (original) +++ cfe/trunk/test/SemaCUDA/call-overloaded-destructor.cu (removed) @@ -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(); -} Modified: cfe/trunk/test/SemaCUDA/function-overload.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=283120&r1=283119&r2=283120&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/function-overload.cu (original) +++ cfe/trunk/test/SemaCUDA/function-overload.cu Mon Oct 3 11:48:23 2016 @@ -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}} Added: cfe/trunk/test/SemaCUDA/no-destructor-overload.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/no-destructor-overload.cu?rev=283120&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/no-destructor-overload.cu (added) +++ cfe/trunk/test/SemaCUDA/no-destructor-overload.cu Mon Oct 3 11:48:23 2016 @@ -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}} +}; + _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits