clang/test/CodeGenCUDA/function-overload.cu

215 строки
7.5 KiB
Plaintext

// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
// Make sure we handle target overloads correctly.
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
// RUN: -fcuda-target-overloads -emit-llvm -o - %s \
// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
// RUN: -fcuda-target-overloads -emit-llvm -o - %s \
// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
// Check target overloads handling with disabled call target checks.
// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
// RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \
// RUN: -fcuda-is-device -o - %s \
// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
// RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
#include "Inputs/cuda.h"
typedef int (*fp_t)(void);
typedef void (*gp_t)(void);
// CHECK-HOST: @hp = global i32 ()* @_Z1hv
// CHECK-HOST: @chp = global i32 ()* @ch
// CHECK-HOST: @dhp = global i32 ()* @_Z2dhv
// CHECK-HOST: @cdhp = global i32 ()* @cdh
// CHECK-HOST: @gp = global void ()* @_Z1gv
// CHECK-BOTH-LABEL: define i32 @_Z2dhv()
__device__ int dh(void) { return 1; }
// CHECK-DEVICE: ret i32 1
__host__ int dh(void) { return 2; }
// CHECK-HOST: ret i32 2
// CHECK-BOTH-LABEL: define i32 @_Z2hdv()
__host__ __device__ int hd(void) { return 3; }
// CHECK-BOTH: ret i32 3
// CHECK-DEVICE-LABEL: define i32 @_Z1dv()
__device__ int d(void) { return 8; }
// CHECK-DEVICE: ret i32 8
// CHECK-HOST-LABEL: define i32 @_Z1hv()
__host__ int h(void) { return 9; }
// CHECK-HOST: ret i32 9
// CHECK-BOTH-LABEL: define void @_Z1gv()
__global__ void g(void) {}
// CHECK-BOTH: ret void
// mangled names of extern "C" __host__ __device__ functions clash
// with those of their __host__/__device__ counterparts, so
// overloading of extern "C" functions can only happen for __host__
// and __device__ functions -- we never codegen them in the same
// compilation and therefore mangled name conflict is not a problem.
// CHECK-BOTH-LABEL: define i32 @cdh()
extern "C" __device__ int cdh(void) {return 10;}
// CHECK-DEVICE: ret i32 10
extern "C" __host__ int cdh(void) {return 11;}
// CHECK-HOST: ret i32 11
// CHECK-DEVICE-LABEL: define i32 @cd()
extern "C" __device__ int cd(void) {return 12;}
// CHECK-DEVICE: ret i32 12
// CHECK-HOST-LABEL: define i32 @ch()
extern "C" __host__ int ch(void) {return 13;}
// CHECK-HOST: ret i32 13
// CHECK-BOTH-LABEL: define i32 @chd()
extern "C" __host__ __device__ int chd(void) {return 14;}
// CHECK-BOTH: ret i32 14
// CHECK-HOST-LABEL: define void @_Z5hostfv()
__host__ void hostf(void) {
#if defined (NOCHECKS)
fp_t dp = d; // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp,
#endif
fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
#if defined (NOCHECKS)
d(); // CHECK-HOST-NC: call i32 @_Z1dv()
cd(); // CHECK-HOST-NC: call i32 @cd()
#endif
h(); // CHECK-HOST: call i32 @_Z1hv()
ch(); // CHECK-HOST: call i32 @ch()
dh(); // CHECK-HOST: call i32 @_Z2dhv()
cdh(); // CHECK-HOST: call i32 @cdh()
g<<<0,0>>>(); // CHECK-HOST: call void @_Z1gv()
}
// CHECK-DEVICE-LABEL: define void @_Z7devicefv()
__device__ void devicef(void) {
fp_t dp = d; // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
#if defined (NOCHECKS)
fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp,
#endif
fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
d(); // CHECK-DEVICE: call i32 @_Z1dv()
cd(); // CHECK-DEVICE: call i32 @cd()
#if defined (NOCHECKS)
h(); // CHECK-DEVICE-NC: call i32 @_Z1hv()
ch(); // CHECK-DEVICE-NC: call i32 @ch()
#endif
dh(); // CHECK-DEVICE: call i32 @_Z2dhv()
cdh(); // CHECK-DEVICE: call i32 @cdh()
}
// CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
__host__ __device__ void hostdevicef(void) {
#if defined (NOCHECKS)
fp_t dp = d; // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp,
fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp,
#endif
fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp,
fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp,
fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp,
fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp,
#endif
#if defined (NOCHECKS)
d(); // CHECK-BOTH-NC: call i32 @_Z1dv()
cd(); // CHECK-BOTH-NC: call i32 @cd()
h(); // CHECK-BOTH-NC: call i32 @_Z1hv()
ch(); // CHECK-BOTH-NC: call i32 @ch()
#endif
dh(); // CHECK-BOTH: call i32 @_Z2dhv()
cdh(); // CHECK-BOTH: call i32 @cdh()
#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
g<<<0,0>>>(); // CHECK-HOST-NC: call void @_Z1gv()
#endif
}
// Test for address of overloaded function resolution in the global context.
fp_t hp = h;
fp_t chp = ch;
fp_t dhp = dh;
fp_t cdhp = cdh;
gp_t gp = g;
int x;
// Check constructors/destructors for D/H functions
struct s_cd_dh {
__host__ s_cd_dh() { x = 11; }
__device__ s_cd_dh() { x = 12; }
__host__ ~s_cd_dh() { x = 21; }
__device__ ~s_cd_dh() { x = 22; }
};
struct s_cd_hd {
__host__ __device__ s_cd_hd() { x = 31; }
__host__ __device__ ~s_cd_hd() { x = 32; }
};
// CHECK-BOTH: define void @_Z7wrapperv
#if defined(__CUDA_ARCH__)
__device__
#else
__host__
#endif
void wrapper() {
s_cd_dh scddh;
// CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
s_cd_hd scdhd;
// CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
// CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
// CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
}
// CHECK-BOTH: ret void
// Now it's time to check what's been generated for the methods we used.
// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
// CHECK-HOST: store i32 11,
// CHECK-DEVICE: store i32 12,
// CHECK-BOTH: ret void
// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
// CHECK-BOTH: store i32 31,
// CHECK-BOTH: ret void
// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
// CHECK-BOTH: store i32 32,
// CHECK-BOTH: ret void
// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
// CHECK-HOST: store i32 21,
// CHECK-DEVICE: store i32 22,
// CHECK-BOTH: ret void