Skip to content

Commit 69a37c2

Browse files
zhaomaosukbenzie
zhaomaosu
and
kbenzie
authored
[DeviceSanitizer] Add e2e tests for detecting out-of-bounds errors on sycl::buffer (#13504)
UR Part: oneapi-src/unified-runtime#1533 --------- Co-authored-by: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
1 parent c74c3f2 commit 69a37c2

File tree

6 files changed

+198
-28
lines changed

6 files changed

+198
-28
lines changed

‎sycl/plugins/unified_runtime/CMakeLists.txt

Lines changed: 10 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -100,34 +100,22 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
100100
endfunction()
101101

102102
set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
103-
# commit c0c607c3a88933b4c5c20a0aca4539781c678411
104-
# Merge: e18c691d 0f2d1f42
105-
# Author: Omar Ahmed <omar.ahmed@codeplay.com>
106-
# Date: Wed May 29 15:02:39 2024 +0100
107-
# Merge pull request #1667 from nrspruit/fix_multi_device_event_cache
108-
# [UR] Fix Multi Device Event Cache for shared Root Device
109-
set(UNIFIED_RUNTIME_TAG c0c607c3a88933b4c5c20a0aca4539781c678411)
103+
# commit 9f783837089c970a22cda08f768aa3dbed38f0d3
104+
# Merge: c015f892 b9442104
105+
# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
106+
# Date: Fri May 31 10:20:23 2024 +0100
107+
# Merge pull request #1533 from AllanZyne/sanitizer-buffer
108+
# [DeviceSanitizer] Support detecting out-of-bounds errors on sycl::buffer
109+
set(UNIFIED_RUNTIME_TAG 9f783837089c970a22cda08f768aa3dbed38f0d3)
110110

111111
fetch_adapter_source(level_zero
112112
${UNIFIED_RUNTIME_REPO}
113-
# commit 1fa0d2302d70a9aa8b717ecb60db0a7d8a4e5d8f
114-
# Merge: e16d01c0 35b7de5c
115-
# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
116-
# Date: Thu May 30 10:14:37 2024 +0100
117-
# Merge pull request #1654 from Bensuo/ewan/debug_log_update
118-
# Debug logging in Level Zero command-buffer update
119-
1fa0d2302d70a9aa8b717ecb60db0a7d8a4e5d8f
113+
${UNIFIED_RUNTIME_TAG}
120114
)
121115

122116
fetch_adapter_source(opencl
123117
${UNIFIED_RUNTIME_REPO}
124-
# commit e16d01c02f1c1c18af19d3c882c08f791c2345d9
125-
# Merge: a3895dbf 483a6325
126-
# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
127-
# Date: Thu May 30 10:04:43 2024 +0100
128-
# Merge pull request #1603 from konradkusiak97/queueFillOPENCLctsFix
129-
# [OpenCL] Modify fill emulation to work for patterns which are not powers of 2
130-
e16d01c02f1c1c18af19d3c882c08f791c2345d9
118+
${UNIFIED_RUNTIME_TAG}
131119
)
132120

133121
fetch_adapter_source(cuda
@@ -137,13 +125,7 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
137125

138126
fetch_adapter_source(hip
139127
${UNIFIED_RUNTIME_REPO}
140-
# commit a3895dbf46a98c8a5cd58704efe85e12fa1050cd
141-
# Merge: 30c5c553 67338dbe
142-
# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
143-
# Date: Thu May 30 10:04:14 2024 +0100
144-
# Merge pull request #1604 from kbenzie/benie/hip-kernel-spec-constants
145-
# [HIP] Implement kernel set spec constant query
146-
a3895dbf46a98c8a5cd58704efe85e12fa1050cd
128+
${UNIFIED_RUNTIME_TAG}
147129
)
148130

149131
fetch_adapter_source(native_cpu

‎sycl/test-e2e/AddressSanitizer/out-of-bounds/buffer/buffer.cpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
// REQUIRES: linux, cpu
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
3+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
5+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_asan_flags -O2 -g -o %t.out
7+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
8+
9+
#include <sycl/detail/core.hpp>
10+
11+
static const int N = 16;
12+
13+
int main() {
14+
sycl::queue q;
15+
16+
std::vector<int> v(N);
17+
18+
{
19+
// We intentionally test sycl::buffer uses host ptr and trigger data write
20+
// back here because in unified runtime we intercept sycl::buffer with usm,
21+
// we need to cover that pattern here.
22+
sycl::buffer<int, 1> buf(v.data(), v.size());
23+
q.submit([&](sycl::handler &h) {
24+
auto A = buf.get_access<sycl::access::mode::read_write>(h);
25+
h.parallel_for<class Test>(
26+
sycl::nd_range<1>(N + 1, 1),
27+
[=](sycl::nd_item<1> item) { A[item.get_global_id()] *= 2; });
28+
}).wait();
29+
// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
30+
// CHECK: {{READ of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(16, 0, 0\)}}
31+
// CHECK: {{#0 .* .*buffer.cpp:}}[[@LINE-4]]
32+
}
33+
34+
return 0;
35+
}

‎sycl/test-e2e/AddressSanitizer/out-of-bounds/buffer/buffer_2d.cpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
// REQUIRES: linux, cpu
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
3+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
5+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
6+
7+
#include <sycl/detail/core.hpp>
8+
9+
int main() {
10+
constexpr size_t size_x = 5;
11+
constexpr size_t size_y = 6;
12+
13+
std::vector<int> v(size_x * size_y);
14+
15+
// We intentionally test sycl::buffer uses host ptr here because in unified
16+
// runtime we intercept sycl::buffer with usm, we need to cover that pattern
17+
// here.
18+
sycl::buffer<int, 2> buf(v.data(), sycl::range<2>(size_x, size_y));
19+
20+
sycl::queue q;
21+
22+
q.submit([&](sycl::handler &cgh) {
23+
auto accessor = buf.get_access<sycl::access::mode::read_write>(cgh);
24+
cgh.parallel_for<class Test>(
25+
sycl::nd_range<2>({size_x, size_y + 1}, {1, 1}),
26+
[=](sycl::nd_item<2> item) {
27+
accessor[item.get_global_id()] = item.get_global_linear_id();
28+
});
29+
}).wait();
30+
// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
31+
// CHECK: {{WRITE of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(6, 4, 0\)}}
32+
// CHECK: {{#0 .* .*buffer_2d.cpp:}}[[@LINE-5]]
33+
34+
return 0;
35+
}

‎sycl/test-e2e/AddressSanitizer/out-of-bounds/buffer/buffer_3d.cpp

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
// REQUIRES: linux, cpu
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
3+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
5+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
6+
7+
#include <sycl/detail/core.hpp>
8+
9+
int main() {
10+
constexpr size_t size_x = 5;
11+
constexpr size_t size_y = 6;
12+
constexpr size_t size_z = 7;
13+
14+
std::vector<int> v(size_x * size_y * size_z);
15+
16+
// We intentionally test sycl::buffer uses host ptr here because in unified
17+
// runtime we intercept sycl::buffer with usm, we need to cover that pattern
18+
// here.
19+
sycl::buffer<int, 3> buf(v.data(), sycl::range<3>(size_x, size_y, size_z));
20+
21+
sycl::queue q;
22+
23+
q.submit([&](sycl::handler &cgh) {
24+
auto accessor = buf.get_access<sycl::access::mode::read_write>(cgh);
25+
26+
cgh.parallel_for<class Test>(
27+
sycl::nd_range<3>({size_x, size_y, size_z + 1}, {1, 1, 1}),
28+
[=](sycl::nd_item<3> item) {
29+
accessor[item.get_global_id()] = item.get_global_linear_id();
30+
});
31+
}).wait();
32+
// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
33+
// CHECK: {{WRITE of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(7, 5, 4\)}}
34+
// CHECK: {{#0 .* .*buffer_3d.cpp:}}[[@LINE-5]]
35+
36+
return 0;
37+
}

‎sycl/test-e2e/AddressSanitizer/out-of-bounds/buffer/buffer_copy_fill.cpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// REQUIRES: linux, cpu
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
3+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
5+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_asan_flags -O2 -g -o %t.out
7+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
8+
9+
#include <sycl/detail/core.hpp>
10+
11+
#include <numeric>
12+
13+
static const int N = 16;
14+
15+
int main() {
16+
sycl::queue q;
17+
18+
std::vector<int> v(N);
19+
std::iota(v.begin(), v.end(), 0);
20+
21+
{
22+
sycl::buffer<int, 1> buf(v.size());
23+
24+
q.submit([&](sycl::handler &h) {
25+
auto A = buf.get_access<sycl::access::mode::write>(h);
26+
h.copy(&v[0], A);
27+
}).wait();
28+
29+
q.submit([&](sycl::handler &h) {
30+
auto A = buf.get_access<sycl::access::mode::write>(h);
31+
h.fill(A, 1);
32+
}).wait();
33+
34+
q.submit([&](sycl::handler &h) {
35+
auto A = buf.get_access<sycl::access::mode::read_write>(h);
36+
h.parallel_for<class Test>(
37+
sycl::nd_range<1>(N + 1, 1),
38+
[=](sycl::nd_item<1> item) { A[item.get_global_id()] *= 2; });
39+
}).wait();
40+
// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
41+
// CHECK: {{READ of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(16, 0, 0\)}}
42+
// CHECK: {{#0 .* .*buffer_copy_fill.cpp:}}[[@LINE-4]]
43+
}
44+
45+
return 0;
46+
}

‎sycl/test-e2e/AddressSanitizer/out-of-bounds/buffer/subbuffer.cpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
// REQUIRES: linux, cpu
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
3+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
5+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_asan_flags -O2 -g -o %t.out
7+
// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
8+
9+
#include <sycl/detail/core.hpp>
10+
11+
int main() {
12+
constexpr size_t size_x = 16;
13+
14+
std::vector<int> v(size_x);
15+
for (size_t i = 0; i < size_x; i++)
16+
v[i] = i;
17+
18+
{
19+
sycl::queue q;
20+
sycl::buffer<int> buf(v.data(), v.size());
21+
sycl::buffer<int> sub_buf(buf, {size_x / 2}, {size_x / 2});
22+
23+
q.submit([&](sycl::handler &cgh) {
24+
auto accessor = sub_buf.get_access<sycl::access::mode::read_write>(cgh);
25+
cgh.parallel_for<class Test>(
26+
sycl::nd_range<1>(size_x / 2 + 1, 1),
27+
[=](sycl::nd_item<1> item) { accessor[item.get_global_id()] *= 2; });
28+
}).wait();
29+
// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
30+
// CHECK: {{READ of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(8, 0, 0\)}}
31+
// CHECK: {{#0 .* .*subbuffer.cpp:}}[[@LINE-4]]
32+
}
33+
34+
return 0;
35+
}

0 commit comments

Comments
 (0)