Author: tra
Date: Wed Feb 24 15:54:45 2016
New Revision: 261778
URL: http://llvm.org/viewvc/llvm-project?rev=261778&view=rev
Log:
[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.
Modified:
cfe/trunk/lib/Sema/SemaOverload.cpp
cfe/trunk/test/SemaCUDA/function-overload.cu
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=261778&r1=261777&r2=261778&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Wed Feb 24 15:54:45 2016
@@ -1129,7 +1129,10 @@ bool Sema::IsOverload(FunctionDecl *New,
// 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
Modified: cfe/trunk/test/SemaCUDA/function-overload.cu
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=261778&r1=261777&r2=261778&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload.cu Wed Feb 24 15:54:45 2016
@@ -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'}}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
