diff options
author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2023-12-01 16:24:01 -0500 |
---|---|---|
committer | GitHub <noreply@github.com> | 2023-12-01 16:24:01 -0500 |
commit | 2b76e20ea782790a78ec58d5f94ce88a173bab7f (patch) | |
tree | 09e1439adf132e88e939c19b885af116b98e46cf | |
parent | a4d85490e029e797d22881b37d476c2050d0d6a2 (diff) |
[CUDA][HIP] allow trivial ctor/dtor in device var init (#73140)
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
-rw-r--r-- | clang/lib/Sema/SemaCUDA.cpp | 9 | ||||
-rw-r--r-- | clang/test/SemaCUDA/trivial-ctor-dtor.cu | 57 |
2 files changed, 66 insertions, 0 deletions
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 318174f7be8f..6a66ecf6f94c 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); diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu new file mode 100644 index 000000000000..34142bcc6212 --- /dev/null +++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu @@ -0,0 +1,57 @@ +// 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() {} +}; + +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 { + explicit C() {}; +}; + +template <> C<int>::C() {}; +__device__ C<int> ci_d; +C<int> ci_h; + +// Check non-trivial ctor specialization +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; |