r354948 - [HIP] change kernel stub name

[Yaxun Liu via cfe-commits]

Author: yaxunl
Date: Tue Feb 26 18:02:52 2019
New Revision: 354948

URL: http://llvm.org/viewvc/llvm-project?rev=354948&view=rev
Log:
[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

Added:
    cfe/trunk/test/CodeGenCUDA/kernel-stub-name.cu
Modified:
    cfe/trunk/lib/CodeGen/CGCUDANV.cpp
    cfe/trunk/lib/CodeGen/CodeGenModule.cpp

Modified: cfe/trunk/lib/CodeGen/CGCUDANV.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCUDANV.cpp?rev=354948&r1=354947&r2=354948&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCUDANV.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCUDANV.cpp Tue Feb 26 18:02:52 2019
@@ -218,6 +218,7 @@ std::string CGNVCUDARuntime::getDeviceSi
 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());
 

Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=354948&r1=354947&r2=354948&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Tue Feb 26 18:02:52 2019
@@ -1048,8 +1048,17 @@ StringRef CodeGenModule::getMangledName(
 
   // 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();
 }
 

Added: cfe/trunk/test/CodeGenCUDA/kernel-stub-name.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/kernel-stub-name.cu?rev=354948&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/kernel-stub-name.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/kernel-stub-name.cu Tue Feb 26 18:02:52 2019
@@ -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]]