From 7208e8d15fbf218deb74483ea8c549c67ca4985e Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 23 Nov 2023 07:49:22 -0500 Subject: [PATCH] [CUDA][HIP] allow trivial ctor/dtor in device var init Treat ctor/dtor in device var init as host device function so that they can be used to initialize file-scope device variables to match nvcc behavior. If they are non-trivial they will be diagnosed. We cannot add implicit host device attrs to non-trivial ctor/dtor since determining whether they are non-trivial needs to know whether they have a trivial body and all their member and base classes' ctor/dtor have trivial body, which is affected by where their bodies are defined or instantiated. Revert "[CUDA][HIP] make trivial ctor/dtor host device" This reverts commit 876f99a2dc86989e4bb88dc1e3bbbf7b18d98e28. Fixes: #72261 Fixes: SWDEV-432412 Fixes: SWDEV-433956 Change-Id: I711db63a2166ce77dea06aad5d04cae10d96ce24 --- clang/include/clang/Sema/Sema.h | 4 --- clang/lib/Sema/SemaCUDA.cpp | 25 +++++++------------ clang/lib/Sema/SemaDecl.cpp | 3 --- clang/lib/Sema/SemaOverload.cpp | 6 ++--- .../test/SemaCUDA/call-host-fn-from-device.cu | 2 +- clang/test/SemaCUDA/default-ctor.cu | 2 +- .../implicit-member-target-collision-cxx11.cu | 2 +- .../implicit-member-target-collision.cu | 2 +- .../implicit-member-target-inherited.cu | 5 ++-- clang/test/SemaCUDA/implicit-member-target.cu | 4 +-- clang/test/SemaCUDA/trivial-ctor-dtor.cu | 9 ++++--- 11 files changed, 24 insertions(+), 40 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 4bc62444930c68..0228d8327b0700 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13193,10 +13193,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 34d989b223467b..3fb66dd1dcad59 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -225,6 +225,15 @@ Sema::CUDAFunctionPreference Sema::IdentifyCUDAPreference(const FunctionDecl *Caller, const FunctionDecl *Callee) { assert(Callee && "Callee must be valid."); + + // Treat ctor/dtor as host device function in device var initializer to allow + // trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor + // will be diagnosed by checkAllowedCUDAInitializer. + if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar && + CurCUDATargetCtx.Target == CFT_Device && + (isa(Callee) || isa(Callee))) + return CFP_HostDevice; + CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller); CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee); @@ -730,22 +739,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(FD)) - IsTrivialCtor = isEmptyCudaConstructor(SourceLocation(), CD); - bool IsTrivialDtor = false; - if (auto *DD = dyn_cast(FD)) - IsTrivialDtor = isEmptyCudaDestructor(SourceLocation(), DD); - if ((IsTrivialCtor || IsTrivialDtor) && !FD->hasAttr() && - !FD->hasAttr()) { - 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 0d0737d9246cdd..7596b51bf05457 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -15884,9 +15884,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/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index fe6cfdd8df74a6..8e81c0d7c8dd7f 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -1404,10 +1404,8 @@ bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old, // Don't allow overloading of destructors. (In theory we could, but it // would be a giant change to clang.) if (!isa(New)) { - CUDAFunctionTarget NewTarget = IdentifyCUDATarget( - New, isa(New)), - OldTarget = IdentifyCUDATarget( - Old, isa(New)); + CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New), + OldTarget = IdentifyCUDATarget(Old); if (NewTarget != CFT_InvalidTarget) { assert((OldTarget != CFT_InvalidTarget) && "Unexpected invalid target."); diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu index b62de92db02d6d..acdd291b664579 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 31971fe6b3863c..cbad7a1774c150 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 edb543f637ccc1..06015ed0d6d8ed 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 16b5978af40872..a50fddaa4615b2 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 ceca0891fc9b03..2178172ed01930 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}} @@ -39,7 +39,6 @@ struct A2_with_device_ctor { }; // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}} // expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}} -// expected-note@-4 {{candidate inherited constructor not viable: call to __device__ function from __host__ function}} struct B2_with_implicit_default_ctor : A2_with_device_ctor { using A2_with_device_ctor::A2_with_device_ctor; @@ -84,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 552f8f2ebd94fd..d87e6962404341 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 index 21d698d28492ac..34142bcc621200 100644 --- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu +++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu @@ -37,12 +37,13 @@ struct TC : TB { ~TC() {} }; +template class TC; + __device__ TC tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} // Check trivial ctor specialization template -struct C { //expected-note {{candidate constructor (the implicit copy constructor) not viable}} - //expected-note@-1 {{candidate constructor (the implicit move constructor) not viable}} +struct C { explicit C() {}; }; @@ -51,6 +52,6 @@ __device__ C ci_d; C ci_h; // Check non-trivial ctor specialization -template <> C::C() { static int nontrivial_ctor = 1; } //expected-note {{candidate constructor not viable: call to __host__ function from __device__ function}} -__device__ C cf_d; //expected-error {{no matching constructor for initialization of 'C'}} +template <> C::C() { static int nontrivial_ctor = 1; } +__device__ C cf_d; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} C cf_h;