forked from OSchip/llvm-project
[CUDA] Merge most of CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu.
Summary: Previously we were using the codegen test to ensure that we choose the right overload. But we can do this within sema, with a bit of cleverness. I left the constructor/destructor checks in CodeGen, because these overloads (particularly on the destructors) are hard to check in Sema. Reviewers: tra Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D18386 llvm-svn: 264207
This commit is contained in:
Justin Lebar
parent
e82caa3055
commit
e5eed04d52
|
|
@ -1,7 +1,9 @@
|
|||
// REQUIRES: x86-registered-target
|
||||
// REQUIRES: nvptx-registered-target
|
||||
|
||||
// Make sure we handle target overloads correctly.
|
||||
// Make sure we handle target overloads correctly. Most of this is checked in
|
||||
// sema, but special functions like constructors and destructors are here.
|
||||
//
|
||||
// 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
|
||||
|
|
@ -23,235 +25,8 @@
|
|||
|
||||
#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
|
||||
|
||||
// HD functions are sometimes allowed to call H or D functions -- this
|
||||
// is an artifact of the source-to-source splitting performed by nvcc
|
||||
// that we need to mimic. During device mode compilation in nvcc, host
|
||||
// functions aren't present at all, so don't participate in
|
||||
// overloading. But in clang, H and D functions are present in both
|
||||
// compilation modes. Clang normally uses the target attribute as a
|
||||
// tiebreaker between overloads with otherwise identical priority, but
|
||||
// in order to match nvcc's behavior, we sometimes need to wholly
|
||||
// discard overloads that would not be present during compilation
|
||||
// under nvcc.
|
||||
|
||||
template <typename T> T template_vs_function(T arg) { return 15; }
|
||||
__device__ float template_vs_function(float arg) { return 16; }
|
||||
|
||||
// Here we expect to call the templated function during host
|
||||
// compilation, even if -fcuda-disable-target-call-checks is passed,
|
||||
// and even though C++ overload rules prefer the non-templated
|
||||
// function.
|
||||
// CHECK-BOTH-LABEL: define void @_Z5hd_tfv()
|
||||
__host__ __device__ void hd_tf(void) {
|
||||
template_vs_function(1.0f);
|
||||
// CHECK-HOST: call float @_Z20template_vs_functionIfET_S0_(float
|
||||
// CHECK-DEVICE: call float @_Z20template_vs_functionf(float
|
||||
template_vs_function(2.0);
|
||||
// CHECK-HOST: call double @_Z20template_vs_functionIdET_S0_(double
|
||||
// CHECK-DEVICE: call float @_Z20template_vs_functionf(float
|
||||
}
|
||||
|
||||
// Calls from __host__ and __device__ functions should always call the
|
||||
// overloaded function that matches their mode.
|
||||
// CHECK-HOST-LABEL: define void @_Z4h_tfv()
|
||||
__host__ void h_tf() {
|
||||
template_vs_function(1.0f);
|
||||
// CHECK-HOST: call float @_Z20template_vs_functionIfET_S0_(float
|
||||
template_vs_function(2.0);
|
||||
// CHECK-HOST: call double @_Z20template_vs_functionIdET_S0_(double
|
||||
}
|
||||
|
||||
// CHECK-DEVICE-LABEL: define void @_Z4d_tfv()
|
||||
__device__ void d_tf() {
|
||||
template_vs_function(1.0f);
|
||||
// CHECK-DEVICE: call float @_Z20template_vs_functionf(float
|
||||
template_vs_function(2.0);
|
||||
// CHECK-DEVICE: call float @_Z20template_vs_functionf(float
|
||||
}
|
||||
|
||||
// In case we have a mix of HD and H-only or D-only candidates in the
|
||||
// overload set, normal C++ overload resolution rules apply first.
|
||||
template <typename T> T template_vs_hd_function(T arg) { return 15; }
|
||||
__host__ __device__ float template_vs_hd_function(float arg) { return 16; }
|
||||
|
||||
// CHECK-BOTH-LABEL: define void @_Z7hd_thdfv()
|
||||
__host__ __device__ void hd_thdf() {
|
||||
template_vs_hd_function(1.0f);
|
||||
// CHECK-HOST: call float @_Z23template_vs_hd_functionf(float
|
||||
// CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
|
||||
template_vs_hd_function(1);
|
||||
// CHECK-HOST: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
|
||||
// CHECK-DEVICE-STRICT: call float @_Z23template_vs_hd_functionf(float
|
||||
// CHECK-DEVICE-NC: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
|
||||
}
|
||||
|
||||
// CHECK-HOST-LABEL: define void @_Z6h_thdfv()
|
||||
__host__ void h_thdf() {
|
||||
template_vs_hd_function(1.0f);
|
||||
// CHECK-HOST: call float @_Z23template_vs_hd_functionf(float
|
||||
template_vs_hd_function(1);
|
||||
// CHECK-HOST: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
|
||||
}
|
||||
|
||||
// CHECK-DEVICE-LABEL: define void @_Z6d_thdfv()
|
||||
__device__ void d_thdf() {
|
||||
template_vs_hd_function(1.0f);
|
||||
// CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
|
||||
template_vs_hd_function(1);
|
||||
// Host-only function template is not callable with strict call checks,
|
||||
// so for device side HD function will be the only choice.
|
||||
// CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
|
||||
}
|
||||
|
||||
// Check that overloads still work the same way on both host and
|
||||
// device side when the overload set contains only functions from one
|
||||
// side of compilation.
|
||||
__device__ float device_only_function(int arg) { return 17; }
|
||||
__device__ float device_only_function(float arg) { return 18; }
|
||||
|
||||
__host__ float host_only_function(int arg) { return 19; }
|
||||
__host__ float host_only_function(float arg) { return 20; }
|
||||
|
||||
// CHECK-BOTH-LABEL: define void @_Z6hd_dofv()
|
||||
__host__ __device__ void hd_dof() {
|
||||
#ifdef NOCHECKS
|
||||
device_only_function(1.0f);
|
||||
// CHECK-BOTH-NC: call float @_Z20device_only_functionf(float
|
||||
device_only_function(1);
|
||||
// CHECK-BOTH-NC: call float @_Z20device_only_functioni(i32
|
||||
host_only_function(1.0f);
|
||||
// CHECK-BOTH-NC: call float @_Z18host_only_functionf(float
|
||||
host_only_function(1);
|
||||
// CHECK-BOTH-NC: call float @_Z18host_only_functioni(i32
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
// CHECK-HOST-LABEL: define void @_Z5hostfv()
|
||||
__host__ void hostf(void) {
|
||||
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,
|
||||
|
||||
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,
|
||||
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()
|
||||
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
|
||||
int x;
|
||||
struct s_cd_dh {
|
||||
__host__ s_cd_dh() { x = 11; }
|
||||
__device__ s_cd_dh() { x = 12; }
|
||||
|
|
@ -300,4 +75,3 @@ void wrapper() {
|
|||
// CHECK-HOST: store i32 21,
|
||||
// CHECK-DEVICE: store i32 22,
|
||||
// CHECK-BOTH: ret void
|
||||
|
||||
|
|
|
|||
|
|
@ -16,58 +16,80 @@
|
|||
|
||||
#include "Inputs/cuda.h"
|
||||
|
||||
typedef int (*fp_t)();
|
||||
typedef void (*gp_t)();
|
||||
// Opaque return types used to check that we pick the right overloads.
|
||||
struct HostReturnTy {};
|
||||
struct HostReturnTy2 {};
|
||||
struct DeviceReturnTy {};
|
||||
struct DeviceReturnTy2 {};
|
||||
struct HostDeviceReturnTy {};
|
||||
struct TemplateReturnTy {};
|
||||
|
||||
typedef HostReturnTy (*HostFnPtr)();
|
||||
typedef DeviceReturnTy (*DeviceFnPtr)();
|
||||
typedef HostDeviceReturnTy (*HostDeviceFnPtr)();
|
||||
typedef void (*GlobalFnPtr)(); // __global__ functions must return void.
|
||||
|
||||
// CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}
|
||||
// compilation.
|
||||
#ifdef __CUDA_ARCH__
|
||||
typedef DeviceReturnTy CurrentReturnTy;
|
||||
#else
|
||||
typedef HostReturnTy CurrentReturnTy;
|
||||
#endif
|
||||
|
||||
// CurrentFnPtr is a function pointer to a {host,device} function during
|
||||
// {host,device} compilation.
|
||||
typedef CurrentReturnTy (*CurrentFnPtr)();
|
||||
|
||||
// Host and unattributed functions can't be overloaded.
|
||||
__host__ void hh() {} // expected-note {{previous definition is here}}
|
||||
void hh() {} // expected-error {{redefinition of 'hh'}}
|
||||
|
||||
// H/D overloading is OK.
|
||||
__host__ int dh() { return 2; }
|
||||
__device__ int dh() { return 2; }
|
||||
__host__ HostReturnTy dh() { return HostReturnTy(); }
|
||||
__device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
|
||||
|
||||
// H/HD and D/HD are not allowed.
|
||||
__host__ __device__ int hdh() { return 5; } // expected-note {{previous definition is here}}
|
||||
__host__ int hdh() { return 4; } // expected-error {{redefinition of 'hdh'}}
|
||||
__host__ __device__ int hdh() { return 0; } // expected-note {{previous definition is here}}
|
||||
__host__ int hdh() { return 0; } // expected-error {{redefinition of 'hdh'}}
|
||||
|
||||
__host__ int hhd() { return 4; } // expected-note {{previous definition is here}}
|
||||
__host__ __device__ int hhd() { return 5; } // expected-error {{redefinition of 'hhd'}}
|
||||
__host__ int hhd() { return 0; } // expected-note {{previous definition is here}}
|
||||
__host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
|
||||
// expected-warning@-1 {{attribute declaration must precede definition}}
|
||||
// expected-note@-3 {{previous definition is here}}
|
||||
|
||||
__host__ __device__ int hdd() { return 7; } // expected-note {{previous definition is here}}
|
||||
__device__ int hdd() { return 6; } // expected-error {{redefinition of 'hdd'}}
|
||||
__host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
|
||||
__device__ int hdd() { return 0; } // expected-error {{redefinition of 'hdd'}}
|
||||
|
||||
__device__ int dhd() { return 6; } // expected-note {{previous definition is here}}
|
||||
__host__ __device__ int dhd() { return 7; } // expected-error {{redefinition of 'dhd'}}
|
||||
__device__ int dhd() { return 0; } // expected-note {{previous definition is here}}
|
||||
__host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
|
||||
// expected-warning@-1 {{attribute declaration must precede definition}}
|
||||
// expected-note@-3 {{previous definition is here}}
|
||||
|
||||
// Same tests for extern "C" functions.
|
||||
extern "C" __host__ int chh() {return 11;} // expected-note {{previous definition is here}}
|
||||
extern "C" int chh() {return 11;} // expected-error {{redefinition of 'chh'}}
|
||||
extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
|
||||
extern "C" int chh() { return 0; } // expected-error {{redefinition of 'chh'}}
|
||||
|
||||
// H/D overloading is OK.
|
||||
extern "C" __device__ int cdh() {return 10;}
|
||||
extern "C" __host__ int cdh() {return 11;}
|
||||
extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
|
||||
extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
|
||||
|
||||
// H/HD and D/HD overloading is not allowed.
|
||||
extern "C" __host__ __device__ int chhd1() {return 12;} // expected-note {{previous definition is here}}
|
||||
extern "C" __host__ int chhd1() {return 13;} // expected-error {{redefinition of 'chhd1'}}
|
||||
extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous definition is here}}
|
||||
extern "C" __host__ int chhd1() { return 0; } // expected-error {{redefinition of 'chhd1'}}
|
||||
|
||||
extern "C" __host__ int chhd2() {return 13;} // expected-note {{previous definition is here}}
|
||||
extern "C" __host__ __device__ int chhd2() {return 12;} // expected-error {{redefinition of 'chhd2'}}
|
||||
extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous definition is here}}
|
||||
extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
|
||||
// expected-warning@-1 {{attribute declaration must precede definition}}
|
||||
// expected-note@-3 {{previous definition is here}}
|
||||
|
||||
// Helper functions to verify calling restrictions.
|
||||
__device__ int d() { return 8; }
|
||||
__device__ DeviceReturnTy d() { return DeviceReturnTy(); }
|
||||
// expected-note@-1 1+ {{'d' declared here}}
|
||||
// expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
|
||||
// expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
|
||||
|
||||
__host__ int h() { return 9; }
|
||||
__host__ HostReturnTy h() { return HostReturnTy(); }
|
||||
// expected-note@-1 1+ {{'h' declared here}}
|
||||
// expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
|
||||
// expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
|
||||
|
|
@ -79,123 +101,112 @@ __global__ void g() {}
|
|||
// expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
|
||||
// expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
|
||||
|
||||
extern "C" __device__ int cd() {return 10;}
|
||||
extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
|
||||
// expected-note@-1 1+ {{'cd' declared here}}
|
||||
// expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
|
||||
// expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
|
||||
|
||||
extern "C" __host__ int ch() {return 11;}
|
||||
extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
|
||||
// expected-note@-1 1+ {{'ch' declared here}}
|
||||
// expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
|
||||
// expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
|
||||
// expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
|
||||
|
||||
__host__ void hostf() {
|
||||
fp_t dp = d;
|
||||
// expected-error@-1 {{reference to __device__ function 'd' in __host__ function}}
|
||||
fp_t cdp = cd;
|
||||
// expected-error@-1 {{reference to __device__ function 'cd' in __host__ function}}
|
||||
fp_t hp = h;
|
||||
fp_t chp = ch;
|
||||
fp_t dhp = dh;
|
||||
fp_t cdhp = cdh;
|
||||
gp_t gp = g;
|
||||
DeviceFnPtr fp_d = d; // expected-error {{reference to __device__ function 'd' in __host__ function}}
|
||||
DeviceReturnTy ret_d = d(); // expected-error {{no matching function for call to 'd'}}
|
||||
DeviceFnPtr fp_cd = cd; // expected-error {{reference to __device__ function 'cd' in __host__ function}}
|
||||
DeviceReturnTy ret_cd = cd(); // expected-error {{no matching function for call to 'cd'}}
|
||||
|
||||
d();
|
||||
// expected-error@-1 {{no matching function for call to 'd'}}
|
||||
cd();
|
||||
// expected-error@-1 {{no matching function for call to 'cd'}}
|
||||
h();
|
||||
ch();
|
||||
dh();
|
||||
cdh();
|
||||
HostFnPtr fp_h = h;
|
||||
HostReturnTy ret_h = h();
|
||||
HostFnPtr fp_ch = ch;
|
||||
HostReturnTy ret_ch = ch();
|
||||
|
||||
HostFnPtr fp_dh = dh;
|
||||
HostReturnTy ret_dh = dh();
|
||||
HostFnPtr fp_cdh = cdh;
|
||||
HostReturnTy ret_cdh = cdh();
|
||||
|
||||
GlobalFnPtr fp_g = g;
|
||||
g(); // expected-error {{call to global function g not configured}}
|
||||
g<<<0,0>>>();
|
||||
g<<<0, 0>>>();
|
||||
}
|
||||
|
||||
__device__ void devicef() {
|
||||
fp_t dp = d;
|
||||
fp_t cdp = cd;
|
||||
fp_t hp = h;
|
||||
// expected-error@-1 {{reference to __host__ function 'h' in __device__ function}}
|
||||
fp_t chp = ch;
|
||||
// expected-error@-1 {{reference to __host__ function 'ch' in __device__ function}}
|
||||
fp_t dhp = dh;
|
||||
fp_t cdhp = cdh;
|
||||
gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
|
||||
DeviceFnPtr fp_d = d;
|
||||
DeviceReturnTy ret_d = d();
|
||||
DeviceFnPtr fp_cd = cd;
|
||||
DeviceReturnTy ret_cd = cd();
|
||||
|
||||
d();
|
||||
cd();
|
||||
h(); // expected-error {{no matching function for call to 'h'}}
|
||||
ch(); // expected-error {{no matching function for call to 'ch'}}
|
||||
dh();
|
||||
cdh();
|
||||
HostFnPtr fp_h = h; // expected-error {{reference to __host__ function 'h' in __device__ function}}
|
||||
HostReturnTy ret_h = h(); // expected-error {{no matching function for call to 'h'}}
|
||||
HostFnPtr fp_ch = ch; // expected-error {{reference to __host__ function 'ch' in __device__ function}}
|
||||
HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
|
||||
|
||||
DeviceFnPtr fp_dh = dh;
|
||||
DeviceReturnTy ret_dh = dh();
|
||||
DeviceFnPtr fp_cdh = cdh;
|
||||
DeviceReturnTy ret_cdh = cdh();
|
||||
|
||||
GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
|
||||
g(); // expected-error {{no matching function for call to 'g'}}
|
||||
g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
|
||||
}
|
||||
|
||||
__global__ void globalf() {
|
||||
fp_t dp = d;
|
||||
fp_t cdp = cd;
|
||||
fp_t hp = h;
|
||||
// expected-error@-1 {{reference to __host__ function 'h' in __global__ function}}
|
||||
fp_t chp = ch;
|
||||
// expected-error@-1 {{reference to __host__ function 'ch' in __global__ function}}
|
||||
fp_t dhp = dh;
|
||||
fp_t cdhp = cdh;
|
||||
gp_t gp = g;
|
||||
// expected-error@-1 {{reference to __global__ function 'g' in __global__ function}}
|
||||
DeviceFnPtr fp_d = d;
|
||||
DeviceReturnTy ret_d = d();
|
||||
DeviceFnPtr fp_cd = cd;
|
||||
DeviceReturnTy ret_cd = cd();
|
||||
|
||||
d();
|
||||
cd();
|
||||
h();
|
||||
// expected-error@-1 {{no matching function for call to 'h'}}
|
||||
ch();
|
||||
// expected-error@-1 {{no matching function for call to 'ch'}}
|
||||
dh();
|
||||
cdh();
|
||||
HostFnPtr fp_h = h; // expected-error {{reference to __host__ function 'h' in __global__ function}}
|
||||
HostReturnTy ret_h = h(); // expected-error {{no matching function for call to 'h'}}
|
||||
HostFnPtr fp_ch = ch; // expected-error {{reference to __host__ function 'ch' in __global__ function}}
|
||||
HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
|
||||
|
||||
DeviceFnPtr fp_dh = dh;
|
||||
DeviceReturnTy ret_dh = dh();
|
||||
DeviceFnPtr fp_cdh = cdh;
|
||||
DeviceReturnTy ret_cdh = cdh();
|
||||
|
||||
GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
|
||||
g(); // expected-error {{no matching function for call to 'g'}}
|
||||
g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
|
||||
}
|
||||
|
||||
__host__ __device__ void hostdevicef() {
|
||||
fp_t dp = d;
|
||||
fp_t cdp = cd;
|
||||
DeviceFnPtr fp_d = d;
|
||||
DeviceReturnTy ret_d = d();
|
||||
DeviceFnPtr fp_cd = cd;
|
||||
DeviceReturnTy ret_cd = cd();
|
||||
#if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
|
||||
// expected-error@-3 {{reference to __device__ function 'd' in __host__ __device__ function}}
|
||||
// expected-error@-3 {{reference to __device__ function 'cd' in __host__ __device__ function}}
|
||||
// expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
|
||||
// expected-error@-5 {{no matching function for call to 'd'}}
|
||||
// expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
|
||||
// expected-error@-5 {{no matching function for call to 'cd'}}
|
||||
#endif
|
||||
|
||||
fp_t hp = h;
|
||||
fp_t chp = ch;
|
||||
HostFnPtr fp_h = h;
|
||||
HostReturnTy ret_h = h();
|
||||
HostFnPtr fp_ch = ch;
|
||||
HostReturnTy ret_ch = ch();
|
||||
#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
|
||||
// expected-error@-3 {{reference to __host__ function 'h' in __host__ __device__ function}}
|
||||
// expected-error@-3 {{reference to __host__ function 'ch' in __host__ __device__ function}}
|
||||
// expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
|
||||
// expected-error@-5 {{no matching function for call to 'h'}}
|
||||
// expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
|
||||
// expected-error@-5 {{no matching function for call to 'ch'}}
|
||||
#endif
|
||||
|
||||
fp_t dhp = dh;
|
||||
fp_t cdhp = cdh;
|
||||
gp_t gp = g;
|
||||
CurrentFnPtr fp_dh = dh;
|
||||
CurrentReturnTy ret_dh = dh();
|
||||
CurrentFnPtr fp_cdh = cdh;
|
||||
CurrentReturnTy ret_cdh = cdh();
|
||||
|
||||
GlobalFnPtr fp_g = g;
|
||||
#if defined(__CUDA_ARCH__)
|
||||
// expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
|
||||
#endif
|
||||
|
||||
d();
|
||||
cd();
|
||||
#if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
|
||||
// expected-error@-3 {{no matching function for call to 'd'}}
|
||||
// expected-error@-3 {{no matching function for call to 'cd'}}
|
||||
#endif
|
||||
|
||||
h();
|
||||
ch();
|
||||
#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
|
||||
// expected-error@-3 {{no matching function for call to 'h'}}
|
||||
// expected-error@-3 {{no matching function for call to 'ch'}}
|
||||
#endif
|
||||
|
||||
dh();
|
||||
cdh();
|
||||
g();
|
||||
g<<<0,0>>>();
|
||||
#if !defined(__CUDA_ARCH__)
|
||||
|
|
@ -207,11 +218,11 @@ __host__ __device__ void hostdevicef() {
|
|||
}
|
||||
|
||||
// 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;
|
||||
HostFnPtr fp_h = h;
|
||||
HostFnPtr fp_ch = ch;
|
||||
CurrentFnPtr fp_dh = dh;
|
||||
CurrentFnPtr fp_cdh = cdh;
|
||||
GlobalFnPtr fp_g = g;
|
||||
|
||||
|
||||
// Test overloading of destructors
|
||||
|
|
@ -305,3 +316,96 @@ private:
|
|||
};
|
||||
__global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
|
||||
void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}
|
||||
|
||||
// HD functions are sometimes allowed to call H or D functions -- this
|
||||
// is an artifact of the source-to-source splitting performed by nvcc
|
||||
// that we need to mimic. During device mode compilation in nvcc, host
|
||||
// functions aren't present at all, so don't participate in
|
||||
// overloading. But in clang, H and D functions are present in both
|
||||
// compilation modes. Clang normally uses the target attribute as a
|
||||
// tiebreaker between overloads with otherwise identical priority, but
|
||||
// in order to match nvcc's behavior, we sometimes need to wholly
|
||||
// discard overloads that would not be present during compilation
|
||||
// under nvcc.
|
||||
|
||||
template <typename T> TemplateReturnTy template_vs_function(T arg) {
|
||||
return TemplateReturnTy();
|
||||
}
|
||||
__device__ DeviceReturnTy template_vs_function(float arg) {
|
||||
return DeviceReturnTy();
|
||||
}
|
||||
|
||||
// Here we expect to call the templated function during host compilation, even
|
||||
// if -fcuda-disable-target-call-checks is passed, and even though C++ overload
|
||||
// rules prefer the non-templated function.
|
||||
__host__ __device__ void test_host_device_calls_template(void) {
|
||||
#ifdef __CUDA_ARCH__
|
||||
typedef DeviceReturnTy ExpectedReturnTy;
|
||||
#else
|
||||
typedef TemplateReturnTy ExpectedReturnTy;
|
||||
#endif
|
||||
|
||||
ExpectedReturnTy ret1 = template_vs_function(1.0f);
|
||||
ExpectedReturnTy ret2 = template_vs_function(2.0);
|
||||
}
|
||||
|
||||
// Calls from __host__ and __device__ functions should always call the
|
||||
// overloaded function that matches their mode.
|
||||
__host__ void test_host_calls_template_fn() {
|
||||
TemplateReturnTy ret1 = template_vs_function(1.0f);
|
||||
TemplateReturnTy ret2 = template_vs_function(2.0);
|
||||
}
|
||||
|
||||
__device__ void test_device_calls_template_fn() {
|
||||
DeviceReturnTy ret1 = template_vs_function(1.0f);
|
||||
DeviceReturnTy ret2 = template_vs_function(2.0);
|
||||
}
|
||||
|
||||
// If we have a mix of HD and H-only or D-only candidates in the overload set,
|
||||
// normal C++ overload resolution rules apply first.
|
||||
template <typename T> TemplateReturnTy template_vs_hd_function(T arg) {
|
||||
return TemplateReturnTy();
|
||||
}
|
||||
__host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
|
||||
return HostDeviceReturnTy();
|
||||
}
|
||||
|
||||
__host__ __device__ void test_host_device_calls_hd_template() {
|
||||
HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
|
||||
|
||||
#if defined(__CUDA_ARCH__) && !defined(NOCHECKS)
|
||||
typedef HostDeviceReturnTy ExpectedReturnTy;
|
||||
#else
|
||||
typedef TemplateReturnTy ExpectedReturnTy;
|
||||
#endif
|
||||
ExpectedReturnTy ret2 = template_vs_hd_function(1);
|
||||
}
|
||||
|
||||
__host__ void test_host_calls_hd_template() {
|
||||
HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
|
||||
TemplateReturnTy ret2 = template_vs_hd_function(1);
|
||||
}
|
||||
|
||||
__device__ void test_device_calls_hd_template() {
|
||||
HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
|
||||
// Host-only function template is not callable with strict call checks,
|
||||
// so for device side HD function will be the only choice.
|
||||
HostDeviceReturnTy ret2 = template_vs_hd_function(1);
|
||||
}
|
||||
|
||||
// Check that overloads still work the same way on both host and
|
||||
// device side when the overload set contains only functions from one
|
||||
// side of compilation.
|
||||
__device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
|
||||
__device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
|
||||
__host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
|
||||
__host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
|
||||
|
||||
__host__ __device__ void test_host_device_nochecks_overloading() {
|
||||
#ifdef NOCHECKS
|
||||
DeviceReturnTy ret1 = device_only_function(1);
|
||||
DeviceReturnTy2 ret2 = device_only_function(1.0f);
|
||||
HostReturnTy ret3 = host_only_function(1);
|
||||
HostReturnTy2 ret4 = host_only_function(1.0f);
|
||||
#endif
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue