[CUDA] disable attribute-based overloading for __global__ functions.

[tra]

Summary

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

[jlebar]

Landed in rL261778.

lib/Sema/SemaOverload.cpp

if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads) {
     if (NewTarget == CFT_InvalidTarget || NewTarget == CFT_Global)
       return false;
 
     assert((OldTarget != CFT_InvalidTarget) && "Unexpected invalid target.");
 
     // 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
     // different CUDA target attributes.
     return NewTarget != OldTarget;
   }
 
   // The signatures match; this is not an overload.

test/SemaCUDA/function-overload.cu

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'}}