From: Justin Lebar Date: Wed, 23 Mar 2016 22:42:30 +0000 (+0000) Subject: [CUDA] Merge most of CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=10fef645a6f8c9c73683fd8f6b95fd4197253e05;p=clang [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@264207 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/test/CodeGenCUDA/function-overload.cu b/test/CodeGenCUDA/function-overload.cu index 3f8f930106..2fa6bb7afe 100644 --- a/test/CodeGenCUDA/function-overload.cu +++ b/test/CodeGenCUDA/function-overload.cu @@ -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 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 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 - diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu index 93a4c9bcfa..8de9e63bd6 100644 --- a/test/SemaCUDA/function-overload.cu +++ b/test/SemaCUDA/function-overload.cu @@ -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; - - 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(); + 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'}} + + 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}} - - d(); - cd(); - h(); // expected-error {{no matching function for call to 'h'}} - ch(); // expected-error {{no matching function for call to 'ch'}} - dh(); - cdh(); + DeviceFnPtr fp_d = d; + DeviceReturnTy ret_d = d(); + DeviceFnPtr fp_cd = cd; + DeviceReturnTy ret_cd = cd(); + + 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}} - - 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(); + DeviceFnPtr fp_d = d; + DeviceReturnTy ret_d = d(); + DeviceFnPtr fp_cd = cd; + DeviceReturnTy ret_cd = cd(); + + 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 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 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 +}