[clang] [PCH, CUDA] Take CUDA attributes into account (PR #125127)

[Artem Belevich via cfe-commits]

https://github.com/Artem-B created https://github.com/llvm/llvm-project/pull/125127

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

>From d6ff7dc3dcc5ca677a811888f5e67cc5aaad9d8f Mon Sep 17 00:00:00 2001
From: Artem Belevich <tra at google.com>
Date: Thu, 30 Jan 2025 14:30:22 -0800
Subject: [PATCH] [PCH, CUDA] Take CUDA attributes into account

During deserialization of CUDA AST we must consider CUDA target
attributes to distinguish overloads from redeclarations.
---
 clang/lib/AST/ASTContext.cpp       | 13 ++++++++++++-
 clang/test/PCH/cuda-kernel-call.cu |  9 +++++++++
 2 files changed, 21 insertions(+), 1 deletion(-)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index cd1bcb3b9a063d..5e7461365c1ff9 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 ffb0c1444fe69a..32b192147fb36e 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