1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
|
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
#include "Inputs/cuda.h"
// CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv
// CHECK-SPIRV: define{{.*}} spir_kernel void @_ZN1A6kernelEv
class A {
public:
static __global__ void kernel(){}
};
// CHECK: define{{.*}} void @_Z10non_kernelv
// CHECK-SPIRV: define{{.*}} void @_Z10non_kernelv
__device__ void non_kernel(){}
// CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneli
// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z6kerneli
__global__ void kernel(int x) {
non_kernel();
}
// CHECK: define{{.*}} amdgpu_kernel void @_Z11EmptyKernelIvEvv
// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z11EmptyKernelIvEvv
template <typename T>
__global__ void EmptyKernel(void) {}
struct Dummy {
/// Type definition of the EmptyKernel kernel entry point
typedef void (*EmptyKernelPtr)();
EmptyKernelPtr Empty() { return EmptyKernel<void>; }
};
// CHECK: define{{.*}} amdgpu_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
template<class T>
__global__ void template_kernel(T x) {}
void launch(void *f);
int main() {
Dummy D;
launch((void*)A::kernel);
launch((void*)kernel);
launch((void*)template_kernel<A>);
launch((void*)D.Empty());
return 0;
}
// CHECK: attributes #[[ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"
|
