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 // Destructors are handled a bit differently, compared to regular functions. 16 // Make sure we do trigger kernel generation on the GPU side even if it's only 17 // referenced by the destructor. 18 template<typename T> __global__ void f(T) {} 19 template<typename T> struct A { 20 ~A() { f<<<1, 1>>>(T()); } 21 }; 22 23 // HOST-LABEL: @a 24 A<int> a; 25 // HOST-LABEL: define linkonce_odr void @_ZN1AIiED1Ev 26 // search further down for the deice-side checks for @_Z1fIiEvT_ 27 28 struct H1D1 { 29 __host__ void operator delete(void *) { host_fn(); }; 30 __device__ void operator delete(void *) { dev_fn(); }; 31 }; 32 33 struct H1D2 { 34 __host__ void operator delete(void *) { host_fn(); }; 35 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 36 }; 37 38 struct H2D1 { 39 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 40 __device__ void operator delete(void *) { dev_fn(); }; 41 }; 42 43 struct H2D2 { 44 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 45 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 46 }; 47 48 struct H1D1D2 { 49 __host__ void operator delete(void *) { host_fn(); }; 50 __device__ void operator delete(void *) { dev_fn(); }; 51 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 52 }; 53 54 struct H1H2D1 { 55 __host__ void operator delete(void *) { host_fn(); }; 56 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 57 __device__ void operator delete(void *) { dev_fn(); }; 58 }; 59 60 struct H1H2D2 { 61 __host__ void operator delete(void *) { host_fn(); }; 62 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 63 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 64 }; 65 66 struct H1H2D1D2 { 67 __host__ void operator delete(void *) { host_fn(); }; 68 __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; 69 __device__ void operator delete(void *) { dev_fn(); }; 70 __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; 71 }; 72 73 74 template <typename T> 75 __host__ __device__ void test_hd(void *p) { 76 T *t = (T *)p; 77 delete t; 78 } 79 80 // Make sure we call the right variant of usual deallocator. 81 __host__ __device__ void tests_hd(void *t) { 82 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D1EvPv 83 // COMMON: call void @_ZN4H1D1dlEPv 84 test_hd<H1D1>(t); 85 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv 86 // DEVICE: call void @_ZN4H1D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1) 87 // HOST: call void @_ZN4H1D2dlEPv(i8* noundef {{.*}}) 88 test_hd<H1D2>(t); 89 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv 90 // DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}}) 91 // HOST: call void @_ZN4H2D1dlEPvj(i8* noundef %3, i32 noundef 1) 92 test_hd<H2D1>(t); 93 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv 94 // COMMON: call void @_ZN4H2D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1) 95 test_hd<H2D2>(t); 96 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv 97 // COMMON: call void @_ZN6H1D1D2dlEPv(i8* noundef %3) 98 test_hd<H1D1D2>(t); 99 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv 100 // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}}) 101 test_hd<H1H2D1>(t); 102 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv 103 // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1) 104 // HOST: call void @_ZN6H1H2D2dlEPv(i8* noundef {{.*}}) 105 test_hd<H1H2D2>(t); 106 // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv 107 // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}}) 108 test_hd<H1H2D1D2>(t); 109 } 110 111 // Make sure that we've generated the kernel used by A::~A. 112 // DEVICE-LABEL: define void @_Z1fIiEvT_ 113 114 // Make sure we've picked deallocator for the correct side of compilation. 115 116 // COMMON-LABEL: define linkonce_odr void @_ZN4H1D1dlEPv(i8* noundef %0) 117 // DEVICE: call void @dev_fn() 118 // HOST: call void @host_fn() 119 120 // DEVICE-LABEL: define linkonce_odr void @_ZN4H1D2dlEPvj(i8* noundef %0, i32 noundef %1) 121 // DEVICE: call void @dev_fn() 122 // HOST-LABEL: define linkonce_odr void @_ZN4H1D2dlEPv(i8* noundef %0) 123 // HOST: call void @host_fn() 124 125 // DEVICE-LABEL: define linkonce_odr void @_ZN4H2D1dlEPv(i8* noundef %0) 126 // DEVICE: call void @dev_fn() 127 // HOST-LABEL: define linkonce_odr void @_ZN4H2D1dlEPvj(i8* noundef %0, i32 noundef %1) 128 // HOST: call void @host_fn() 129 130 // COMMON-LABEL: define linkonce_odr void @_ZN4H2D2dlEPvj(i8* noundef %0, i32 noundef %1) 131 // DEVICE: call void @dev_fn() 132 // HOST: call void @host_fn() 133 134 // COMMON-LABEL: define linkonce_odr void @_ZN6H1D1D2dlEPv(i8* noundef %0) 135 // DEVICE: call void @dev_fn() 136 // HOST: call void @host_fn() 137 138 // COMMON-LABEL: define linkonce_odr void @_ZN6H1H2D1dlEPv(i8* noundef %0) 139 // DEVICE: call void @dev_fn() 140 // HOST: call void @host_fn() 141 142 // DEVICE-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPvj(i8* noundef %0, i32 noundef %1) 143 // DEVICE: call void @dev_fn() 144 // HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(i8* noundef %0) 145 // HOST: call void @host_fn() 146 147 // COMMON-LABEL: define linkonce_odr void @_ZN8H1H2D1D2dlEPv(i8* noundef %0) 148 // DEVICE: call void @dev_fn() 149 // HOST: call void @host_fn() 150 151 // DEVICE: !0 = !{void (i32)* @_Z1fIiEvT_, !"kernel", i32 1} 152