1 // REQUIRES: x86-registered-target 2 // REQUIRES: nvptx-registered-target 3 4 // Make sure we handle target overloads correctly. 5 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ 6 // RUN: -fcuda-target-overloads -emit-llvm -o - %s \ 7 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s 8 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \ 9 // RUN: -fcuda-target-overloads -emit-llvm -o - %s \ 10 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s 11 12 // Check target overloads handling with disabled call target checks. 13 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \ 14 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \ 15 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \ 16 // RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s 17 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \ 18 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \ 19 // RUN: -fcuda-is-device -o - %s \ 20 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \ 21 // RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s 22 23 #include "Inputs/cuda.h" 24 25 typedef int (*fp_t)(void); 26 typedef void (*gp_t)(void); 27 28 // CHECK-HOST: @hp = global i32 ()* @_Z1hv 29 // CHECK-HOST: @chp = global i32 ()* @ch 30 // CHECK-HOST: @dhp = global i32 ()* @_Z2dhv 31 // CHECK-HOST: @cdhp = global i32 ()* @cdh 32 // CHECK-HOST: @gp = global void ()* @_Z1gv 33 34 // CHECK-BOTH-LABEL: define i32 @_Z2dhv() 35 __device__ int dh(void) { return 1; } 36 // CHECK-DEVICE: ret i32 1 37 __host__ int dh(void) { return 2; } 38 // CHECK-HOST: ret i32 2 39 40 // CHECK-BOTH-LABEL: define i32 @_Z2hdv() 41 __host__ __device__ int hd(void) { return 3; } 42 // CHECK-BOTH: ret i32 3 43 44 // CHECK-DEVICE-LABEL: define i32 @_Z1dv() 45 __device__ int d(void) { return 8; } 46 // CHECK-DEVICE: ret i32 8 47 48 // CHECK-HOST-LABEL: define i32 @_Z1hv() 49 __host__ int h(void) { return 9; } 50 // CHECK-HOST: ret i32 9 51 52 // CHECK-BOTH-LABEL: define void @_Z1gv() 53 __global__ void g(void) {} 54 // CHECK-BOTH: ret void 55 56 // mangled names of extern "C" __host__ __device__ functions clash 57 // with those of their __host__/__device__ counterparts, so 58 // overloading of extern "C" functions can only happen for __host__ 59 // and __device__ functions -- we never codegen them in the same 60 // compilation and therefore mangled name conflict is not a problem. 61 62 // CHECK-BOTH-LABEL: define i32 @cdh() 63 extern "C" __device__ int cdh(void) {return 10;} 64 // CHECK-DEVICE: ret i32 10 65 extern "C" __host__ int cdh(void) {return 11;} 66 // CHECK-HOST: ret i32 11 67 68 // CHECK-DEVICE-LABEL: define i32 @cd() 69 extern "C" __device__ int cd(void) {return 12;} 70 // CHECK-DEVICE: ret i32 12 71 72 // CHECK-HOST-LABEL: define i32 @ch() 73 extern "C" __host__ int ch(void) {return 13;} 74 // CHECK-HOST: ret i32 13 75 76 // CHECK-BOTH-LABEL: define i32 @chd() 77 extern "C" __host__ __device__ int chd(void) {return 14;} 78 // CHECK-BOTH: ret i32 14 79 80 // CHECK-HOST-LABEL: define void @_Z5hostfv() 81 __host__ void hostf(void) { 82 #if defined (NOCHECKS) 83 fp_t dp = d; // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp, 84 fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp, 85 #endif 86 fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp, 87 fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp, 88 fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp, 89 fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp, 90 fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp, 91 fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp, 92 gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp, 93 94 #if defined (NOCHECKS) 95 d(); // CHECK-HOST-NC: call i32 @_Z1dv() 96 cd(); // CHECK-HOST-NC: call i32 @cd() 97 #endif 98 h(); // CHECK-HOST: call i32 @_Z1hv() 99 ch(); // CHECK-HOST: call i32 @ch() 100 dh(); // CHECK-HOST: call i32 @_Z2dhv() 101 cdh(); // CHECK-HOST: call i32 @cdh() 102 g<<<0,0>>>(); // CHECK-HOST: call void @_Z1gv() 103 } 104 105 // CHECK-DEVICE-LABEL: define void @_Z7devicefv() 106 __device__ void devicef(void) { 107 fp_t dp = d; // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp, 108 fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp, 109 #if defined (NOCHECKS) 110 fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp, 111 fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp, 112 #endif 113 fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp, 114 fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp, 115 fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp, 116 fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp, 117 118 d(); // CHECK-DEVICE: call i32 @_Z1dv() 119 cd(); // CHECK-DEVICE: call i32 @cd() 120 #if defined (NOCHECKS) 121 h(); // CHECK-DEVICE-NC: call i32 @_Z1hv() 122 ch(); // CHECK-DEVICE-NC: call i32 @ch() 123 #endif 124 dh(); // CHECK-DEVICE: call i32 @_Z2dhv() 125 cdh(); // CHECK-DEVICE: call i32 @cdh() 126 } 127 128 // CHECK-BOTH-LABEL: define void @_Z11hostdevicefv() 129 __host__ __device__ void hostdevicef(void) { 130 #if defined (NOCHECKS) 131 fp_t dp = d; // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp, 132 fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp, 133 fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp, 134 fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp, 135 #endif 136 fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp, 137 fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp, 138 fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp, 139 fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp, 140 #if defined (NOCHECKS) && !defined(__CUDA_ARCH__) 141 gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp, 142 #endif 143 144 #if defined (NOCHECKS) 145 d(); // CHECK-BOTH-NC: call i32 @_Z1dv() 146 cd(); // CHECK-BOTH-NC: call i32 @cd() 147 h(); // CHECK-BOTH-NC: call i32 @_Z1hv() 148 ch(); // CHECK-BOTH-NC: call i32 @ch() 149 #endif 150 dh(); // CHECK-BOTH: call i32 @_Z2dhv() 151 cdh(); // CHECK-BOTH: call i32 @cdh() 152 #if defined (NOCHECKS) && !defined(__CUDA_ARCH__) 153 g<<<0,0>>>(); // CHECK-HOST-NC: call void @_Z1gv() 154 #endif 155 } 156 157 // Test for address of overloaded function resolution in the global context. 158 fp_t hp = h; 159 fp_t chp = ch; 160 fp_t dhp = dh; 161 fp_t cdhp = cdh; 162 gp_t gp = g; 163 164 int x; 165 // Check constructors/destructors for D/H functions 166 struct s_cd_dh { 167 __host__ s_cd_dh() { x = 11; } 168 __device__ s_cd_dh() { x = 12; } 169 __host__ ~s_cd_dh() { x = 21; } 170 __device__ ~s_cd_dh() { x = 22; } 171 }; 172 173 struct s_cd_hd { 174 __host__ __device__ s_cd_hd() { x = 31; } 175 __host__ __device__ ~s_cd_hd() { x = 32; } 176 }; 177 178 // CHECK-BOTH: define void @_Z7wrapperv 179 #if defined(__CUDA_ARCH__) 180 __device__ 181 #else 182 __host__ 183 #endif 184 void wrapper() { 185 s_cd_dh scddh; 186 // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev( 187 s_cd_hd scdhd; 188 // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev 189 190 // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev( 191 // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev( 192 } 193 // CHECK-BOTH: ret void 194 195 // Now it's time to check what's been generated for the methods we used. 196 197 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev( 198 // CHECK-HOST: store i32 11, 199 // CHECK-DEVICE: store i32 12, 200 // CHECK-BOTH: ret void 201 202 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev( 203 // CHECK-BOTH: store i32 31, 204 // CHECK-BOTH: ret void 205 206 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev( 207 // CHECK-BOTH: store i32 32, 208 // CHECK-BOTH: ret void 209 210 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev( 211 // CHECK-HOST: store i32 21, 212 // CHECK-DEVICE: store i32 22, 213 // CHECK-BOTH: ret void 214 215