From 2e53b8068e17d5347f6fc63e68faa2bf8531fb28 Mon Sep 17 00:00:00 2001
From: Yaxun Liu <Yaxun.Liu@amd.com>
Date: Wed, 27 Feb 2019 02:02:52 +0000
Subject: [PATCH] [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
---
 lib/CodeGen/CGCUDANV.cpp             |  1 +
 lib/CodeGen/CodeGenModule.cpp        | 13 +++++++++++--
 test/CodeGenCUDA/kernel-stub-name.cu | 20 ++++++++++++++++++++
 3 files changed, 32 insertions(+), 2 deletions(-)
 create mode 100644 test/CodeGenCUDA/kernel-stub-name.cu

diff --git a/lib/CodeGen/CGCUDANV.cpp b/lib/CodeGen/CGCUDANV.cpp
index 62661039a3..0602601f5e 100644
--- 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 65116dba09..8ea8128cea 100644
--- 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
index 0000000000..539d7eec1b
--- /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]]
-- 
2.50.1