diff options
author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2023-11-16 08:42:54 -0500 |
---|---|---|
committer | GitHub <noreply@github.com> | 2023-11-16 08:42:54 -0500 |
commit | 27e6e4a4d0e3296cebad8db577ec0469a286795e (patch) | |
tree | 49f7b49ab9cca163c827277f560660d1bcc6e19d /clang | |
parent | ea84897ba3e7727a3aa3fbd6d84b6b4ab573c70d (diff) | |
download | llvm-27e6e4a4d0e3296cebad8db577ec0469a286795e.zip llvm-27e6e4a4d0e3296cebad8db577ec0469a286795e.tar.gz llvm-27e6e4a4d0e3296cebad8db577ec0469a286795e.tar.bz2 |
[CUDA][HIP] make trivial ctor/dtor host device (#72394)
Make trivial ctor/dtor implicitly host device functions so that they can
be used to initialize file-scope
device variables to match nvcc behavior.
Fixes: https://github.com/llvm/llvm-project/issues/72261
Fixes: SWDEV-432412
Diffstat (limited to 'clang')
-rw-r--r-- | clang/include/clang/Sema/Sema.h | 4 | ||||
-rw-r--r-- | clang/lib/Sema/SemaCUDA.cpp | 16 | ||||
-rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 3 | ||||
-rw-r--r-- | clang/test/SemaCUDA/call-host-fn-from-device.cu | 2 | ||||
-rw-r--r-- | clang/test/SemaCUDA/default-ctor.cu | 2 | ||||
-rw-r--r-- | clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu | 2 | ||||
-rw-r--r-- | clang/test/SemaCUDA/implicit-member-target-collision.cu | 2 | ||||
-rw-r--r-- | clang/test/SemaCUDA/implicit-member-target-inherited.cu | 4 | ||||
-rw-r--r-- | clang/test/SemaCUDA/implicit-member-target.cu | 4 | ||||
-rw-r--r-- | clang/test/SemaCUDA/trivial-ctor-dtor.cu | 40 |
10 files changed, 71 insertions, 8 deletions
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 731d088..501dc01 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13453,6 +13453,10 @@ public: 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 318174f..b94f448 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -772,6 +772,22 @@ 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 3876eb5..a6cd0bb 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -16232,6 +16232,9 @@ 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 acdd291..b62de92 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() {} + S() { static int nontrivial_ctor = 1; } // 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 cbad7a1..31971fe 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; // expected-error{{no matching constructor for initialization of 'Out'}} + Out 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 06015ed..edb543f 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() {} + A1_with_host_ctor() { static int nontrivial_ctor = 1; } }; 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 a50fdda..16b5978 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() {} + A1_with_host_ctor() { static int nontrivial_ctor = 1; } }; 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 2178172..781199b 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() {} + A1_with_host_ctor() { static int nontrivial_ctor = 1; } }; // 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() {} + A4_with_host_ctor() { static int nontrivial_ctor = 1; } }; 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 d87e696..552f8f2 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() {} + A1_with_host_ctor() { static int nontrivial_ctor = 1; } }; // 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() {} + A4_with_host_ctor() { static int nontrivial_ctor = 1; } }; struct B4_with_implicit_default_ctor { diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu new file mode 100644 index 0000000..1df8adc --- /dev/null +++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu @@ -0,0 +1,40 @@ +// 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}} |