1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
2 // RUN:     -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCN %s
3 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda- -fcuda-is-device \
4 // RUN:     -emit-llvm %s -o - | FileCheck -check-prefix=NVPTX %s
5 #include "Inputs/cuda.h"
6 
7 struct A {
8   int a[32];
9 };
10 
11 // AMDGCN: define amdgpu_kernel void @_Z6kernel1A(%struct.A addrspace(4)* byref(%struct.A) align 4 %{{.+}})
12 // NVPTX: define void @_Z6kernel1A(%struct.A* byval(%struct.A) align 4 %x)
13 __global__ void kernel(A x) {
14 }
15 
16 class Kernel {
17 public:
18   // AMDGCN: define amdgpu_kernel void @_ZN6Kernel12memberKernelE1A(%struct.A addrspace(4)* byref(%struct.A) align 4 %{{.+}})
19   // NVPTX: define void @_ZN6Kernel12memberKernelE1A(%struct.A* byval(%struct.A) align 4 %x)
20   static __global__ void memberKernel(A x){}
21   template<typename T> static __global__ void templateMemberKernel(T x) {}
22 };
23 
24 
25 template <typename T>
26 __global__ void templateKernel(T x) {}
27 
28 void launch(void*);
29 
30 void test() {
31   Kernel K;
32   // AMDGCN: define amdgpu_kernel void @_Z14templateKernelI1AEvT_(%struct.A addrspace(4)* byref(%struct.A) align 4 %{{.+}}
33   // NVPTX: define void @_Z14templateKernelI1AEvT_(%struct.A* byval(%struct.A) align 4 %x)
34   launch((void*)templateKernel<A>);
35 
36   // AMDGCN: define amdgpu_kernel void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A addrspace(4)* byref(%struct.A) align 4 %{{.+}}
37   // NVPTX: define void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A* byval(%struct.A) align 4 %x)
38   launch((void*)Kernel::templateMemberKernel<A>);
39 }
40