Clang Project

clang_source_code/test/CodeGenCUDA/function-overload.cu
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
15int x;
16struct s_cd_dh {
17  __host__ s_cd_dh() { x = 11; }
18  __device__ s_cd_dh() { x = 12; }
19};
20
21struct s_cd_hd {
22  __host__ __device__ s_cd_hd() { x = 31; }
23  __host__ __device__ ~s_cd_hd() { x = 32; }
24};
25
26// CHECK-BOTH: define void @_Z7wrapperv
27#if defined(__CUDA_ARCH__)
28__device__
29#else
30__host__
31#endif
32void wrapper() {
33  s_cd_dh scddh;
34  // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
35  s_cd_hd scdhd;
36  // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
37
38  // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
39}
40// CHECK-BOTH: ret void
41
42// Now it's time to check what's been generated for the methods we used.
43
44// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
45// CHECK-HOST:   store i32 11,
46// CHECK-DEVICE: store i32 12,
47// CHECK-BOTH: ret void
48
49// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
50// CHECK-BOTH:   store i32 31,
51// CHECK-BOTH: ret void
52
53// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
54// CHECK-BOTH: store i32 32,
55// CHECK-BOTH: ret void
56