[HIP] change kernel stub name

Add .stub to kernel stub function name so that it is different from kernel
name in device code. This is necessary to let debugger find correct symbol
for kernel.

Differential Revision: https://reviews.llvm.org/D58518

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@354948 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/CodeGen/CGCUDANV.cpp b/lib/CodeGen/CGCUDANV.cpp
index 62661039a32a1a089b287d44cbb17aa41b807482..0602601f5e11a7cfd2316c832ac252c7f24769bb 100644 (file)
--- a/lib/CodeGen/CGCUDANV.cpp
+++ b/lib/CodeGen/CGCUDANV.cpp
@@ -218,6 +218,7 @@ std::string CGNVCUDARuntime::getDeviceSideName(const Decl *D) {
 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                      FunctionArgList &Args) {
   assert(getDeviceSideName(CGF.CurFuncDecl) == CGF.CurFn->getName() ||
+         getDeviceSideName(CGF.CurFuncDecl) + ".stub" == CGF.CurFn->getName() ||
          CGF.CGM.getContext().getTargetInfo().getCXXABI() !=
              CGF.CGM.getContext().getAuxTargetInfo()->getCXXABI());
 

diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp
index 65116dba095b0f57356c5fe6d9ce98bdee46568b..8ea8128ceabb36704d65c519d0ce88ceed570425 100644 (file)
--- a/lib/CodeGen/CodeGenModule.cpp
+++ b/lib/CodeGen/CodeGenModule.cpp
@@ -1048,8 +1048,17 @@ StringRef CodeGenModule::getMangledName(GlobalDecl GD) {
 
   // Keep the first result in the case of a mangling collision.
   const auto *ND = cast<NamedDecl>(GD.getDecl());
-  auto Result =
-      Manglings.insert(std::make_pair(getMangledNameImpl(*this, GD, ND), GD));
+  std::string MangledName = getMangledNameImpl(*this, GD, ND);
+
+  // Postfix kernel stub names with .stub to differentiate them from kernel
+  // names in device binaries. This is to facilitate the debugger to find
+  // the correct symbols for kernels in the device binary.
+  if (auto *FD = dyn_cast<FunctionDecl>(GD.getDecl()))
+    if (getLangOpts().HIP && !getLangOpts().CUDAIsDevice &&
+        FD->hasAttr<CUDAGlobalAttr>())
+      MangledName = MangledName + ".stub";
+
+  auto Result = Manglings.insert(std::make_pair(MangledName, GD));
   return MangledDeclNames[CanonicalGD] = Result.first->first();
 }
 

diff --git a/test/CodeGenCUDA/kernel-stub-name.cu b/test/CodeGenCUDA/kernel-stub-name.cu
new file mode 100644 (file)
index 0000000..539d7ee
--- /dev/null
+++ b/test/CodeGenCUDA/kernel-stub-name.cu
@@ -0,0 +1,20 @@
+// RUN: echo "GPU binary would be here" > %t
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN:     -fcuda-include-gpubinary %t -o - -x hip\
+// RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CHECK
+
+#include "Inputs/cuda.h"
+
+template<class T>
+__global__ void kernelfunc() {}
+
+// CHECK-LABEL: define{{.*}}@_Z8hostfuncv()
+// CHECK: call void @[[STUB:_Z10kernelfuncIiEvv.stub]]()
+void hostfunc(void) { kernelfunc<int><<<1, 1>>>(); }
+
+// CHECK: define{{.*}}@[[STUB]]
+// CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[STUB]]
+
+// CHECK-LABEL: define{{.*}}@__hip_register_globals
+// CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[STUB]]