https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/128926
>From 30bde986afda876d877d63212d8e4a34d2927aad Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <yaxun....@amd.com> Date: Wed, 26 Feb 2025 11:43:28 -0500 Subject: [PATCH] [CUDA][HIP] fix virtual dtor host/device attr Currently if CUDA/HIP users use template class with virtual dtor and std::string data member with C++20 and MSVC. When the template class is explicitly instantiated, there is error about host function called by host device function (used to be undefined symbols in linking stage before member destructors were checked by deferred diagnostics). It was caused by clang inferring host/device attributes for default dtors. Since all dtors of member and parent classes have implicit host device attrs, clang infers the virtual dtor have implicit host and device attrs. Since virtual dtor of explicitly instantiated template class must be emitted, this causes constexpr dtor of std::string emitted, which calls a host function which was not emitted on device side. This is a serious issue since it prevents users from using std::string with C++20 on Windows. When inferring host device attr of virtual dtor of explicit template class instantiation, clang should be conservative since it is sure to be emitted. Since an implicit host device function may call a host function, clang cannot assume it is always available on device. This guarantees dtors that may call host functions not to have implicit device attr, therefore will not be emitted on device side. Fixes: https://github.com/llvm/llvm-project/issues/108548 Fixes: SWDEV-517435 --- clang/docs/HIPSupport.rst | 20 ++++++++++++++++++++ clang/lib/Sema/SemaCUDA.cpp | 8 ++++++-- clang/test/SemaCUDA/dtor.cu | 18 ++++++++++-------- 3 files changed, 36 insertions(+), 10 deletions(-) diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst index 481ed39230813..8f473c21e1918 100644 --- a/clang/docs/HIPSupport.rst +++ b/clang/docs/HIPSupport.rst @@ -286,6 +286,26 @@ Example Usage basePtr->virtualFunction(); // Allowed since obj is constructed in device code } +Host and Device Attributes of Default Destructors +=================================================== + +If a default destructor does not have explicit host or device attributes, +clang infers these attributes based on the destructors of its data members +and base classes. If any conflicts are detected among these destructors, +clang diagnoses the issue. Otherwise, clang adds an implicit host or device +attribute according to whether the data members's and base classes's +destructors can execute on the host or device side. + +For explicit template classes with virtual destructors, which must be emitted, +the inference adopts a conservative approach. In this case, implicit host or +device attributes from member and base class destructors are ignored. This +precaution is necessary because, although a constexpr destructor carries +implicit host or device attributes, a constexpr function may call a +non-constexpr function, which is by default a host function. + +Users can override the inferred host and device attributes of default +destructors by adding explicit host and device attributes to them. + C++ Standard Parallelism Offload Support: Compiler And Runtime ============================================================== diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 473956c37bb51..0e5fc5e1a40b4 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -437,7 +437,9 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, if (!SMOR.getMethod()) continue; - CUDAFunctionTarget BaseMethodTarget = IdentifyTarget(SMOR.getMethod()); + CUDAFunctionTarget BaseMethodTarget = + IdentifyTarget(SMOR.getMethod(), IsExpVDtor); + if (!InferredTarget) { InferredTarget = BaseMethodTarget; } else { @@ -481,7 +483,9 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, if (!SMOR.getMethod()) continue; - CUDAFunctionTarget FieldMethodTarget = IdentifyTarget(SMOR.getMethod()); + CUDAFunctionTarget FieldMethodTarget = + IdentifyTarget(SMOR.getMethod(), IsExpVDtor); + if (!InferredTarget) { InferredTarget = FieldMethodTarget; } else { diff --git a/clang/test/SemaCUDA/dtor.cu b/clang/test/SemaCUDA/dtor.cu index c266e51f5c29e..cc37837e70791 100644 --- a/clang/test/SemaCUDA/dtor.cu +++ b/clang/test/SemaCUDA/dtor.cu @@ -32,22 +32,24 @@ public: template class B<float>; } -// The implicit host/device attrs of virtual dtor B<float>::~B() is inferred to -// have implicit device attr since dtors of its members and parent classes can -// be executed on device. This causes a diagnostic since B<float>::~B() must -// be emitted, and it eventually causes host_fun() called on device side. +// The implicit host/device attrs of virtual dtor ~B() should be +// conservatively inferred, where constexpr member dtor's should +// not be considered device since they may call host functions. +// Therefore B<float>::~B() should not have implicit device attr. +// However C<float>::~C() should have implicit device attr since +// it is trivial. namespace ExplicitInstantiationDtorNoAttr { -void host_fun() // dev-note {{'host_fun' declared here}} +void host_fun() {} template <unsigned> constexpr void hd_fun() { - host_fun(); // dev-error{{reference to __host__ function 'host_fun' in __host__ __device__ function}} + host_fun(); } struct A { - constexpr ~A() { // dev-note {{called by '~B'}} - hd_fun<8>(); // dev-note {{called by '~A'}} + constexpr ~A() { + hd_fun<8>(); } }; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
