From: Artem Belevich <tra@google.com>
Date: Wed, 24 Feb 2016 21:54:45 +0000 (+0000)
Subject: [CUDA] do not allow attribute-based overloading for __global__ functions.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=a4841b214ebe36b744a229a82d701a1f0cfe6dcd;p=clang

[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.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@261778 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp
index f1b74ad554..a80a87ed43 100644
--- a/lib/Sema/SemaOverload.cpp
+++ b/lib/Sema/SemaOverload.cpp
@@ -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
diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu
index bc9fe2a1eb..bdbd550e95 100644
--- a/test/SemaCUDA/function-overload.cu
+++ b/test/SemaCUDA/function-overload.cu
@@ -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'}}