1 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ 2 // RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE 3 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \ 4 // RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST 5 // RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \ 6 // RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE 7 // RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \ 8 // RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST 9 10 #include "Inputs/cuda.h" 11 extern "C" __host__ void host_fn(); 12 extern "C" __device__ void dev_fn(); 13 extern "C" __host__ __device__ void hd_fn(); 14 15 struct H1D1 { 16 __host__ void operator delete(void *) { host_fn(); }; 17 __device__ void operator delete(void *) { dev_fn(); }; 18 }; 19 20 struct H1D2 { 21 __host__ void operator delete(void *) { host_fn(); }; 22 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 23 }; 24 25 struct H2D1 { 26 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 27 __device__ void operator delete(void *) { dev_fn(); }; 28 }; 29 30 struct H2D2 { 31 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 32 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 33 }; 34 35 struct H1D1D2 { 36 __host__ void operator delete(void *) { host_fn(); }; 37 __device__ void operator delete(void *) { dev_fn(); }; 38 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 39 }; 40 41 struct H1H2D1 { 42 __host__ void operator delete(void *) { host_fn(); }; 43 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 44 __device__ void operator delete(void *) { dev_fn(); }; 45 }; 46 47 struct H1H2D2 { 48 __host__ void operator delete(void *) { host_fn(); }; 49 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 50 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 51 }; 52 53 struct H1H2D1D2 { 54 __host__ void operator delete(void *) { host_fn(); }; 55 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 56 __device__ void operator delete(void *) { dev_fn(); }; 57 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 58 }; 59 60 61 template <typename T> 62 __host__ __device__ void test_hd(void *p) { 63 T *t = (T *)p; 64 delete t; 65 } 66 67 // Make sure we call the right variant of usual deallocator. 68 __host__ __device__ void tests_hd(void *t) { 69 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D1EvPv 70 // COMMON: call void @_ZN4H1D1dlEPv 71 test_hd<H1D1>(t); 72 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv 73 // DEVICE: call void @_ZN4H1D2dlEPvj(i8* {{.*}}, i32 1) 74 // HOST: call void @_ZN4H1D2dlEPv(i8* {{.*}}) 75 test_hd<H1D2>(t); 76 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv 77 // DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}}) 78 // HOST: call void @_ZN4H2D1dlEPvj(i8* %3, i32 1) 79 test_hd<H2D1>(t); 80 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv 81 // COMMON: call void @_ZN4H2D2dlEPvj(i8* {{.*}}, i32 1) 82 test_hd<H2D2>(t); 83 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv 84 // COMMON: call void @_ZN6H1D1D2dlEPv(i8* %3) 85 test_hd<H1D1D2>(t); 86 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv 87 // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}}) 88 test_hd<H1H2D1>(t); 89 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv 90 // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* {{.*}}, i32 1) 91 // HOST: call void @_ZN6H1H2D2dlEPv(i8* {{.*}}) 92 test_hd<H1H2D2>(t); 93 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv 94 // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}}) 95 test_hd<H1H2D1D2>(t); 96 } 97 98 // Make sure we've picked deallocator for the correct side of compilation. 99 100 // COMMON-LABEL: define linkonce_odr void @_ZN4H1D1dlEPv(i8* %0) 101 // DEVICE: call void @dev_fn() 102 // HOST: call void @host_fn() 103 104 // DEVICE-LABEL: define linkonce_odr void @_ZN4H1D2dlEPvj(i8* %0, i32 %1) 105 // DEVICE: call void @dev_fn() 106 // HOST-LABEL: define linkonce_odr void @_ZN4H1D2dlEPv(i8* %0) 107 // HOST: call void @host_fn() 108 109 // DEVICE-LABEL: define linkonce_odr void @_ZN4H2D1dlEPv(i8* %0) 110 // DEVICE: call void @dev_fn() 111 // HOST-LABEL: define linkonce_odr void @_ZN4H2D1dlEPvj(i8* %0, i32 %1) 112 // HOST: call void @host_fn() 113 114 // COMMON-LABEL: define linkonce_odr void @_ZN4H2D2dlEPvj(i8* %0, i32 %1) 115 // DEVICE: call void @dev_fn() 116 // HOST: call void @host_fn() 117 118 // COMMON-LABEL: define linkonce_odr void @_ZN6H1D1D2dlEPv(i8* %0) 119 // DEVICE: call void @dev_fn() 120 // HOST: call void @host_fn() 121 122 // COMMON-LABEL: define linkonce_odr void @_ZN6H1H2D1dlEPv(i8* %0) 123 // DEVICE: call void @dev_fn() 124 // HOST: call void @host_fn() 125 126 // DEVICE-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPvj(i8* %0, i32 %1) 127 // DEVICE: call void @dev_fn() 128 // HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(i8* %0) 129 // HOST: call void @host_fn() 130 131 // COMMON-LABEL: define linkonce_odr void @_ZN8H1H2D1D2dlEPv(i8* %0) 132 // DEVICE: call void @dev_fn() 133 // HOST: call void @host_fn() 134