Author: Yaxun (Sam) Liu Date: 2023-11-22T21:20:53-05:00 New Revision: 6b3470b4b83195aeeda60b101e8d3bf8800c321c
URL: https://github.com/llvm/llvm-project/commit/6b3470b4b83195aeeda60b101e8d3bf8800c321c DIFF: https://github.com/llvm/llvm-project/commit/6b3470b4b83195aeeda60b101e8d3bf8800c321c.diff LOG: Revert "[CUDA][HIP] make trivial ctor/dtor host device (#72394)" This reverts commit 27e6e4a4d0e3296cebad8db577ec0469a286795e. This patch is reverted due to regression. A testcase is: `template <class T> struct ptr { ~ptr() { static int x = 1;} }; template <class T> struct Abc : ptr<T> { public: Abc(); ~Abc() {} }; template class Abc<int>; ` Added: Modified: clang/include/clang/Sema/Sema.h clang/lib/Sema/SemaCUDA.cpp clang/lib/Sema/SemaDecl.cpp clang/test/SemaCUDA/call-host-fn-from-device.cu clang/test/SemaCUDA/default-ctor.cu clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu clang/test/SemaCUDA/implicit-member-target-collision.cu clang/test/SemaCUDA/implicit-member-target-inherited.cu clang/test/SemaCUDA/implicit-member-target.cu Removed: clang/test/SemaCUDA/trivial-ctor-dtor.cu ################################################################################ diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 59806bcbcbb2dbc..e8914f5fcddf19e 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13466,10 +13466,6 @@ class Sema final { void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD, const LookupResult &Previous); - /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to a - /// trivial cotr/dtor that does not have host and device attributes. - void maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FunctionDecl *FD); - /// May add implicit CUDAConstantAttr attribute to VD, depending on VD /// and current compilation settings. void MaybeAddCUDAConstantAttr(VarDecl *VD); diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index b94f448dabe7517..318174f7be8fa95 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -772,22 +772,6 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD, NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context)); } -// If a trivial ctor/dtor has no host/device -// attributes, make it implicitly host device function. -void Sema::maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FunctionDecl *FD) { - bool IsTrivialCtor = false; - if (auto *CD = dyn_cast<CXXConstructorDecl>(FD)) - IsTrivialCtor = isEmptyCudaConstructor(SourceLocation(), CD); - bool IsTrivialDtor = false; - if (auto *DD = dyn_cast<CXXDestructorDecl>(FD)) - IsTrivialDtor = isEmptyCudaDestructor(SourceLocation(), DD); - if ((IsTrivialCtor || IsTrivialDtor) && !FD->hasAttr<CUDAHostAttr>() && - !FD->hasAttr<CUDADeviceAttr>()) { - FD->addAttr(CUDAHostAttr::CreateImplicit(Context)); - FD->addAttr(CUDADeviceAttr::CreateImplicit(Context)); - } -} - // TODO: `__constant__` memory may be a limited resource for certain targets. // A safeguard may be needed at the end of compilation pipeline if // `__constant__` memory usage goes beyond limit. diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 4e1857b931cc868..23dd8ae15c16583 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -16255,9 +16255,6 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, if (FD && !FD->isDeleted()) checkTypeSupport(FD->getType(), FD->getLocation(), FD); - if (LangOpts.CUDA) - maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FD); - return dcl; } diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu index b62de92db02d6de..acdd291b664579b 100644 --- a/clang/test/SemaCUDA/call-host-fn-from-device.cu +++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu @@ -12,7 +12,7 @@ extern "C" void host_fn() {} struct Dummy {}; struct S { - S() { static int nontrivial_ctor = 1; } + S() {} // expected-note@-1 2 {{'S' declared here}} ~S() { host_fn(); } // expected-note@-1 {{'~S' declared here}} diff --git a/clang/test/SemaCUDA/default-ctor.cu b/clang/test/SemaCUDA/default-ctor.cu index 31971fe6b3863c7..cbad7a1774c1501 100644 --- a/clang/test/SemaCUDA/default-ctor.cu +++ b/clang/test/SemaCUDA/default-ctor.cu @@ -25,7 +25,7 @@ __device__ void fd() { InD ind; InH inh; // expected-error{{no matching constructor for initialization of 'InH'}} InHD inhd; - Out out; + Out out; // expected-error{{no matching constructor for initialization of 'Out'}} OutD outd; OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}} OutHD outhd; diff --git a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu index edb543f637ccc18..06015ed0d6d8edc 100644 --- a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu +++ b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu @@ -6,7 +6,7 @@ // Test 1: collision between two bases struct A1_with_host_ctor { - A1_with_host_ctor() { static int nontrivial_ctor = 1; } + A1_with_host_ctor() {} }; struct B1_with_device_ctor { diff --git a/clang/test/SemaCUDA/implicit-member-target-collision.cu b/clang/test/SemaCUDA/implicit-member-target-collision.cu index 16b5978af40872b..a50fddaa4615b22 100644 --- a/clang/test/SemaCUDA/implicit-member-target-collision.cu +++ b/clang/test/SemaCUDA/implicit-member-target-collision.cu @@ -6,7 +6,7 @@ // Test 1: collision between two bases struct A1_with_host_ctor { - A1_with_host_ctor() { static int nontrivial_ctor = 1; } + A1_with_host_ctor() {} }; struct B1_with_device_ctor { diff --git a/clang/test/SemaCUDA/implicit-member-target-inherited.cu b/clang/test/SemaCUDA/implicit-member-target-inherited.cu index 781199bba6b5a11..2178172ed01930d 100644 --- a/clang/test/SemaCUDA/implicit-member-target-inherited.cu +++ b/clang/test/SemaCUDA/implicit-member-target-inherited.cu @@ -6,7 +6,7 @@ // Test 1: infer inherited default ctor to be host. struct A1_with_host_ctor { - A1_with_host_ctor() { static int nontrivial_ctor = 1; } + A1_with_host_ctor() {} }; // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}} // expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}} @@ -83,7 +83,7 @@ void hostfoo3() { // Test 4: infer inherited default ctor from a field, not a base struct A4_with_host_ctor { - A4_with_host_ctor() { static int nontrivial_ctor = 1; } + A4_with_host_ctor() {} }; struct B4_with_inherited_host_ctor : A4_with_host_ctor{ diff --git a/clang/test/SemaCUDA/implicit-member-target.cu b/clang/test/SemaCUDA/implicit-member-target.cu index 552f8f2ebd94fd5..d87e69624043419 100644 --- a/clang/test/SemaCUDA/implicit-member-target.cu +++ b/clang/test/SemaCUDA/implicit-member-target.cu @@ -6,7 +6,7 @@ // Test 1: infer default ctor to be host. struct A1_with_host_ctor { - A1_with_host_ctor() { static int nontrivial_ctor = 1; } + A1_with_host_ctor() {} }; // The implicit default constructor is inferred to be host because it only needs @@ -75,7 +75,7 @@ void hostfoo3() { // Test 4: infer default ctor from a field, not a base struct A4_with_host_ctor { - A4_with_host_ctor() { static int nontrivial_ctor = 1; } + A4_with_host_ctor() {} }; struct B4_with_implicit_default_ctor { diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu deleted file mode 100644 index 1df8adc62bab590..000000000000000 --- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu +++ /dev/null @@ -1,40 +0,0 @@ -// RUN: %clang_cc1 -isystem %S/Inputs -fsyntax-only -verify %s -// RUN: %clang_cc1 -isystem %S/Inputs -fcuda-is-device -fsyntax-only -verify %s - -#include <cuda.h> - -// Check trivial ctor/dtor -struct A { - int x; - A() {} - ~A() {} -}; - -__device__ A a; - -// Check trivial ctor/dtor of template class -template<typename T> -struct TA { - T x; - TA() {} - ~TA() {} -}; - -__device__ TA<int> ta; - -// Check non-trivial ctor/dtor in parent template class -template<typename T> -struct TB { - T x; - TB() { static int nontrivial_ctor = 1; } - ~TB() {} -}; - -template<typename T> -struct TC : TB<T> { - T x; - TC() {} - ~TC() {} -}; - -__device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
