74 lines · plain
1// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -emit-llvm -fexceptions \2// RUN: -o - -x hip %s | FileCheck %s3 4#include "Inputs/cuda.h"5 6int* hvar;7__device__ int* dvar;8 9// CHECK-LABEL: define {{.*}}@_Znwm10// CHECK: load ptr, ptr @hvar11void* operator new(unsigned long size) {12 return hvar;13}14// CHECK-LABEL: define {{.*}}@_ZdlPv15// CHECK: store ptr inttoptr (i64 1 to ptr), ptr @hvar16void operator delete(void *p) {17 hvar = (int*)1;18}19 20__device__ void* operator new(unsigned long size) {21 return dvar;22}23 24__device__ void operator delete(void *p) {25 dvar = (int*)11;26}27 28class A {29 int x;30public:31 A(){32 x = 123;33 }34};35 36template<class T>37class shared_ptr {38 int id;39 T *ptr;40public:41 shared_ptr(T *p) {42 id = 2;43 ptr = p;44 }45};46 47// The constructor of B calls the delete operator to clean up48// the memory allocated by the new operator when exceptions happen.49// Make sure the host delete operator is used on host side.50//51// No need to do similar checks on the device side since it does52// not support exception.53 54// CHECK-LABEL: define {{.*}}@main55// CHECK: call void @_ZN1BC1Ev56 57// CHECK-LABEL: define {{.*}}@_ZN1BC1Ev58// CHECK: call void @_ZN1BC2Ev59 60// CHECK-LABEL: define {{.*}}@_ZN1BC2Ev61// CHECK: call {{.*}}@_Znwm62// CHECK: invoke void @_ZN1AC1Ev63// CHECK: call void @_ZN10shared_ptrI1AEC1EPS0_64// CHECK: cleanup65// CHECK: call void @_ZdlPv66 67struct B{68 shared_ptr<A> pa{new A};69};70 71int main() {72 B b;73}74