Aligned allocation versus CUDA: make deallocation function preference order

match other CUDA preference orders, per discussion with jlebar. We now model
this in an attempt to match overload resolution as closely as possible:

- First, we throw out all non-callable (due to CUDA host/device mismatch)
  operator delete functions.
- Then we apply sizedness / alignedness preferences based on whether the type
  is overaligned and whether the deallocation function is a member.
- Finally, we use the CUDA callability preference as a tiebreaker.

llvm-svn: 283830

1 files changed, 13 insertions, 0 deletions

diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
index 782c13d5445..b83beba58db 100644
--- a/clang/test/SemaCUDA/call-host-fn-from-device.cu
+++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -46,6 +46,14 @@ struct T {
 
   operator Dummy() { return Dummy(); }
   // expected-note@-1 {{'operator Dummy' declared here}}
+
+  __host__ void operator delete(void*);
+  __device__ void operator delete(void*, size_t);
+};
+
+struct U {
+  __device__ void operator delete(void*, size_t) = delete;
+  __host__ __device__ void operator delete(void*);
 };
 
 __host__ __device__ void T::hd3() {
@@ -82,6 +90,11 @@ __host__ __device__ void explicit_destructor(S *s) {
   // expected-error@-1 {{reference to __host__ function '~S' in __host__ __device__ function}}
 }
 
+__host__ __device__ void class_specific_delete(T *t, U *u) {
+  delete t; // ok, call sized device delete even though host has preferable non-sized version
+  delete u; // ok, call non-sized HD delete rather than sized D delete
+}
+
 __host__ __device__ void hd_member_fn() {
   T t;
   // Necessary to trigger an error on T::hd.  It's (implicitly) inline, so