[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

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

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.

lib/CodeGen/CGCUDANV.cpp

assert(getDeviceSideName(CGF.CurFuncDecl) == CGF.CurFn->getName() ||
              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
     emitDeviceStubBodyLegacy(CGF, Args);
+
+  // 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 (CGF.getLangOpts().HIP)
+    CGF.CurFn->setName(CGF.CurFn->getName() + ".stub");
 }
 
 // CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
 // array and kernels are launched using cudaLaunchKernel().
 void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
                                             FunctionArgList &Args) {
   // Build the shadow stack entry at the very start of the function.
 

test/CodeGenCUDA/device-stub.cu

 // LNX: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
 // * Alias to global symbol containing the NVModuleID.
 // RDC: @__fatbinwrap[[MODULE_ID]] = alias { i32, i32, i8*, i8* }
 // RDC-SAME: { i32, i32, i8*, i8* }* @__[[PREFIX]]_fatbin_wrapper
 
 // Test that we build the correct number of calls to cudaSetupArgument followed
 // by a call to cudaLaunch.
 
-// LNX: define{{.*}}kernelfunc
+// CUDA-LABEL: define{{.*}}kernelfunc
+// HIP-LABEL: define{{.*}}@_Z10kernelfunciii.stub
 
 // New launch sequence stores arguments into local buffer and passes array of
 // pointers to them directly to cudaLaunchKernel
 // CUDA-NEW: alloca
 // CUDA-NEW: store
 // CUDA-NEW: store
 // CUDA-NEW: store
 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration