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)
kernel(A 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)
memberKernel(A x)21 static __global__ void memberKernel(A x){}
templateMemberKernel(T x)22 template<typename T> static __global__ void templateMemberKernel(T x) {}
23 };
24
25
26 template <typename T>
templateKernel(T x)27 __global__ void templateKernel(T x) {}
28
29 void launch(void*);
30
test()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