// 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;
}