llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Yaxun (Sam) Liu (yxsamliu) <details> <summary>Changes</summary> 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. Fixes: #<!-- -->72261 Fixes: SWDEV-432412 --- Full diff: https://github.com/llvm/llvm-project/pull/73140.diff 11 Files Affected: - (modified) clang/include/clang/Sema/Sema.h (-4) - (modified) clang/lib/Sema/SemaCUDA.cpp (+9-16) - (modified) clang/lib/Sema/SemaDecl.cpp (-3) - (modified) clang/lib/Sema/SemaOverload.cpp (+2-4) - (modified) clang/test/SemaCUDA/call-host-fn-from-device.cu (+1-1) - (modified) clang/test/SemaCUDA/default-ctor.cu (+1-1) - (modified) clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu (+1-1) - (modified) clang/test/SemaCUDA/implicit-member-target-collision.cu (+1-1) - (modified) clang/test/SemaCUDA/implicit-member-target-inherited.cu (+2-3) - (modified) clang/test/SemaCUDA/implicit-member-target.cu (+2-2) - (modified) clang/test/SemaCUDA/trivial-ctor-dtor.cu (+5-4) ``````````diff 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..6a66ecf6f94c178 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<CXXConstructorDecl>(Callee) || isa<CXXDestructorDecl>(Callee))) + return CFP_HostDevice; + CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller); CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee); @@ -772,22 +781,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/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index 64607e28b8b35e6..9800d7f1c9cfee9 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -1491,10 +1491,8 @@ static bool IsOverloadOrOverrideImpl(Sema &SemaRef, FunctionDecl *New, // Don't allow overloading of destructors. (In theory we could, but it // would be a giant change to clang.) if (!isa<CXXDestructorDecl>(New)) { - Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget( - New, isa<CXXConstructorDecl>(New)), - OldTarget = SemaRef.IdentifyCUDATarget( - Old, isa<CXXConstructorDecl>(New)); + Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(New), + OldTarget = SemaRef.IdentifyCUDATarget(Old); if (NewTarget != Sema::CFT_InvalidTarget) { assert((OldTarget != Sema::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 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 ceca0891fc9b03c..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}} @@ -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 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 index 21d698d28492ac3..34142bcc621200f 100644 --- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu +++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu @@ -37,12 +37,13 @@ struct TC : TB<T> { ~TC() {} }; +template class TC<int>; + __device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} // Check trivial ctor specialization template <typename T> -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<int> ci_d; C<int> ci_h; // Check non-trivial ctor specialization -template <> C<float>::C() { static int nontrivial_ctor = 1; } //expected-note {{candidate constructor not viable: call to __host__ function from __device__ function}} -__device__ C<float> cf_d; //expected-error {{no matching constructor for initialization of 'C<float>'}} +template <> C<float>::C() { static int nontrivial_ctor = 1; } +__device__ C<float> cf_d; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} C<float> cf_h; `````````` </details> https://github.com/llvm/llvm-project/pull/73140 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
