[PATCH] D113491: [HIP] Fix device stub name for Windows

[Yaxun Liu via Phabricator via cfe-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
yaxunl marked an inline comment as done.
Closed by commit rG38211bbab1d9: [HIP] Fix device stub name for Windows (authored by yaxunl).
Herald added a project: clang.

Changed prior to commit:
  https://reviews.llvm.org/D113491?vs=385836&id=389228#toc

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D113491/new/

https://reviews.llvm.org/D113491

Files:
  clang/include/clang/AST/GlobalDecl.h
  clang/lib/AST/MicrosoftMangle.cpp
  clang/test/CodeGenCUDA/kernel-stub-name.cu


Index: clang/test/CodeGenCUDA/kernel-stub-name.cu
===================================================================
--- clang/test/CodeGenCUDA/kernel-stub-name.cu
+++ 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 at ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?nskernel at ns@@YAXXZ"]], align 8
-// MSVC: @[[HTKERN:"\?\?\$kernelfunc at H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$kernelfunc at H@@YAXXZ.*"]], comdat, align 8
+// MSVC: @[[HNSKERN:"\?nskernel at ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?__device_stub__nskernel at ns@@YAXXZ"]], align 8
+// MSVC: @[[HTKERN:"\?\?\$kernelfunc at H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$__device_stub__kernelfunc at 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 @@
 // 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>>>();
Index: clang/lib/AST/MicrosoftMangle.cpp
===================================================================
--- clang/lib/AST/MicrosoftMangle.cpp
+++ clang/lib/AST/MicrosoftMangle.cpp
@@ -962,7 +962,19 @@
   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;
       }
 
Index: clang/include/clang/AST/GlobalDecl.h
===================================================================
--- clang/include/clang/AST/GlobalDecl.h
+++ 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 @@
   }
 
   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());
   }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D113491.389228.patch
Type: text/x-patch
Size: 3517 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20211123/85961451/attachment.bin>