| commit 127091bfd5edf10495fee4724fd21c666e5d79c1 |
| Author: Artem Belevich <tra@google.com> |
| Date: Thu Jan 14 16:05:33 2021 -0800 |
| |
| [CUDA] Normalize handling of defauled dtor. |
| |
| Defaulted destructor was treated inconsistently, compared to other |
| compiler-generated functions. |
| |
| When Sema::IdentifyCUDATarget() got called on just-created dtor which didn't |
| have implicit __host__ __device__ attributes applied yet, it would treat it as a |
| host function. That happened to (sometimes) hide the error when dtor referred |
| to a host-only functions. |
| |
| Even when we had identified defaulted dtor as a HD function, we still treated it |
| inconsistently during selection of usual deallocators, where we did not allow |
| referring to wrong-side functions, while it is allowed for other HD functions. |
| |
| This change brings handling of defaulted dtors in line with other HD functions. |
| |
| Differential Revision: https://reviews.llvm.org/D94732 |
| |
| diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp |
| index 0f06adf38f7a..ee91eb4c5deb 100644 |
| --- a/clang/lib/Sema/SemaCUDA.cpp |
| +++ b/clang/lib/Sema/SemaCUDA.cpp |
| @@ -123,7 +123,8 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D, |
| return CFT_Device; |
| } else if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr)) { |
| return CFT_Host; |
| - } else if (D->isImplicit() && !IgnoreImplicitHDAttr) { |
| + } else if ((D->isImplicit() || !D->isUserProvided()) && |
| + !IgnoreImplicitHDAttr) { |
| // Some implicit declarations (like intrinsic functions) are not marked. |
| // Set the most lenient target on them for maximal flexibility. |
| return CFT_HostDevice; |
| diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp |
| index 1ee52107c3da..d91db60f17a0 100644 |
| --- a/clang/lib/Sema/SemaExprCXX.cpp |
| +++ b/clang/lib/Sema/SemaExprCXX.cpp |
| @@ -1527,9 +1527,24 @@ Sema::BuildCXXTypeConstructExpr(TypeSourceInfo *TInfo, |
| bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) { |
| // [CUDA] Ignore this function, if we can't call it. |
| const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext); |
| - if (getLangOpts().CUDA && |
| - IdentifyCUDAPreference(Caller, Method) <= CFP_WrongSide) |
| - return false; |
| + if (getLangOpts().CUDA) { |
| + auto CallPreference = IdentifyCUDAPreference(Caller, Method); |
| + // If it's not callable at all, it's not the right function. |
| + if (CallPreference < CFP_WrongSide) |
| + return false; |
| + if (CallPreference == CFP_WrongSide) { |
| + // Maybe. We have to check if there are better alternatives. |
| + DeclContext::lookup_result R = |
| + Method->getDeclContext()->lookup(Method->getDeclName()); |
| + for (const auto *D : R) { |
| + if (const auto *FD = dyn_cast<FunctionDecl>(D)) { |
| + if (IdentifyCUDAPreference(Caller, FD) > CFP_WrongSide) |
| + return false; |
| + } |
| + } |
| + // We've found no better variants. |
| + } |
| + } |
| |
| SmallVector<const FunctionDecl*, 4> PreventedBy; |
| bool Result = Method->isUsualDeallocationFunction(PreventedBy); |
| diff --git a/clang/test/CodeGenCUDA/usual-deallocators.cu b/clang/test/CodeGenCUDA/usual-deallocators.cu |
| index 7e7752497f34..6f4cc267a23f 100644 |
| --- a/clang/test/CodeGenCUDA/usual-deallocators.cu |
| +++ b/clang/test/CodeGenCUDA/usual-deallocators.cu |
| @@ -12,6 +12,19 @@ extern "C" __host__ void host_fn(); |
| extern "C" __device__ void dev_fn(); |
| extern "C" __host__ __device__ void hd_fn(); |
| |
| +// Destructors are handled a bit differently, compared to regular functions. |
| +// Make sure we do trigger kernel generation on the GPU side even if it's only |
| +// referenced by the destructor. |
| +template<typename T> __global__ void f(T) {} |
| +template<typename T> struct A { |
| + ~A() { f<<<1, 1>>>(T()); } |
| +}; |
| + |
| +// HOST-LABEL: @a |
| +A<int> a; |
| +// HOST-LABEL: define linkonce_odr void @_ZN1AIiED1Ev |
| +// search further down for the deice-side checks for @_Z1fIiEvT_ |
| + |
| struct H1D1 { |
| __host__ void operator delete(void *) { host_fn(); }; |
| __device__ void operator delete(void *) { dev_fn(); }; |
| @@ -95,6 +108,9 @@ __host__ __device__ void tests_hd(void *t) { |
| test_hd<H1H2D1D2>(t); |
| } |
| |
| +// Make sure that we've generated the kernel used by A::~A. |
| +// DEVICE-LABEL: define dso_local void @_Z1fIiEvT_ |
| + |
| // Make sure we've picked deallocator for the correct side of compilation. |
| |
| // COMMON-LABEL: define linkonce_odr void @_ZN4H1D1dlEPv(i8* %0) |
| @@ -131,3 +147,5 @@ __host__ __device__ void tests_hd(void *t) { |
| // COMMON-LABEL: define linkonce_odr void @_ZN8H1H2D1D2dlEPv(i8* %0) |
| // DEVICE: call void @dev_fn() |
| // HOST: call void @host_fn() |
| + |
| +// DEVICE: !0 = !{void (i32)* @_Z1fIiEvT_, !"kernel", i32 1} |
| diff --git a/clang/test/SemaCUDA/usual-deallocators.cu b/clang/test/SemaCUDA/usual-deallocators.cu |
| index a0238649c6dc..3670a3bf32c3 100644 |
| --- a/clang/test/SemaCUDA/usual-deallocators.cu |
| +++ b/clang/test/SemaCUDA/usual-deallocators.cu |
| @@ -93,3 +93,12 @@ __host__ __device__ void tests_hd(void *t) { |
| test_hd<H1H2D2>(t); |
| test_hd<H1H2D1D2>(t); |
| } |
| + |
| +// This should produce no errors. Defaulted destructor should be treated as HD, |
| +// which allows referencing host-only `operator delete` with a deferred |
| +// diagnostics that would fire if we ever attempt to codegen it on device.. |
| +struct H { |
| + virtual ~H() = default; |
| + static void operator delete(void *) {} |
| +}; |
| +H h; |