[PCH, CUDA] Take CUDA attributes into account

[Artem-B]

During deserialization of CUDA AST we must consider CUDA target attributes to distinguish overloads from redeclarations.

Fixes #106394

[llvmbot]

@llvm/pr-subscribers-clang

Author: Artem Belevich (Artem-B)

Changes

During deserialization of CUDA AST we must consider CUDA target attributes to distinguish overloads from redeclarations.

---

Full diff: https://github.com/llvm/llvm-project/pull/125127.diff

2 Files Affected:

• (modified) clang/lib/AST/ASTContext.cpp (+12-1)
• (modified) clang/test/PCH/cuda-kernel-call.cu (+9)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index cd1bcb3b9a063d8..5e7461365c1ff91 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -7224,6 +7224,17 @@ 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 +7265,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 {
diff --git a/clang/test/PCH/cuda-kernel-call.cu b/clang/test/PCH/cuda-kernel-call.cu
index ffb0c1444fe69a6..32b192147fb36ea 100644
--- a/clang/test/PCH/cuda-kernel-call.cu
+++ b/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,19 @@ void kcall(void (*kp)()) {
 __global__ void kern() {
 }
 
+__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

[kadircet]

thanks a lot, LGTM!