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

[SYCL] Add optimized/non-optimized CHECKs for SYCL/XPTI/kernel/* #1000

Merged
merged 2 commits into from
Apr 21, 2022
Merged

[SYCL] Add optimized/non-optimized CHECKs for SYCL/XPTI/kernel/* #1000

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c23ee2e
[SYCL] Add optimized/non-optimized CHECKs for SYCL/XPTI/kernel/*
aelovikov-intel Apr 19, 2022
b843cf1
Use -fno-sycl-dead-args-optimization instead of -O0
aelovikov-intel Apr 20, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
39 changes: 22 additions & 17 deletions SYCL/XPTI/kernel/basic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,7 +1,9 @@
// REQUIRES: xptifw, opencl
// RUN: %clangxx %s -DXPTI_COLLECTOR -DXPTI_CALLBACK_API_EXPORTS %xptifw_lib %shared_lib %fPIC %cxx_std_optionc++17 -o %t_collector.dll
// RUN: %clangxx -fsycl %s -o %t.out
// 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
// RUN: %clangxx -fsycl -O2 %s -o %t.opt.out
// 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
// RUN: %clangxx -fsycl -fno-sycl-dead-args-optimization %s -o %t.noopt.out
// 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

#ifdef XPTI_COLLECTOR

Expand Down Expand Up @@ -56,19 +58,20 @@ int main() {
auto A1 = Buf.get_access<mode::read_write>(cgh);
// CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID2:.*]]|2016|1026|{{.*}}.cpp:[[# @LINE + 1]]:65
sycl::accessor<int, 1, mode::read_write, target::local> A2(Range, cgh);
// CHECK:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 6
// CHECK-OPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 6
// CHECK-NOOPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 7 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 12
cgh.parallel_for<class FillBuffer>(
Range, [=](sycl::id<1> WIid, sycl::kernel_handler kh) {
// CHECK: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
// CHECK-OPT: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
int h = Val;
// CHECK: arg1 : {1, {{.*}}0, 20, 1}
// CHECK-OPT: arg1 : {1, {{.*}}0, 20, 1}
A2[WIid[0]] = h;
// CHECK: arg2 : {0, [[ACCID1]], 4062, 2}
// CHECK: arg3 : {1, [[ACCID1]], 8, 3}
// CHECK-OPT: arg2 : {0, [[ACCID1]], 4062, 2}
// CHECK-OPT: arg3 : {1, [[ACCID1]], 8, 3}
A1[WIid[0]] = A2[WIid[0]];
// CHECK: arg4 : {3, {{.*}}, 8, 4}
// CHECK-OPT: arg4 : {3, {{.*}}, 8, 4}
PtrDevice[WIid[0]] = WIid[0];
// CHECK: arg5 : {3, {{.*}}, 8, 5}
// CHECK-OPT: arg5 : {3, {{.*}}, 8, 5}
PtrShared[WIid[0]] = PtrDevice[WIid[0]];
});
})
Expand All @@ -81,22 +84,24 @@ int main() {
// CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID4:.+]]|2014|1026|{{.*}}.cpp:[[# @LINE + 1]]:16
auto Acc = Buf.get_access<mode::read_write>(cgh);
Functor1 F(Val, Acc);
// CHECK: Node create|{{.*}}Functor1|{{.*}}.cpp:[[# @LINE - 4 ]]:3|{1, 1, 1}, {0, 0, 0}, {0, 0, 0}, 3
// CHECK-OPT: Node create|{{.*}}Functor1|{{.*}}.cpp:[[# @LINE - 4 ]]:3|{1, 1, 1}, {0, 0, 0}, {0, 0, 0}, 3
// CHECK-NOOPT: Node create|{{.*}}Functor1|{{.*}}.cpp:[[# @LINE - 5 ]]:3|{1, 1, 1}, {0, 0, 0}, {0, 0, 0}, 5
cgh.single_task(F);
// CHECK: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
// CHECK: arg1 : {0, [[ACCID4]], 4062, 1}
// CHECK: arg2 : {1, [[ACCID4]], 8, 2}
// CHECK-OPT: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
// CHECK-OPT: arg1 : {0, [[ACCID4]], 4062, 1}
// CHECK-OPT: arg2 : {1, [[ACCID4]], 8, 2}
});

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

return 0;
Expand Down
12 changes: 8 additions & 4 deletions SYCL/XPTI/kernel/content.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,7 +1,9 @@
// REQUIRES: xptifw, opencl
// RUN: %clangxx %s -DXPTI_COLLECTOR -DXPTI_CALLBACK_API_EXPORTS %xptifw_lib %shared_lib %fPIC %cxx_std_optionc++17 -o %t_collector.dll
// RUN: %clangxx -fsycl %s -o %t.out
// 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
// RUN: %clangxx -fsycl -O2 %s -o %t.opt.out
// 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
// RUN: %clangxx -fsycl -fno-sycl-dead-args-optimization %s -o %t.noopt.out
// 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

#ifdef XPTI_COLLECTOR

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

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