59 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: nvptx-registered-target3 4// Make sure we handle target overloads correctly. Most of this is checked in5// 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 %s9// 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 %s11 12#include "Inputs/cuda.h"13 14// Check constructors/destructors for D/H functions15#ifdef __CUDA_ARCH__16__device__17#endif18int x;19struct s_cd_dh {20 __host__ s_cd_dh() { x = 11; }21 __device__ s_cd_dh() { x = 12; }22};23 24struct s_cd_hd {25 __host__ __device__ s_cd_hd() { x = 31; }26 __host__ __device__ ~s_cd_hd() { x = 32; }27};28 29// CHECK-BOTH: define{{.*}} void @_Z7wrapperv30#if defined(__CUDA_ARCH__)31__device__32#else33__host__34#endif35void wrapper() {36 s_cd_dh scddh;37 // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(38 s_cd_hd scdhd;39 // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev40 41 // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(42}43// CHECK-BOTH: ret void44 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 void51 52// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(53// CHECK-BOTH: store i32 31,54// CHECK-BOTH: ret void55 56// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(57// CHECK-BOTH: store i32 32,58// CHECK-BOTH: ret void59