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