1 // RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
2 #include "Inputs/cuda.h"
4 // CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv
7 static __global__ void kernel(){}
10 // CHECK: define{{.*}} void @_Z10non_kernelv
11 __device__ void non_kernel(){}
13 // CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneli
14 __global__ void kernel(int x) {
18 // CHECK: define{{.*}} amdgpu_kernel void @_Z11EmptyKernelIvEvv
20 __global__ void EmptyKernel(void) {}
23 /// Type definition of the EmptyKernel kernel entry point
24 typedef void (*EmptyKernelPtr)();
25 EmptyKernelPtr Empty() { return EmptyKernel<void>; }
28 // CHECK: define{{.*}} amdgpu_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
30 __global__ void template_kernel(T x) {}
36 launch((void*)A::kernel);
37 launch((void*)kernel);
38 launch((void*)template_kernel<A>);
39 launch((void*)D.Empty());
42 // CHECK: attributes #[[ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"