[HIP] change kernel stub name

[yaxunl]

Summary

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

[tra]

My guess is that this is needed because HIP debugger can see symbols from both host and device executables at the same time. Is that so?

If that's the case, I guess HIP may have similar naming problem for __host__ __device__ foo() if it's used on both host and device.

lib/CodeGen/CGCUDANV.cpp

230–231 ↗

(On Diff #187815)

It may be worth adding a comment why kernel stub in HIP needs a different name.

[yaxunl]

In D58518#1406124, @tra wrote:

My guess is that this is needed because HIP debugger can see symbols from both host and device executables at the same time. Is that so?

If that's the case, I guess HIP may have similar naming problem for __host__ __device__ foo() if it's used on both host and device.

Probably, will fix it in seperate patch if it is true.

lib/CodeGen/CGCUDANV.cpp

230–231 ↗

(On Diff #187815)

will do when commit

[t-tye]

Yes this relates to supporting the debugger.

For the same function being present on both host and device, having the same name is correct as the debugger must set a breakpoint at both places. This is similar to needing to set a breakpoint at every place a function is inlined.

[tra]

In D58518#1406202, @t-tye wrote:

Yes this relates to supporting the debugger.

For the same function being present on both host and device, having the same name is correct as the debugger must set a breakpoint at both places. This is similar to needing to set a breakpoint at every place a function is inlined.

I'm confused. Are you saying that HIP does *not* need a different name for the stub then?

[t-tye]

To clarify, I am saying that the stub does have a different name since it is conceptually part of the implementation of doing the call to the device function implementation, and is not in fact the the device function being called itself. However, when we generate code for a function that is present on both the host and device, both copies of the code are for the same source level function and so can have the same symbol name (which was a question that was asked).

[tra]

In D58518#1406274, @t-tye wrote:

To clarify, I am saying that the stub does have a different name since it is conceptually part of the implementation of doing the call to the device function implementation, and is not in fact the the device function being called itself. However, when we generate code for a function that is present on both the host and device, both copies of the code are for the same source level function and so can have the same symbol name (which was a question that was asked)

Got it. Agreed.

[yaxunl]

Fixed regressions.

[tra]

lib/CodeGen/CodeGenModule.cpp

1059 ↗

(On Diff #187980)

Changing mangled name exposes this change to a wider scope of potential issues. Is the mangled name still valid after this change? I.e. will external demanglers have problem with it? Is . a valid symbol in mangled names on all platforms we support?

I think changing the name here is way too late and we should figure out how to change the stub name when we generate it.

@echristo Eric, what do you think?

[yaxunl]

lib/CodeGen/CodeGenModule.cpp

1059 ↗

(On Diff #187980)

The external demangler can still demangle this name. e.g. c++filt will demangle this name and add [clone .stub] after that.

As far as I can see this function is only called in codegen to map FunctionDecl names to LLVM function names. I've tried this change with real ML frameworks and it works.

Changing at this place is not too late. The stub function name is requested at multiple places in codegen, not just at the emitting of stub function definition. For template kernel function, the emitting of stub function definition is deferred after emitting of the call of the stub function. Basically, codegen needs to find the corresponding LLVM stub function by getMangledName first, then by GetOrCreateLLVMFunction. If we do not change getMangledName, codegen will not get the correct stub function name consistently at all places. That's why the previous patch does not work.

[tra]

lib/CodeGen/CodeGenModule.cpp

1059 ↗

(On Diff #187980)

I stand corrected. @jyknight and @bkramer pointed out that appending .WHATEVER is currently used for cloning functions and should be OK to do.

cfe/trunk/lib/CodeGen/CGCUDANV.cpp

std::string CGNVCUDARuntime::getDeviceSideName(const Decl *D) {
   } else
     DeviceSideName = ND->getIdentifier()->getName();
   return DeviceSideName;
 }
 
 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());
 
   EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                          CudaFeature::CUDA_USES_NEW_LAUNCH))
     emitDeviceStubBodyNew(CGF, Args);
   else

cfe/trunk/lib/CodeGen/CodeGenModule.cpp

StringRef CodeGenModule::getMangledName(GlobalDecl GD) {
   }
 
   auto FoundName = MangledDeclNames.find(CanonicalGD);
   if (FoundName != MangledDeclNames.end())
     return FoundName->second;
 
   // 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();
 }
 
 StringRef CodeGenModule::getBlockMangledName(GlobalDecl GD,
                                              const BlockDecl *BD) {
   MangleContext &MangleCtx = getCXXABI().getMangleContext();
   const Decl *D = GD.getDecl();
 

cfe/trunk/test/CodeGenCUDA/kernel-stub-name.cu

+// 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]]