[PCH, CUDA] Take CUDA attributes into account

clang/lib/AST/ASTContext.cpp

@@ -7224,6 +7224,16 @@ static bool isSameQualifier(const NestedNameSpecifier *X,
   return !PX && !PY;
 }
 
+static bool hasSameCudaAttrs(const FunctionDecl *A, const FunctionDecl *B) {
+  if (!A->getASTContext().getLangOpts().CUDA)
+    return true; // Target attributes are overloadable in CUDA compilation only.
+  if (A->hasAttr<CUDADeviceAttr>() != B->hasAttr<CUDADeviceAttr>())
+    return false;
+  if (A->hasAttr<CUDADeviceAttr>() && B->hasAttr<CUDADeviceAttr>())
+    return A->hasAttr<CUDAHostAttr>() == B->hasAttr<CUDAHostAttr>();
+  return true; // unattributed and __host__ functions are the same.
+}
+
 /// Determine whether the attributes we can overload on are identical for A and
 /// B. Will ignore any overloadable attrs represented in the type of A and B.
 static bool hasSameOverloadableAttrs(const FunctionDecl *A,
@@ -7254,7 +7264,7 @@ static bool hasSameOverloadableAttrs(const FunctionDecl *A,
     if (Cand1ID != Cand2ID)
       return false;
   }
-  return true;
+  return hasSameCudaAttrs(A, B);
 }
 
 bool ASTContext::isSameEntity(const NamedDecl *X, const NamedDecl *Y) const {

clang/test/PCH/cuda-kernel-call.cu

@@ -1,5 +1,7 @@
 // RUN: %clang_cc1 -emit-pch -o %t %s
 // RUN: %clang_cc1 -include-pch %t -fsyntax-only %s 
+// RUN: %clang_cc1 -emit-pch -fcuda-is-device -o %t-device %s
+// RUN: %clang_cc1 -fcuda-is-device -include-pch %t-device -fsyntax-only %s
 
 #ifndef HEADER
 #define HEADER
@@ -14,12 +16,21 @@ void kcall(void (*kp)()) {
 __global__ void kern() {
 }
 
+// Make sure that target overloaded functions remain
+// available as overloads after PCH deserialization.
+__host__ int overloaded_func();
+__device__ int overloaded_func();
+
 #else
 // Using the header.
 
 void test() {
   kcall(kern);
   kern<<<1, 1>>>();
+  overloaded_func();
 }
 
+__device__ void test () {
+  overloaded_func();
+}
 #endif