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