diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2019-02-27 02:02:52 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2019-02-27 02:02:52 +0000 |
| commit | e739ac0e255597d818c907223034ddf3bc18a593 (patch) | |
| tree | 7c4ea703ba9503f46b9233991fc0a27418209f03 /clang/test/CodeGenCUDA | |
| parent | c4eff2111abd4785c70578631179e77c88513fb1 (diff) | |
| download | bcm5719-llvm-e739ac0e255597d818c907223034ddf3bc18a593.tar.gz bcm5719-llvm-e739ac0e255597d818c907223034ddf3bc18a593.zip | |
[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
llvm-svn: 354948
Diffstat (limited to 'clang/test/CodeGenCUDA')
| -rw-r--r-- | clang/test/CodeGenCUDA/kernel-stub-name.cu | 20 |
1 files changed, 20 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/kernel-stub-name.cu b/clang/test/CodeGenCUDA/kernel-stub-name.cu new file mode 100644 index 00000000000..539d7eec1ba --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-stub-name.cu @@ -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]] |
