1 // RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
2 // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
3 #include "Inputs/cuda.h"
5 // CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv
6 // CHECK-SPIRV: define{{.*}} spir_kernel void @_ZN1A6kernelEv
9 static __global__ void kernel(){}
12 // CHECK: define{{.*}} void @_Z10non_kernelv
13 // CHECK-SPIRV: define{{.*}} void @_Z10non_kernelv
14 __device__ void non_kernel(){}
16 // CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneli
17 // CHECK-SPIRV: define{{.*}} spir_kernel void @_Z6kerneli
18 __global__ void kernel(int x) {
22 // CHECK: define{{.*}} amdgpu_kernel void @_Z11EmptyKernelIvEvv
23 // CHECK-SPIRV: define{{.*}} spir_kernel void @_Z11EmptyKernelIvEvv
25 __global__ void EmptyKernel(void) {}
28 /// Type definition of the EmptyKernel kernel entry point
29 typedef void (*EmptyKernelPtr)();
30 EmptyKernelPtr Empty() { return EmptyKernel<void>; }
33 // CHECK: define{{.*}} amdgpu_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
34 // CHECK-SPIRV: define{{.*}} spir_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
36 __global__ void template_kernel(T x) {}
42 launch((void*)A::kernel);
43 launch((void*)kernel);
44 launch((void*)template_kernel<A>);
45 launch((void*)D.Empty());
48 // CHECK: attributes #[[ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"