llvm-project/clang/test/CodeGenCUDA
Yaxun Liu c18e9ecd4f [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
2019-02-14 02:00:09 +00:00
..
Inputs [CUDA] add support for the new kernel launch API in CUDA-9.2+. 2019-01-31 21:34:03 +00:00
address-spaces.cu [CUDA] Let device-side shared variables be initialized with undef 2018-04-02 17:38:24 +00:00
alias.cu Disable emitting static extern C aliases for amdgcn target for CUDA 2018-03-29 14:50:00 +00:00
amdgpu-kernel-attrs.cu [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes 2018-06-12 23:58:59 +00:00
builtins-amdgcn.cu Try to make builtin address space declarations not useless 2018-08-02 12:14:28 +00:00
convergent.cu [FileCheck] Add -allow-deprecated-dag-overlap to failing clang tests 2018-07-11 20:26:20 +00:00
cuda-builtin-vars.cu
debug-info-address-class.cu [DEBUG_INFO][NVPTX] Generate correct data about variable address class. 2019-02-05 19:45:57 +00:00
device-stub.cu [CUDA][HIP] Use device side kernel and variable names when registering them 2019-02-14 02:00:09 +00:00
device-var-init.cu [CUDA] Make all host-side shadows of device-side variables undef. 2018-12-13 21:43:04 +00:00
device-vtable.cu [CUDA][HIP] Do not emit type info when compiling for device 2018-06-05 15:11:02 +00:00
filter-decl.cu Really fix test on windows. 2018-02-23 19:38:41 +00:00
flush-denormals.cu [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn 2018-07-21 02:02:22 +00:00
fp-contract.cu
function-overload.cu
kernel-amdgcn.cu [CUDA] Set LLVM calling convention for CUDA kernel 2018-04-20 17:01:03 +00:00
kernel-args-alignment.cu [CUDA] add support for the new kernel launch API in CUDA-9.2+. 2019-01-31 21:34:03 +00:00
kernel-args.cu [CUDA][HIP] Set kernel calling convention before arrange function 2018-06-12 00:16:33 +00:00
kernel-call.cu [CUDA] add support for the new kernel launch API in CUDA-9.2+. 2019-01-31 21:34:03 +00:00
launch-bounds.cu
library-builtin.cu [CUDA] CUDA has no device-side library builtins. 2018-01-23 19:08:18 +00:00
link-device-bitcode.cu Rename -mlink-cuda-bitcode to -mlink-builtin-bitcode 2018-08-20 18:16:48 +00:00
llvm-used.cu
nothrow.cu
printf-aggregate.cu
printf.cu
propagate-metadata.cu Rename -mlink-cuda-bitcode to -mlink-builtin-bitcode 2018-08-20 18:16:48 +00:00
ptx-kernels.cu
types.cu Do not copy long double and 128-bit fp format from aux target for AMDGPU 2019-01-31 21:57:51 +00:00
usual-deallocators.cu [CUDA] Ignore uncallable functions when we check for usual deallocators. 2018-09-21 17:29:33 +00:00