replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[CUDA] do not allow attribute-based overloading for __global__ functions. 

__global__ functions are present on both host and device side,
so providing __host__ or __device__ overloads is not going to
do anything useful.

llvm-svn: 261778
This commit is contained in:
Artem Belevich 2016-02-24 21:54:45 +00:00
parent d7a35492ad
commit 1ef9b59284
2 changed files with 14 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
clang/lib/Sema/SemaOverload.cpp
View File

@ -1129,7 +1129,10 @@ bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old,
// Don't allow mixing of HD with other kinds. This guarantees that
// we have only one viable function with this signature on any
// side of CUDA compilation .
if ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice))
// __global__ functions can't be overloaded based on attribute
// difference because, like HD, they also exist on both sides.
if ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice) ||
(NewTarget == CFT_Global) || (OldTarget == CFT_Global))
return false;
// Allow overloading of functions with same signature, but

10
clang/test/SemaCUDA/function-overload.cu
View File

@ -302,3 +302,13 @@ 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'}}