diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index cd1bcb3b9a063..7f54fe36a0e08 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() != B->hasAttr()) + return false; + if (A->hasAttr() && B->hasAttr()) + return A->hasAttr() == B->hasAttr(); + 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 ffb0c1444fe69..da9d81c531c41 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