caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

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

Fixes #106394
This commit is contained in:
Artem Belevich
2025-02-03 10:27:11 -08:00
committed by GitHub
parent 9fbd5fbcc6
commit db60244519
2 changed files with 22 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
clang/lib/AST/ASTContext.cpp
View File

@@ -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 {

11
clang/test/PCH/cuda-kernel-call.cu
View File

@@ -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