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" 4 5 // CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv 6 // CHECK-SPIRV: define{{.*}} spir_kernel void @_ZN1A6kernelEv 7 class A { 8 public: 9 static __global__ void kernel(){} 10 }; 11 12 // CHECK: define{{.*}} void @_Z10non_kernelv 13 // CHECK-SPIRV: define{{.*}} void @_Z10non_kernelv 14 __device__ void non_kernel(){} 15 16 // CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneli 17 // CHECK-SPIRV: define{{.*}} spir_kernel void @_Z6kerneli 18 __global__ void kernel(int x) { 19 non_kernel(); 20 } 21 22 // CHECK: define{{.*}} amdgpu_kernel void @_Z11EmptyKernelIvEvv 23 // CHECK-SPIRV: define{{.*}} spir_kernel void @_Z11EmptyKernelIvEvv 24 template <typename T> 25 __global__ void EmptyKernel(void) {} 26 27 struct Dummy { 28 /// Type definition of the EmptyKernel kernel entry point 29 typedef void (*EmptyKernelPtr)(); 30 EmptyKernelPtr Empty() { return EmptyKernel<void>; } 31 }; 32 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]*]] 35 template<class T> 36 __global__ void template_kernel(T x) {} 37 38 void launch(void *f); 39 40 int main() { 41 Dummy D; 42 launch((void*)A::kernel); 43 launch((void*)kernel); 44 launch((void*)template_kernel<A>); 45 launch((void*)D.Empty()); 46 return 0; 47 } 48 // CHECK: attributes #[[ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024" 49