Author: Yaxun (Sam) Liu Date: 2021-11-23T12:03:49-05:00 New Revision: 38211bbab1d949f682271abba0171424a5a335ab
URL: https://github.com/llvm/llvm-project/commit/38211bbab1d949f682271abba0171424a5a335ab DIFF: https://github.com/llvm/llvm-project/commit/38211bbab1d949f682271abba0171424a5a335ab.diff LOG: [HIP] Fix device stub name for Windows This is a follow up of https://reviews.llvm.org/D68578 where device stub name is changed for Itanium mangling but not Microsoft mangling. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D113491 Added: Modified: clang/include/clang/AST/GlobalDecl.h clang/lib/AST/MicrosoftMangle.cpp clang/test/CodeGenCUDA/kernel-stub-name.cu Removed: ################################################################################ diff --git a/clang/include/clang/AST/GlobalDecl.h b/clang/include/clang/AST/GlobalDecl.h index 8cb56fb4ae90b..88abba28c991d 100644 --- a/clang/include/clang/AST/GlobalDecl.h +++ b/clang/include/clang/AST/GlobalDecl.h @@ -18,6 +18,7 @@ #include "clang/AST/DeclCXX.h" #include "clang/AST/DeclObjC.h" #include "clang/AST/DeclOpenMP.h" +#include "clang/AST/DeclTemplate.h" #include "clang/Basic/ABI.h" #include "clang/Basic/LLVM.h" #include "llvm/ADT/DenseMapInfo.h" @@ -129,8 +130,12 @@ class GlobalDecl { } KernelReferenceKind getKernelReferenceKind() const { - assert(isa<FunctionDecl>(getDecl()) && - cast<FunctionDecl>(getDecl())->hasAttr<CUDAGlobalAttr>() && + assert(((isa<FunctionDecl>(getDecl()) && + cast<FunctionDecl>(getDecl())->hasAttr<CUDAGlobalAttr>()) || + (isa<FunctionTemplateDecl>(getDecl()) && + cast<FunctionTemplateDecl>(getDecl()) + ->getTemplatedDecl() + ->hasAttr<CUDAGlobalAttr>())) && "Decl is not a GPU kernel!"); return static_cast<KernelReferenceKind>(Value.getInt()); } diff --git a/clang/lib/AST/MicrosoftMangle.cpp b/clang/lib/AST/MicrosoftMangle.cpp index beca1be78d99e..79a448a2435cd 100644 --- a/clang/lib/AST/MicrosoftMangle.cpp +++ b/clang/lib/AST/MicrosoftMangle.cpp @@ -962,7 +962,19 @@ void MicrosoftCXXNameMangler::mangleUnqualifiedName(GlobalDecl GD, switch (Name.getNameKind()) { case DeclarationName::Identifier: { if (const IdentifierInfo *II = Name.getAsIdentifierInfo()) { - mangleSourceName(II->getName()); + bool IsDeviceStub = + ND && + ((isa<FunctionDecl>(ND) && ND->hasAttr<CUDAGlobalAttr>()) || + (isa<FunctionTemplateDecl>(ND) && + cast<FunctionTemplateDecl>(ND) + ->getTemplatedDecl() + ->hasAttr<CUDAGlobalAttr>())) && + GD.getKernelReferenceKind() == KernelReferenceKind::Stub; + if (IsDeviceStub) + mangleSourceName( + (llvm::Twine("__device_stub__") + II->getName()).str()); + else + mangleSourceName(II->getName()); break; } diff --git a/clang/test/CodeGenCUDA/kernel-stub-name.cu b/clang/test/CodeGenCUDA/kernel-stub-name.cu index 8e82c3612e323..71856bee98902 100644 --- a/clang/test/CodeGenCUDA/kernel-stub-name.cu +++ b/clang/test/CodeGenCUDA/kernel-stub-name.cu @@ -28,8 +28,8 @@ // GNU: @[[HDKERN:_Z11kernel_declv]] = external constant void ()*, align 8 // MSVC: @[[HCKERN:ckernel]] = dso_local constant void ()* @[[CSTUB:__device_stub__ckernel]], align 8 -// MSVC: @[[HNSKERN:"\?nskernel@ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?nskernel@ns@@YAXXZ"]], align 8 -// MSVC: @[[HTKERN:"\?\?\$kernelfunc@H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$kernelfunc@H@@YAXXZ.*"]], comdat, align 8 +// MSVC: @[[HNSKERN:"\?nskernel@ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?__device_stub__nskernel@ns@@YAXXZ"]], align 8 +// MSVC: @[[HTKERN:"\?\?\$kernelfunc@H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$__device_stub__kernelfunc@H@@YAXXZ.*"]], comdat, align 8 // MSVC: @[[HDKERN:"\?kernel_decl@@YAXXZ.*"]] = external dso_local constant void ()*, align 8 extern "C" __global__ void ckernel() {} @@ -69,7 +69,7 @@ extern "C" void launch(void *kern); // CHECK: call void @[[NSSTUB]]() // CHECK: call void @[[TSTUB]]() // GNU: call void @[[DSTUB:_Z26__device_stub__kernel_declv]]() -// MSVC: call void @[[DSTUB:"\?kernel_decl@@YAXXZ"]]() +// MSVC: call void @[[DSTUB:"\?__device_stub__kernel_decl@@YAXXZ"]]() extern "C" void fun1(void) { ckernel<<<1, 1>>>(); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
