[CUDA] Normalize handling of defauled dtor.

[tra]

Summary

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.

[rsmith]

clang/lib/Sema/SemaDeclCXX.cpp

15162–15170 ↗

(On Diff #316808)

Is this safe? What happens if the destructor for the variable is a template, and instantiating that template results in a reference to a device function? Eg:

template<typename T> __device__ void f() {}
template<typename T> struct A {
  ~A() { f<<<>>>(); }
};
A a;

clang/test/SemaCUDA/default-ctor.cu

28 ↗

(On Diff #316808)

I don't think we should accept this -- only functions that are defaulted on their first declaration should get special treatment. Instead of checking for isDefaulted() above, you should check for !isUserProvided() instead.

[tra]

Addressed Richard's comments.

[tra]

clang/lib/Sema/SemaDeclCXX.cpp

15162–15170 ↗

(On Diff #316808)

This is business as usual -- we catch it during host compilation, where a is instantiated.

h.cu:3:10: error: no matching function for call to 'f'
  ~A() { f<T>(); }
         ^~~~
h.cu:5:8: note: in instantiation of member function 'A<int>::~A' requested here
A<int> a;
       ^
h.cu:1:51: note: candidate function not viable: call to __device__ function from __host__ function
template<typename T> __attribute__((device)) void f() {}

1 error generated when compiling for host.

If it were a __device__ A<int> a; , then we catch it during GPU compilation and also complain that we can't have dynamic initializers.

clang/test/SemaCUDA/default-ctor.cu

28 ↗

(On Diff #316808)

Fixed.

[rsmith]

clang/lib/Sema/SemaDeclCXX.cpp

15162–15170 ↗

(On Diff #316808)

Sorry, testcase wasn't quite right; I meant for f to be __global__ not __device__ so that the kernel call to it works. Fixed example:

extern "C" int cudaConfigureCall(int a, int b);
template<typename T> __attribute__((__global__)) void f(T) {}
template<typename T> struct A {
  ~A() { f<<<1, 1>>>(T()); }
};
A<int> a;

I think that this is valid. In order for it to work, we need to trigger instantiation of f<int> on the device side of the compilation. In order to do that, we need to trigger instantiation of A<int>::~A(), so we need to mark it referenced on the device side. (This is, I think, in line with the general principle that we want to do the same template instantiations of host functions on both sides of the compilation, so that both sides agree on which kernel functions are referenced.)

[tra]

clang/lib/Sema/SemaDeclCXX.cpp

15162–15170 ↗

(On Diff #316808)

You're right. To think of it this particular change is not needed at all any more. The real issue is fixed by the better selection of the usual deallocator. We do not need to skip dtor checks here.

[tra]

Removed unneeded changes.

[tra]

Added a test for the corner case Richard has pointed out in the comments.

clang/lib/Sema/SemaCUDA.cpp

if (D->hasAttr<CUDAGlobalAttr>())
     return CFT_Global;
 
   if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
     if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr))
       return CFT_HostDevice;
     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;
   }
 
   return CFT_Host;
 }
 

clang/lib/Sema/SemaExprCXX.cpp

Sema::BuildCXXTypeConstructExpr(TypeSourceInfo *TInfo,
   }
 
   return Result;
 }
 
 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)
+  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);
 
   if (Result || !getLangOpts().CUDA || PreventedBy.empty())
     return Result;
 
   // In case of CUDA, return true if none of the 1-argument deallocator

clang/test/CodeGenCUDA/usual-deallocators.cu

 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE
 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST
 // RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE
 // RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST
 
 #include "Inputs/cuda.h"
 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(); };
 };
 
 struct H1D2 {
   __host__ void operator delete(void *) { host_fn(); };
   __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); };
__host__ __device__ void tests_hd(void *t) {
   // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* {{.*}}, i32 1)
   // HOST:   call void @_ZN6H1H2D2dlEPv(i8* {{.*}})
   test_hd<H1H2D2>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv
   // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}})
   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)
 // DEVICE: call void @dev_fn()
 // HOST:   call void @host_fn()
 
 // DEVICE-LABEL: define  linkonce_odr void @_ZN4H1D2dlEPvj(i8* %0, i32 %1)
 // DEVICE: call void @dev_fn()

 // DEVICE-LABEL: define  linkonce_odr void @_ZN6H1H2D2dlEPvj(i8* %0, i32 %1)
 // DEVICE: call void @dev_fn()
 // HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(i8* %0)
 // HOST: call void @host_fn()
 
 // 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}

clang/test/SemaCUDA/usual-deallocators.cu

__host__ __device__ void tests_hd(void *t) {
   test_hd<H2D1>(t);
   test_hd<H2D2>(t);
   test_hd<H1D1D2>(t);
   test_hd<H1H2D1>(t);
   test_hd<H1H2D1>(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;