49 lines · plain
1// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s2// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV3#include "Inputs/cuda.h"4 5// CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv6// CHECK-SPIRV: define{{.*}} spir_kernel void @_ZN1A6kernelEv7class A {8public:9 static __global__ void kernel(){}10};11 12// CHECK: define{{.*}} void @_Z10non_kernelv13// CHECK-SPIRV: define{{.*}} void @_Z10non_kernelv14__device__ void non_kernel(){}15 16// CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneli17// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z6kerneli18__global__ void kernel(int x) {19 non_kernel();20}21 22// CHECK: define{{.*}} amdgpu_kernel void @_Z11EmptyKernelIvEvv23// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z11EmptyKernelIvEvv24template <typename T>25__global__ void EmptyKernel(void) {}26 27struct Dummy {28 /// Type definition of the EmptyKernel kernel entry point29 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]*]]35template<class T>36__global__ void template_kernel(T x) {}37 38void launch(void *f);39 40int 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