Skip to content

Commit 459fa6f

Browse files
authored
[SYCL][NFC] Add missing LIT test (#5304)
After community pulldown of llvm/llvm-project@627868263cd4 the LIT test associated to the patch didn't get to be added (or might have been deleted inadvertently. This patch is to add back the LIT test.
1 parent fcd5a98 commit 459fa6f

File tree

1 file changed

+60
-0
lines changed

1 file changed

+60
-0
lines changed
Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
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>
17+
KERNEL void parallel_for(const KernelType &KernelFunc) {
18+
KernelFunc();
19+
}
20+
21+
void my_kernel(int my_param) {
22+
int my_local = 0;
23+
my_local = my_param;
24+
}
25+
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: call void @llvm.dbg.declare(
38+
// CHECK-SAME: metadata i32* %my_param.addr,
39+
// CHECK-SAME: metadata [[MY_PARAM:![0-9]+]],
40+
// CHECK-SAME: metadata !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
41+
// CHECK-SAME: )
42+
// CHECK: call void @llvm.dbg.declare(
43+
// CHECK-SAME: metadata i32* %my_local,
44+
// CHECK-SAME: metadata [[MY_LOCAL:![0-9]+]],
45+
// CHECK-SAME: metadata !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: )

0 commit comments

Comments
 (0)