1 // RUN: %clang_cc1 %s -o - -O0 -emit-llvm \
2 // RUN: -triple spir64-unknown-unknown \
3 // RUN: -aux-triple x86_64-unknown-linux-gnu \
4 // RUN: -fsycl-is-device \
5 // RUN: -finclude-default-header \
6 // RUN: -debug-info-kind=limited -gno-column-info \
7 // RUN: | FileCheck %s
8 //
9 // In spir functions, validate the llvm.dbg.declare intrinsics created for
10 // parameters and locals refer to the stack allocation in the alloca address
11 // space.
12 //
13
14 #define KERNEL __attribute__((sycl_kernel))
15
16 template <typename KernelName, typename KernelType>
parallel_for(const KernelType & KernelFunc)17 KERNEL void parallel_for(const KernelType &KernelFunc) {
18 KernelFunc();
19 }
20
my_kernel(int my_param)21 void my_kernel(int my_param) {
22 int my_local = 0;
23 my_local = my_param;
24 }
25
my_host()26 int my_host() {
27 parallel_for<class K>([=]() { my_kernel(42); });
28 return 0;
29 }
30
31 // CHECK: define {{.*}}spir_func void @_Z9my_kerneli(
32 // CHECK-SAME i32 %my_param
33 // CHECK-SAME: !dbg [[MY_KERNEL:![0-9]+]]
34 // CHECK-SAME: {
35 // CHECK: %my_param.addr = alloca i32, align 4
36 // CHECK: %my_local = alloca i32, align 4
37 // CHECK: #dbg_declare(
38 // CHECK-SAME: ptr %my_param.addr,
39 // CHECK-SAME: [[MY_PARAM:![0-9]+]],
40 // CHECK-SAME: !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
41 // CHECK-SAME: )
42 // CHECK: #dbg_declare(
43 // CHECK-SAME: ptr %my_local,
44 // CHECK-SAME: [[MY_LOCAL:![0-9]+]],
45 // CHECK-SAME: !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
46 // CHECK-SAME: )
47 // CHECK: }
48
49 // CHECK: [[MY_KERNEL]] = distinct !DISubprogram(
50 // CHECK-SAME: name: "my_kernel"
51 // CHECK-SAME: )
52 // CHECK: [[MY_PARAM]] = !DILocalVariable(
53 // CHECK-SAME: name: "my_param"
54 // CHECK-SAME: arg: 1
55 // CHECK-SAME: scope: [[MY_KERNEL]]
56 // CHECK-SAME: )
57 // CHECK: [[MY_LOCAL]] = !DILocalVariable(
58 // CHECK-SAME: name: "my_local"
59 // CHECK-SAME: scope: [[MY_KERNEL]]
60 // CHECK-SAME: )
61