Author: Yaxun (Sam) Liu Date: 2023-12-01T16:24:01-05:00 New Revision: 2b76e20ea782790a78ec58d5f94ce88a173bab7f
URL: https://github.com/llvm/llvm-project/commit/2b76e20ea782790a78ec58d5f94ce88a173bab7f DIFF: https://github.com/llvm/llvm-project/commit/2b76e20ea782790a78ec58d5f94ce88a173bab7f.diff LOG: [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 Added: clang/test/SemaCUDA/trivial-ctor-dtor.cu Modified: clang/lib/Sema/SemaCUDA.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 318174f7be8fa95..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); diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu new file mode 100644 index 000000000000000..34142bcc621200f --- /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; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
