[CUDA][HIP] Fix overloading resolution of delete operator

[yaxunl]

Summary

Currently clang does not consider host/device preference when resolving delete operator in the file scope, which causes device operator delete selected for class member initialization.

clang/lib/Sema/SemaExprCXX.cpp

UsualDeallocFnInfo(Sema &S, DeclAccessPair Found)
       if (NumBaseParams < FD->getNumParams() &&
           FD->getParamDecl(NumBaseParams)->getType()->isAlignValT()) {
         ++NumBaseParams;
         HasAlignValT = true;
       }
 
       // In CUDA, determine how much we'd like / dislike to call this.
       if (S.getLangOpts().CUDA)
-        if (auto *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true))
-          CUDAPref = S.IdentifyCUDAPreference(Caller, FD);
+        CUDAPref = S.IdentifyCUDAPreference(
+            S.getCurFunctionDecl(/*AllowLambda=*/true), FD);
     }
 
     explicit operator bool() const { return FD; }
 
     bool isBetterThan(const UsualDeallocFnInfo &Other, bool WantSize,
                       bool WantAlign) const {
       // C++ P0722:
       //   A destroying operator delete is preferred over a non-destroying

clang/test/CodeGenCUDA/member-init.cu

+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -emit-llvm -fexceptions \
+// RUN:   -o - -x hip %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+int* hvar;
+__device__ int* dvar;
+
+// CHECK-LABEL: define {{.*}}@_Znwm
+// CHECK:    load ptr, ptr @hvar
+void* operator new(unsigned long size) {
+  return hvar;
+}
+// CHECK-LABEL: define {{.*}}@_ZdlPv
+// CHECK:    store ptr inttoptr (i64 1 to ptr), ptr @hvar
+void operator delete(void *p) {
+  hvar = (int*)1;
+}
+
+__device__ void* operator new(unsigned long size) {
+  return dvar;
+}
+
+__device__ void operator delete(void *p) {
+  dvar = (int*)11;
+}
+
+class A {
+  int x;
+public:
+  A(){
+    x = 123;
+  }
+};
+
+template<class T>
+class shared_ptr {
+   int id;
+   T *ptr;
+public:
+  shared_ptr(T *p) {
+    id = 2;
+    ptr = p;
+  }
+};
+
+// The constructor of B calls the delete operator to clean up
+// the memory allocated by the new operator when exceptions happen.
+// Make sure the host delete operator is used on host side.
+//
+// No need to do similar checks on the device side since it does
+// not support exception.
+
+// CHECK-LABEL: define {{.*}}@main
+// CHECK:    call void @_ZN1BC1Ev
+
+// CHECK-LABEL: define {{.*}}@_ZN1BC1Ev
+// CHECK:    call void @_ZN1BC2Ev
+
+// CHECK-LABEL: define {{.*}}@_ZN1BC2Ev
+// CHECK: call {{.*}}@_Znwm
+// CHECK:  invoke void @_ZN1AC1Ev
+// CHECK:  call void @_ZN10shared_ptrI1AEC1EPS0_
+// CHECK:  cleanup
+// CHECK:  call void @_ZdlPv
+
+struct B{
+  shared_ptr<A> pa{new A};
+};
+
+int main() {
+  B b;
+}

clang/test/SemaCUDA/member-init.cu

+// RUN: %clang_cc1 -fsyntax-only -verify -fexceptions %s
+// expected-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+__device__ void operator delete(void *p) {}
+
+class A {
+  int x;
+public:
+  A() {
+  x = 123;
+  }
+};
+
+template<class T>
+class shared_ptr {
+  T *ptr;
+public:
+  shared_ptr(T *p) {
+    ptr = p;
+  }
+};
+
+// The constructor of B calls the delete operator to clean up
+// the memory allocated by the new operator when exceptions happen.
+// Make sure that there are no diagnostics due to the device delete
+// operator is used.
+//
+// No need to do similar checks on the device side since it does
+// not support exception.
+struct B{
+  shared_ptr<A> pa{new A};
+};
+
+int main() {
+  B b;
+}