caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
a6714922d100fc9ca90cfc4dc4bb0df875206197
clang-p2996/clang/test/SemaCUDA/function-overload.cu
Justin Lebar d3fd70dedd [CUDA] Rework tests now that we emit deferred diagnostics during sema. Test-only change. 
Summary:
Previously we had to split out a lot of our tests into a test that
checked only immediate errors and a test that checked only deferred
errors.  This was because, if you emitted any immediate errors, we
wouldn't run codegen, where the deferred errors were emitted.

We've fixed this, and now emit deferred errors during sema.  This lets
us merge a bunch of tests, and lets us convert some other tests to
-fsyntax-only.

Reviewers: tra

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D25755

llvm-svn: 284553
2016-10-19 00:06:49 +00:00

399 lines
17 KiB
Plaintext
Raw Blame History

// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
#include "Inputs/cuda.h"
// 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__ HostReturnTy dh() { return HostReturnTy(); }
__device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
// H/HD and D/HD are not allowed.
__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 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 0; } // expected-note {{previous definition is here}}
__device__ int hdd() { return 0; } // expected-error {{redefinition of 'hdd'}}
__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 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__ 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 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 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__ 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__ 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}}
// expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
__global__ void g() {}
// expected-note@-1 1+ {{'g' declared here}}
// expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
// 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__ 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__ 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() {
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>>>();
}
__device__ void devicef() {
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() {
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() {
DeviceFnPtr fp_d = d;
DeviceReturnTy ret_d = d();
DeviceFnPtr fp_cd = cd;
DeviceReturnTy ret_cd = cd();
#if !defined(__CUDA_ARCH__)
// expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
// expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
// expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
// expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
#endif
HostFnPtr fp_h = h;
HostReturnTy ret_h = h();
HostFnPtr fp_ch = ch;
HostReturnTy ret_ch = ch();
#if defined(__CUDA_ARCH__)
// expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
// expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
// expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
// expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
#endif
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
g();
#if defined (__CUDA_ARCH__)
// expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
#else
// expected-error@-4 {{call to global function g not configured}}
#endif
g<<<0,0>>>();
#if defined(__CUDA_ARCH__)
// expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
#endif
}
// Test for address of overloaded function resolution in the global context.
HostFnPtr fp_h = h;
HostFnPtr fp_ch = ch;
CurrentFnPtr fp_dh = dh;
CurrentFnPtr fp_cdh = cdh;
GlobalFnPtr fp_g = g;
// Test overloading of destructors
// Can't mix H and unattributed destructors
struct d_h {
~d_h() {} // expected-note {{previous declaration is here}}
__host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
};
// HD is OK
struct d_hd {
__host__ __device__ ~d_hd() {}
};
// Test overloading of member functions
struct m_h {
void operator delete(void *ptr); // expected-note {{previous declaration is here}}
__host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
};
// D/H overloading is OK
struct m_dh {
__device__ void operator delete(void *ptr);
__host__ void operator delete(void *ptr);
};
// HD by itself is OK
struct m_hd {
__device__ __host__ void operator delete(void *ptr);
};
struct m_hhd {
__host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
__host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
};
struct m_hdh {
__host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
__host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
};
struct m_dhd {
__device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
__host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
};
struct m_hdd {
__host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
__device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
};
// __global__ functions can't be overloaded based on attribute
// difference.
struct G {
friend void friend_of_g(G &arg);
private:
int x;
};
__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)
#ifdef __CUDA_ARCH__
//expected-note@-2 {{declared here}}
#endif
{
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);
TemplateReturnTy ret2 = template_vs_hd_function(1);
#ifdef __CUDA_ARCH__
// expected-error@-2 {{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
#endif
}
__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(); }
#ifndef __CUDA_ARCH__
// expected-note@-3 {{'device_only_function' declared here}}
// expected-note@-3 {{'device_only_function' declared here}}
#endif
__host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
__host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
#ifdef __CUDA_ARCH__
// expected-note@-3 {{'host_only_function' declared here}}
// expected-note@-3 {{'host_only_function' declared here}}
#endif
__host__ __device__ void test_host_device_single_side_overloading() {
DeviceReturnTy ret1 = device_only_function(1);
DeviceReturnTy2 ret2 = device_only_function(1.0f);
#ifndef __CUDA_ARCH__
// expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
// expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
#endif
HostReturnTy ret3 = host_only_function(1);
HostReturnTy2 ret4 = host_only_function(1.0f);
#ifdef __CUDA_ARCH__
// expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
// expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
#endif
}
// Verify that we allow overloading function templates.
template <typename T> __host__ T template_overload(const T &a) { return a; };
template <typename T> __device__ T template_overload(const T &a) { return a; };
__host__ void test_host_template_overload() {
template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
}
__device__ void test_device_template_overload() {
template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
}
Reference in New Issue View Git Blame Copy Permalink