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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
