[HIP] Add the interface deriving the stub name of device kernels.

[hliao]

Summary

• Revise the interface to derive the stub name and simplify the assertion of it.

[tra]

Is there particular reason you need to switch to this naming scheme?

One issue with this patch is that demanglers will no longer be able to deal with the name. While they do know to ignore .stub suffix, they can't deal with __device_stub_ prefix.
E.g:

% c++filt __device_stub___Z10kernelfuncIiEvv
__device_stub___Z10kernelfuncIiEvv
% c++filt _Z10kernelfuncIiEvv.stub
void kernelfunc<int>() [clone .stub]

clang/lib/CodeGen/CGCUDANV.cpp

222–226

I'm not sure I understand what exactly this assertion checks.
The condition appears to be true is host/device ABIs are different OR the name of the current function is the same as the (possibly mangled) device-side name + __device_stub_ prefix.

While the first part makes sense, I'm not sure I understand the name comparison part.
Could you tell me more and, maybe, add a comment explaining what's going on here.

789

I suspect return "__device_stub__" + Name; would do. StringRef will convert to std::string and copy elision should avoid unnecessary copy.

[hliao]

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

clang/lib/CodeGen/CGCUDANV.cpp

222–226

The second is to ensure, if, under the same ABI, kernel stub name derived from device-side name mangling should be the same the sub name generated from host-side, CGF.CurFn->getName() is the mangled named from host compilation

[hliao]

clang/lib/CodeGen/CGCUDANV.cpp

222–226

previous assertion expression gets the same goal, if ABI is different, the stub name from device-side should match the stub name from the host-side compilation. As we add a dedicated interface to the derive stub name, we could simplify the comparison to a single one.
Also, we put the simple condition checking ahead (a common practice) to reduce the overhead of string comparison

[tra]

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

clang/lib/CodeGen/CGCUDANV.cpp

222–226

This definitely needs a comment.

[hliao]

clang/lib/CodeGen/CGCUDANV.cpp

789

"devicestub__" + Name results in Twine, where not copy is generated. Only the final str() converts Twine into std::string involving copies. Otherwise, there's one copy from Name to std::string and another copy by std::string operator+, right?

[hliao]

In D63335#1543854, @tra wrote:

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

it's based on debugger people told me, with ".stub", the debugger still could find it match the original device kernel even though it could find both of them. But, they want to match the original one only and leave the stub one intentionally unmatched.

[hliao]

In D63335#1543854, @tra wrote:

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

Yeah, I understand that un-demangleable name causes lots of frustration. But, based on what I learned, CUDA generated the similar thing, e.g. __device_stub__Z15transformKernelPfiif is the stub function from cuda 10.1

[hliao]

In D63335#1543854, @tra wrote:

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

Is it OK for us to mangle __device_stub __ as the nested name into the original one, says, we prepend _ZN15__device_stub__E, so that we have _ZN15__device_stub__E10kernelfuncIiEvv

and

$ c++filt _ZN15__device_stub__E10kernelfuncIiEvv
__device_stub__(kernelfunc<int>, void, void)

[tra]

In D63335#1544019, @hliao wrote:

In D63335#1543854, @tra wrote:

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

it's based on debugger people told me, with ".stub", the debugger still could find it match the original device kernel even though it could find both of them. But, they want to match the original one only and leave the stub one intentionally unmatched.

Sorry, I still don't think I understand the reasons for this change. The stub and the kernel do have a different name now. I don't quite get it why the debugger can differentiate the names when they differ by prefix, but can't when they differ by suffix. It sounds like an attempt to work around a problem somewhere else.

Could you talk to the folks requesting the change and get more details on what exactly we need to do here and, more importantly, why.

[tra]

In D63335#1544026, @hliao wrote:

Is it OK for us to mangle __device_stub __ as the nested name into the original one, says, we prepend _ZN15__device_stub__E, so that we have _ZN15__device_stub__E10kernelfuncIiEvv

and

$ c++filt _ZN15__device_stub__E10kernelfuncIiEvv
__device_stub__(kernelfunc<int>, void, void)

I don't think it's a good idea. While it demangles to something, it's not what the demangled name should be. Stub's signature should match that of the kernel.

In D63335#1544021, @hliao wrote:

Yeah, I understand that un-demangleable name causes lots of frustration. But, based on what I learned, CUDA generated the similar thing, e.g. __device_stub__Z15transformKernelPfiif is the stub function from cuda 10.1

NVCC does a lot of things differently. It does not mean it's a good reason for us to copy *all* of their choices.
Let's figure out the underlying reasons for this change and then we can figure out about what's the right thing to do here.

[hliao]

In D63335#1544311, @tra wrote:

In D63335#1544019, @hliao wrote:

In D63335#1543854, @tra wrote:

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

it's based on debugger people told me, with ".stub", the debugger still could find it match the original device kernel even though it could find both of them. But, they want to match the original one only and leave the stub one intentionally unmatched.

Sorry, I still don't think I understand the reasons for this change. The stub and the kernel do have a different name now. I don't quite get it why the debugger can differentiate the names when they differ by prefix, but can't when they differ by suffix. It sounds like an attempt to work around a problem somewhere else.

Could you talk to the folks requesting the change and get more details on what exactly we need to do here and, more importantly, why.

But, after unmangling, debugger still could match both as they are almost identical excep the final variants, like clone. The debugger will set all locations matching that specified kernel name.

[tra]

In D63335#1544315, @hliao wrote:

Sorry, I still don't think I understand the reasons for this change. The stub and the kernel do have a different name now. I don't quite get it why the debugger can differentiate the names when they differ by prefix, but can't when they differ by suffix. It sounds like an attempt to work around a problem somewhere else.

Could you talk to the folks requesting the change and get more details on what exactly we need to do here and, more importantly, why.

But, after unmangling, debugger still could match both as they are almost identical excep the final variants, like clone. The debugger will set all locations matching that specified kernel name.

OK, so the real issue is that demangled name looks identical to debugger.
One way to deal with that is to , essentially, break mangling in compiler.
Another would be to teach debugger how to distinguish the stub from the kernel using additional information likely available to debugger (i.e. mangled name or the location of the symbol -- is it in the host binary or in the GPU binary).

I would argue that breaking mangling is not the best choice here.
I think debugger does have sufficient information to deal with this and that would be the right place to deal with the issue.

[hliao]

In D63335#1544320, @tra wrote:

In D63335#1544315, @hliao wrote:

Sorry, I still don't think I understand the reasons for this change. The stub and the kernel do have a different name now. I don't quite get it why the debugger can differentiate the names when they differ by prefix, but can't when they differ by suffix. It sounds like an attempt to work around a problem somewhere else.

Could you talk to the folks requesting the change and get more details on what exactly we need to do here and, more importantly, why.

But, after unmangling, debugger still could match both as they are almost identical excep the final variants, like clone. The debugger will set all locations matching that specified kernel name.

OK, so the real issue is that demangled name looks identical to debugger.
One way to deal with that is to , essentially, break mangling in compiler.
Another would be to teach debugger how to distinguish the stub from the kernel using additional information likely available to debugger (i.e. mangled name or the location of the symbol -- is it in the host binary or in the GPU binary).

I would argue that breaking mangling is not the best choice here.
I think debugger does have sufficient information to deal with this and that would be the right place to deal with the issue.

em, I did push the later as well, :(. OK, I will simplify the patch to change any functionality but move the calculation of device name into a common interface. So that, vendor could adjust that internally with minimal change. OK?

[hliao]

Just revise the interface for device kernel stubbing.

[tra]

In D63335#1544324, @hliao wrote:

I think debugger does have sufficient information to deal with this and that would be the right place to deal with the issue.

em, I did push the later as well, :(. OK, I will simplify the patch to change any functionality but move the calculation of device name into a common interface. So that, vendor could adjust that internally with minimal change. OK?

:-( Sorry about that. I realize how frustrating that can be.

Perhaps it's worth trying once more. You can argue that this change will have trouble being upstreamed without a good technical explanation why it must be done in the compiler. Perhaps they do have compelling reasons why it's hard to do in the debugger, but without specific details from their end it appears indistinguishable from a (possibly misguided) quick fix. It may help if you could get the debugger folks to chime in directly on the review.

[hliao]

In D63335#1544428, @tra wrote:

In D63335#1544324, @hliao wrote:

I think debugger does have sufficient information to deal with this and that would be the right place to deal with the issue.

em, I did push the later as well, :(. OK, I will simplify the patch to change any functionality but move the calculation of device name into a common interface. So that, vendor could adjust that internally with minimal change. OK?

:-( Sorry about that. I realize how frustrating that can be.

Perhaps it's worth trying once more. You can argue that this change will have trouble being upstreamed without a good technical explanation why it must be done in the compiler. Perhaps they do have compelling reasons why it's hard to do in the debugger, but without specific details from their end it appears indistinguishable from a (possibly misguided) quick fix. It may help if you could get the debugger folks to chime in directly on the review.

shall we review code refactoring first, so that that change could be just a single line change. Yes, I could post that later and drag in necessary stake holders.

[tra]

LGTM. This is a cleaner way to provide stub name tweaks.

clang/lib/CodeGen/CGCUDANV.cpp

223

says .... -> (e.g. ....)

225

... should be the same ...
or, perhaps, identical

clang/lib/CodeGen/CodeGenModule.cpp

1091

Adjust kernel stub mangling as we may need to be able to differentiate them from the kernel itself (e.g. for HIP).

[yaxunl]

LGTM. Thanks.

clang/lib/CodeGen/CGCUDANV.cpp

void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
                          unsigned Flags) override {
     DeviceVars.push_back({&Var, VD, Flags});
   }
 
   /// Creates module constructor function
   llvm::Function *makeModuleCtorFunction() override;
   /// Creates module destructor function
   llvm::Function *makeModuleDtorFunction() override;
+  /// Construct and return the stub name of a kernel.
+  std::string getDeviceStubName(llvm::StringRef Name) const override;
 };
 
 }
 
 std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const {
   if (CGM.getLangOpts().HIP)
     return ((Twine("hip") + Twine(FuncName)).str());
   return ((Twine("cuda") + Twine(FuncName)).str());
if (DeviceMC->shouldMangleDeclName(ND)) {
     DeviceSideName = Out.str();
   } 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());
+  assert((CGF.CGM.getContext().getAuxTargetInfo() &&
+          (CGF.CGM.getContext().getAuxTargetInfo()->getCXXABI() !=
+           CGF.CGM.getContext().getTargetInfo().getCXXABI())) ||
+         getDeviceStubName(getDeviceSideName(CGF.CurFuncDecl)) ==
+             CGF.CurFn->getName());
 
   EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                          CudaFeature::CUDA_USES_NEW_LAUNCH))
     emitDeviceStubBodyNew(CGF, Args);
   else
     emitDeviceStubBodyLegacy(CGF, Args);
 }
if (CGM.getLangOpts().HIP) {
     DtorBuilder.SetInsertPoint(ExitBlock);
   } else {
     DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
   }
   DtorBuilder.CreateRetVoid();
   return ModuleDtorFunc;
 }
 
+std::string CGNVCUDARuntime::getDeviceStubName(llvm::StringRef Name) const {
+  if (!CGM.getLangOpts().HIP)
+    return Name;
+  return std::move(("__device_stub__" + Name).str());
+}
+
 CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
   return new CGNVCUDARuntime(CGM);
 }

clang/lib/CodeGen/CGCUDARuntime.h

 // subclasses of this implement code generation for specific CUDA
 // runtime libraries.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_LIB_CODEGEN_CGCUDARUNTIME_H
 #define LLVM_CLANG_LIB_CODEGEN_CGCUDARUNTIME_H
 
+#include "llvm/ADT/StringRef.h"
+
 namespace llvm {
 class Function;
 class GlobalVariable;
 }
 
 namespace clang {
 
 class CUDAKernelCallExpr;
public:
 
   /// Constructs and returns a module initialization function or nullptr if it's
   /// not needed. Must be called after all kernels have been emitted.
   virtual llvm::Function *makeModuleCtorFunction() = 0;
 
   /// Returns a module cleanup function or nullptr if it's not needed.
   /// Must be called after ModuleCtorFunction
   virtual llvm::Function *makeModuleDtorFunction() = 0;
+
+  /// Construct and return the stub name of a kernel.
+  virtual std::string getDeviceStubName(llvm::StringRef Name) const = 0;
 };
 
 /// Creates an instance of a CUDA runtime class.
 CGCUDARuntime *CreateNVCUDARuntime(CodeGenModule &CGM);
 
 }
 }
 
 #endif

clang/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());
   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.
+  // Derive the kernel stub from CUDA runtime.
   if (auto *FD = dyn_cast<FunctionDecl>(GD.getDecl()))
-    if (getLangOpts().HIP && !getLangOpts().CUDAIsDevice &&
-        FD->hasAttr<CUDAGlobalAttr>())
-      MangledName = MangledName + ".stub";
+    if (!getLangOpts().CUDAIsDevice && FD->hasAttr<CUDAGlobalAttr>())
+      MangledName = getCUDARuntime().getDeviceStubName(MangledName);
 
   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();

clang/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]]()
+// CHECK: call void @[[STUB:__device_stub___Z10kernelfuncIiEvv]]()
 void hostfunc(void) { kernelfunc<int><<<1, 1>>>(); }
 
 // CHECK: define{{.*}}@[[STUB]]
 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[STUB]]
 
 // CHECK-LABEL: define{{.*}}@__hip_register_globals
 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[STUB]]