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]] _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
