hliao created this revision. hliao added reviewers: yaxunl, tra. Herald added a project: clang. Herald added a subscriber: cfe-commits.
- Prefix kernel stub with `__device_stub__` to avoid potential symbol name conflicts in debugger. - Revise the interface to derive the stub name and simplify the assertion of it. Repository: rG LLVM Github Monorepo https://reviews.llvm.org/D63335 Files: clang/lib/CodeGen/CGCUDANV.cpp clang/lib/CodeGen/CGCUDARuntime.h clang/lib/CodeGen/CodeGenModule.cpp clang/test/CodeGenCUDA/kernel-stub-name.cu Index: clang/test/CodeGenCUDA/kernel-stub-name.cu =================================================================== --- clang/test/CodeGenCUDA/kernel-stub-name.cu +++ clang/test/CodeGenCUDA/kernel-stub-name.cu @@ -10,7 +10,7 @@ __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]] Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -1088,13 +1088,10 @@ 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(); Index: clang/lib/CodeGen/CGCUDARuntime.h =================================================================== --- clang/lib/CodeGen/CGCUDARuntime.h +++ clang/lib/CodeGen/CGCUDARuntime.h @@ -15,6 +15,8 @@ #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; @@ -63,6 +65,9 @@ /// 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. Index: clang/lib/CodeGen/CGCUDANV.cpp =================================================================== --- clang/lib/CodeGen/CGCUDANV.cpp +++ clang/lib/CodeGen/CGCUDANV.cpp @@ -132,6 +132,8 @@ 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; }; } @@ -217,10 +219,11 @@ 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(), @@ -780,6 +783,12 @@ 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); }
Index: clang/test/CodeGenCUDA/kernel-stub-name.cu =================================================================== --- clang/test/CodeGenCUDA/kernel-stub-name.cu +++ clang/test/CodeGenCUDA/kernel-stub-name.cu @@ -10,7 +10,7 @@ __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]] Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -1088,13 +1088,10 @@ 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(); Index: clang/lib/CodeGen/CGCUDARuntime.h =================================================================== --- clang/lib/CodeGen/CGCUDARuntime.h +++ clang/lib/CodeGen/CGCUDARuntime.h @@ -15,6 +15,8 @@ #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; @@ -63,6 +65,9 @@ /// 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. Index: clang/lib/CodeGen/CGCUDANV.cpp =================================================================== --- clang/lib/CodeGen/CGCUDANV.cpp +++ clang/lib/CodeGen/CGCUDANV.cpp @@ -132,6 +132,8 @@ 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; }; } @@ -217,10 +219,11 @@ 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(), @@ -780,6 +783,12 @@ 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); }
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits