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

[via cfe-commits]

Author: Artem Belevich
Date: 2025-02-03T10:27:11-08:00
New Revision: db60244519023a2b083caa3ed3a27a6b59eb03d8

URL: https://github.com/llvm/llvm-project/commit/db60244519023a2b083caa3ed3a27a6b59eb03d8
DIFF: https://github.com/llvm/llvm-project/commit/db60244519023a2b083caa3ed3a27a6b59eb03d8.diff

LOG: [PCH, CUDA] Take CUDA attributes into account (#125127)

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

Fixes #106394

Added: 
    

Modified: 
    clang/lib/AST/ASTContext.cpp
    clang/test/PCH/cuda-kernel-call.cu

Removed: 
    


################################################################################
diff  --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 2dc96691f1da70..f3aedbc39d12ae 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/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 {

diff  --git a/clang/test/PCH/cuda-kernel-call.cu b/clang/test/PCH/cuda-kernel-call.cu
index ffb0c1444fe69a..da9d81c531c415 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,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