From c18e9ecd4fc1f8bce7cf1ad0a63cb64a3fc40040 Mon Sep 17 00:00:00 2001 From: Yaxun Liu Date: Thu, 14 Feb 2019 02:00:09 +0000 Subject: [CUDA][HIP] Use device side kernel and variable names when registering them __hipRegisterFunction and __hipRegisterVar need to accept device side kernel and variable names so that HIP runtime can associate kernel stub functions in host code with kernel symbols in fat binaries, and associate shadow variables in host code with device variables in fat binaries. Currently, clang assumes kernel functions and device variables have the same name as the kernel stub functions and shadow variables. However, when host is compiled in windows with MSVC C++ ABI and device is compiled with Itanium C++ ABI (e.g. AMDGPU), kernels and device symbols in fat binary are mangled differently than host. This patch gets the device side kernel and variable name by mangling them in the mangle context of aux target. Differential Revision: https://reviews.llvm.org/D58163 llvm-svn: 354004 --- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'clang/lib/CodeGen/CodeGenModule.cpp') diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ece26deff017..972d2afa8e6c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3635,7 +3635,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // Extern global variables will be registered in the TU where they are // defined. if (!D->hasExternalStorage()) - getCUDARuntime().registerDeviceVar(*GV, Flags); + getCUDARuntime().registerDeviceVar(D, *GV, Flags); } else if (D->hasAttr()) // __shared__ variables are odd. Shadows do get created, but // they are not registered with the CUDA runtime, so they -- cgit v1.2.3