Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 3bd46fd

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL] Add optimized/non-optimized CHECKs for SYCL/XPTI/kernel/* (#1000)
If the test-suite was configured to pass -O0 the tests were changing their behavior because no dead kernel arguments optimization was performed, resulting in failing CHECKs. Be explicit in the "RUN:" commands and expected "CHECK:" lines.
1 parent 1de20a0 commit 3bd46fd

File tree

2 files changed

+30
-21
lines changed

2 files changed

+30
-21
lines changed

‎SYCL/XPTI/kernel/basic.cpp

Lines changed: 22 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,9 @@
11
// REQUIRES: xptifw, opencl
22
// RUN: %clangxx %s -DXPTI_COLLECTOR -DXPTI_CALLBACK_API_EXPORTS %xptifw_lib %shared_lib %fPIC %cxx_std_optionc++17 -o %t_collector.dll
3-
// RUN: %clangxx -fsycl %s -o %t.out
4-
// RUN: env XPTI_TRACE_ENABLE=1 XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher XPTI_SUBSCRIBERS=%t_collector.dll %BE_RUN_PLACEHOLDER %t.out | FileCheck %s 2>&1
3+
// RUN: %clangxx -fsycl -O2 %s -o %t.opt.out
4+
// RUN: env XPTI_TRACE_ENABLE=1 XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher XPTI_SUBSCRIBERS=%t_collector.dll %BE_RUN_PLACEHOLDER %t.opt.out | FileCheck %s --check-prefixes=CHECK,CHECK-OPT
5+
// RUN: %clangxx -fsycl -fno-sycl-dead-args-optimization %s -o %t.noopt.out
6+
// RUN: env XPTI_TRACE_ENABLE=1 XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher XPTI_SUBSCRIBERS=%t_collector.dll %BE_RUN_PLACEHOLDER %t.noopt.out | FileCheck %s --check-prefixes=CHECK,CHECK-NOOPT
57

68
#ifdef XPTI_COLLECTOR
79

@@ -56,19 +58,20 @@ int main() {
5658
auto A1 = Buf.get_access<mode::read_write>(cgh);
5759
// CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID2:.*]]|2016|1026|{{.*}}.cpp:[[# @LINE + 1]]:65
5860
sycl::accessor<int, 1, mode::read_write, target::local> A2(Range, cgh);
59-
// CHECK:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 6
61+
// CHECK-OPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 6
62+
// CHECK-NOOPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 7 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 12
6063
cgh.parallel_for<class FillBuffer>(
6164
Range, [=](sycl::id<1> WIid, sycl::kernel_handler kh) {
62-
// CHECK: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
65+
// CHECK-OPT: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
6366
int h = Val;
64-
// CHECK: arg1 : {1, {{.*}}0, 20, 1}
67+
// CHECK-OPT: arg1 : {1, {{.*}}0, 20, 1}
6568
A2[WIid[0]] = h;
66-
// CHECK: arg2 : {0, [[ACCID1]], 4062, 2}
67-
// CHECK: arg3 : {1, [[ACCID1]], 8, 3}
69+
// CHECK-OPT: arg2 : {0, [[ACCID1]], 4062, 2}
70+
// CHECK-OPT: arg3 : {1, [[ACCID1]], 8, 3}
6871
A1[WIid[0]] = A2[WIid[0]];
69-
// CHECK: arg4 : {3, {{.*}}, 8, 4}
72+
// CHECK-OPT: arg4 : {3, {{.*}}, 8, 4}
7073
PtrDevice[WIid[0]] = WIid[0];
71-
// CHECK: arg5 : {3, {{.*}}, 8, 5}
74+
// CHECK-OPT: arg5 : {3, {{.*}}, 8, 5}
7275
PtrShared[WIid[0]] = PtrDevice[WIid[0]];
7376
});
7477
})
@@ -81,22 +84,24 @@ int main() {
8184
// CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID4:.+]]|2014|1026|{{.*}}.cpp:[[# @LINE + 1]]:16
8285
auto Acc = Buf.get_access<mode::read_write>(cgh);
8386
Functor1 F(Val, Acc);
84-
// CHECK: Node create|{{.*}}Functor1|{{.*}}.cpp:[[# @LINE - 4 ]]:3|{1, 1, 1}, {0, 0, 0}, {0, 0, 0}, 3
87+
// CHECK-OPT: Node create|{{.*}}Functor1|{{.*}}.cpp:[[# @LINE - 4 ]]:3|{1, 1, 1}, {0, 0, 0}, {0, 0, 0}, 3
88+
// CHECK-NOOPT: Node create|{{.*}}Functor1|{{.*}}.cpp:[[# @LINE - 5 ]]:3|{1, 1, 1}, {0, 0, 0}, {0, 0, 0}, 5
8589
cgh.single_task(F);
86-
// CHECK: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
87-
// CHECK: arg1 : {0, [[ACCID4]], 4062, 1}
88-
// CHECK: arg2 : {1, [[ACCID4]], 8, 2}
90+
// CHECK-OPT: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
91+
// CHECK-OPT: arg1 : {0, [[ACCID4]], 4062, 1}
92+
// CHECK-OPT: arg2 : {1, [[ACCID4]], 8, 2}
8993
});
9094

9195
Queue.submit([&](cl::sycl::handler &cgh) {
9296
// CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.+]]|2014|1026|{{.*}}.cpp:[[# @LINE + 1]]:16
9397
auto Acc = Buf.get_access<mode::read_write>(cgh);
9498
Functor2 F(Val, Acc);
95-
// CHECK: Node create|{{.*}}Functor2|{{.*}}.cpp:[[# @LINE - 4 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 3
99+
// CHECK-OPT: Node create|{{.*}}Functor2|{{.*}}.cpp:[[# @LINE - 4 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 3
100+
// CHECK-NOOPT: Node create|{{.*}}Functor2|{{.*}}.cpp:[[# @LINE - 5 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 5
96101
cgh.parallel_for(Range, F);
97-
// CHECK: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
98-
// CHECK: arg1 : {0, [[ACCID5]], 4062, 1}
99-
// CHECK: arg2 : {1, [[ACCID5]], 8, 2}
102+
// CHECK-OPT: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
103+
// CHECK-OPT: arg1 : {0, [[ACCID5]], 4062, 1}
104+
// CHECK-OPT: arg2 : {1, [[ACCID5]], 8, 2}
100105
});
101106

102107
return 0;

‎SYCL/XPTI/kernel/content.cpp

Lines changed: 8 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,9 @@
11
// REQUIRES: xptifw, opencl
22
// RUN: %clangxx %s -DXPTI_COLLECTOR -DXPTI_CALLBACK_API_EXPORTS %xptifw_lib %shared_lib %fPIC %cxx_std_optionc++17 -o %t_collector.dll
3-
// RUN: %clangxx -fsycl %s -o %t.out
4-
// RUN: env XPTI_TRACE_ENABLE=1 XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher XPTI_SUBSCRIBERS=%t_collector.dll %BE_RUN_PLACEHOLDER %t.out | FileCheck %s 2>&1
3+
// RUN: %clangxx -fsycl -O2 %s -o %t.opt.out
4+
// RUN: env XPTI_TRACE_ENABLE=1 XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher XPTI_SUBSCRIBERS=%t_collector.dll %BE_RUN_PLACEHOLDER %t.opt.out | FileCheck %s --check-prefix=CHECK-OPT
5+
// RUN: %clangxx -fsycl -fno-sycl-dead-args-optimization %s -o %t.noopt.out
6+
// RUN: env XPTI_TRACE_ENABLE=1 XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher XPTI_SUBSCRIBERS=%t_collector.dll %BE_RUN_PLACEHOLDER %t.noopt.out | FileCheck %s --check-prefix=CHECK-NOOPT
57

68
#ifdef XPTI_COLLECTOR
79

@@ -27,7 +29,8 @@ int main() {
2729

2830
auto sumR = reduction(sumBuf, cgh, plus<>());
2931
// Reduction kernel is used
30-
// CHECK:Node create|{{.*}}reduction{{.*}}test1{{.*}}|{{.*}}.cpp:[[# @LINE - 5 ]]:3|{1024, 1, 1}, {{{.*}}, 1, 1}, {0, 0, 0}, 5
32+
// CHECK-OPT:Node create|{{.*}}reduction{{.*}}test1{{.*}}|{{.*}}.cpp:[[# @LINE - 5 ]]:3|{1024, 1, 1}, {{{.*}}, 1, 1}, {0, 0, 0}, 5
33+
// CHECK-NOOPT:Node create|{{.*}}reduction{{.*}}test1{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{1024, 1, 1}, {{{.*}}, 1, 1}, {0, 0, 0}, 13
3134
cgh.parallel_for<class test1>(
3235
range<1>{1024}, sumR,
3336
[=](id<1> idx, auto &sum) { sum += inputValues[idx]; });
@@ -42,7 +45,8 @@ int main() {
4245
myQueue.submit([&](handler &cgh) {
4346
auto in = in_buf.template get_access<access::mode::read>(cgh);
4447
auto out = out_buf.template get_access<access::mode::read_write>(cgh);
45-
// CHECK:Node create|{{.*}}test2{{.*}}|{{.*}}.cpp:[[# @LINE - 3 ]]:5|{128, 4, 2}, {32, 2, 1}, {16, 1, 0}, 2
48+
// CHECK-OPT:Node create|{{.*}}test2{{.*}}|{{.*}}.cpp:[[# @LINE - 3 ]]:5|{128, 4, 2}, {32, 2, 1}, {16, 1, 0}, 2
49+
// CHECK-NOOPT:Node create|{{.*}}test2{{.*}}|{{.*}}.cpp:[[# @LINE - 4 ]]:5|{128, 4, 2}, {32, 2, 1}, {16, 1, 0}, 8
4650
cgh.parallel_for<class test2>(
4751
nd_range<3>({128, 4, 2}, {32, 2, 1}, {16, 1, 0}), [=](nd_item<3> it) {
4852
auto sg = it.get_sub_group();

0 commit comments

Comments
 (0)