[PATCH] D17581: [CUDA] disable attribute-based overloading for __global__ functions.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 24 13:31:21 PST 2016
tra created this revision.
tra added a reviewer: jlebar.
tra added a subscriber: cfe-commits.
__global__ functions are present on both host and device side,
so providing __host__ or __device__ overloads is not going to
do anything useful.
http://reviews.llvm.org/D17581
Files:
lib/Sema/SemaOverload.cpp
test/SemaCUDA/function-overload.cu
Index: test/SemaCUDA/function-overload.cu
===================================================================
--- test/SemaCUDA/function-overload.cu
+++ test/SemaCUDA/function-overload.cu
@@ -302,3 +302,13 @@
__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'}}
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1129,7 +1129,10 @@
// 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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17581.48973.patch
Type: text/x-patch
Size: 1577 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160224/4dfcbea1/attachment.bin>
More information about the cfe-commits
mailing list