diff options
author | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-03-29 15:02:08 +0000 |
---|---|---|
committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-03-29 15:02:08 +0000 |
commit | 96325393f64511a3fae168b7f3c3d31733e82b45 (patch) | |
tree | 6363344453081aa3fc7dcf764d322669d99ef900 /test/CodeGenCUDA | |
parent | 8a1c33e1b7dbdf4162ae4c8aea1145c5d5cbd129 (diff) |
Set calling convention for CUDA kernel
This patch sets target specific calling convention for CUDA kernels in IR.
Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.
Differential Revision: https://reviews.llvm.org/D44747
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@328795 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r-- | test/CodeGenCUDA/kernel-amdgcn.cu | 29 |
1 files changed, 29 insertions, 0 deletions
diff --git a/test/CodeGenCUDA/kernel-amdgcn.cu b/test/CodeGenCUDA/kernel-amdgcn.cu new file mode 100644 index 0000000000..a7369f901b --- /dev/null +++ b/test/CodeGenCUDA/kernel-amdgcn.cu @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s +#include "Inputs/cuda.h" + +// CHECK: define amdgpu_kernel void @_ZN1A6kernelEv +class A { +public: + static __global__ void kernel(){} +}; + +// CHECK: define void @_Z10non_kernelv +__device__ void non_kernel(){} + +// CHECK: define amdgpu_kernel void @_Z6kerneli +__global__ void kernel(int x) { + non_kernel(); +} + +// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_ +template<class T> +__global__ void template_kernel(T x) {} + +void launch(void *f); + +int main() { + launch((void*)A::kernel); + launch((void*)kernel); + launch((void*)template_kernel<A>); + return 0; +} |