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 \ 8 // RUN: -fcuda-target-overloads -emit-llvm -o - %s \ 9 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s 10 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \ 11 // RUN: -fcuda-target-overloads -emit-llvm -o - %s \ 12 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \ 13 // RUN: -check-prefix=CHECK-DEVICE-STRICT %s 14 15 // Check target overloads handling with disabled call target checks. 16 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \ 17 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \ 18 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \ 19 // RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s 20 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \ 21 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \ 22 // RUN: -fcuda-is-device -o - %s \ 23 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \ 24 // RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s 25 26 #include "Inputs/cuda.h" 27 28 // Check constructors/destructors for D/H functions 29 int x; 30 struct s_cd_dh { 31 __host__ s_cd_dh() { x = 11; } 32 __device__ s_cd_dh() { x = 12; } 33 __host__ ~s_cd_dh() { x = 21; } 34 __device__ ~s_cd_dh() { x = 22; } 35 }; 36 37 struct s_cd_hd { 38 __host__ __device__ s_cd_hd() { x = 31; } 39 __host__ __device__ ~s_cd_hd() { x = 32; } 40 }; 41 42 // CHECK-BOTH: define void @_Z7wrapperv 43 #if defined(__CUDA_ARCH__) 44 __device__ 45 #else 46 __host__ 47 #endif 48 void wrapper() { 49 s_cd_dh scddh; 50 // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev( 51 s_cd_hd scdhd; 52 // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev 53 54 // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev( 55 // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev( 56 } 57 // CHECK-BOTH: ret void 58 59 // Now it's time to check what's been generated for the methods we used. 60 61 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev( 62 // CHECK-HOST: store i32 11, 63 // CHECK-DEVICE: store i32 12, 64 // CHECK-BOTH: ret void 65 66 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev( 67 // CHECK-BOTH: store i32 31, 68 // CHECK-BOTH: ret void 69 70 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev( 71 // CHECK-BOTH: store i32 32, 72 // CHECK-BOTH: ret void 73 74 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev( 75 // CHECK-HOST: store i32 21, 76 // CHECK-DEVICE: store i32 22, 77 // CHECK-BOTH: ret void 78