llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Artem Belevich (Artem-B) <details> <summary>Changes</summary> 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 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 `````````` </details> https://github.com/llvm/llvm-project/pull/125127 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
