1 // REQUIRES: x86-registered-target 2 // REQUIRES: nvptx-registered-target 3 4 // Make sure we handle target overloads correctly. Most of this is checked in 5 // sema, but special functions like constructors and destructors are here. 6 // 7 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \ 8 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s 9 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \ 10 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s 11 12 #include "Inputs/cuda.h" 13 14 // Check constructors/destructors for D/H functions 15 int x; 16 struct s_cd_dh { 17 __host__ s_cd_dh() { x = 11; } 18 __device__ s_cd_dh() { x = 12; } 19 __host__ ~s_cd_dh() { x = 21; } 20 __device__ ~s_cd_dh() { x = 22; } 21 }; 22 23 struct s_cd_hd { 24 __host__ __device__ s_cd_hd() { x = 31; } 25 __host__ __device__ ~s_cd_hd() { x = 32; } 26 }; 27 28 // CHECK-BOTH: define void @_Z7wrapperv 29 #if defined(__CUDA_ARCH__) 30 __device__ 31 #else 32 __host__ 33 #endif 34 void wrapper() { 35 s_cd_dh scddh; 36 // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev( 37 s_cd_hd scdhd; 38 // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev 39 40 // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev( 41 // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev( 42 } 43 // CHECK-BOTH: ret void 44 45 // Now it's time to check what's been generated for the methods we used. 46 47 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev( 48 // CHECK-HOST: store i32 11, 49 // CHECK-DEVICE: store i32 12, 50 // CHECK-BOTH: ret void 51 52 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev( 53 // CHECK-BOTH: store i32 31, 54 // CHECK-BOTH: ret void 55 56 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev( 57 // CHECK-BOTH: store i32 32, 58 // CHECK-BOTH: ret void 59 60 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev( 61 // CHECK-HOST: store i32 21, 62 // CHECK-DEVICE: store i32 22, 63 // CHECK-BOTH: ret void 64
