[DeviceSanitizer] Add e2e tests for detecting out-of-bounds errors on sycl::buffer

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -100,34 +100,22 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit c0c607c3a88933b4c5c20a0aca4539781c678411
-  # Merge: e18c691d 0f2d1f42
-  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
-  # Date:   Wed May 29 15:02:39 2024 +0100
-  #     Merge pull request #1667 from nrspruit/fix_multi_device_event_cache
-  #     [UR] Fix Multi Device Event Cache for shared Root Device
-  set(UNIFIED_RUNTIME_TAG c0c607c3a88933b4c5c20a0aca4539781c678411)
+  # commit 9f783837089c970a22cda08f768aa3dbed38f0d3
+  # Merge: c015f892 b9442104
+  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
+  # Date:   Fri May 31 10:20:23 2024 +0100
+  #     Merge pull request #1533 from AllanZyne/sanitizer-buffer
+  #     [DeviceSanitizer] Support detecting out-of-bounds errors on sycl::buffer
+  set(UNIFIED_RUNTIME_TAG 9f783837089c970a22cda08f768aa3dbed38f0d3)
 
   fetch_adapter_source(level_zero
     ${UNIFIED_RUNTIME_REPO}
-    # commit 1fa0d2302d70a9aa8b717ecb60db0a7d8a4e5d8f
-    # Merge: e16d01c0 35b7de5c
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Thu May 30 10:14:37 2024 +0100
-    #     Merge pull request #1654 from Bensuo/ewan/debug_log_update
-    #     Debug logging in Level Zero command-buffer update
-    1fa0d2302d70a9aa8b717ecb60db0a7d8a4e5d8f
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(opencl
     ${UNIFIED_RUNTIME_REPO}
-    # commit e16d01c02f1c1c18af19d3c882c08f791c2345d9
-    # Merge: a3895dbf 483a6325
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Thu May 30 10:04:43 2024 +0100
-    #     Merge pull request #1603 from konradkusiak97/queueFillOPENCLctsFix
-    #     [OpenCL] Modify fill emulation to work for patterns which are not powers of 2
-    e16d01c02f1c1c18af19d3c882c08f791c2345d9
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(cuda
@@ -137,13 +125,7 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
 
   fetch_adapter_source(hip
     ${UNIFIED_RUNTIME_REPO}
-    # commit a3895dbf46a98c8a5cd58704efe85e12fa1050cd
-    # Merge: 30c5c553 67338dbe
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Thu May 30 10:04:14 2024 +0100
-    #     Merge pull request #1604 from kbenzie/benie/hip-kernel-spec-constants
-    #     [HIP] Implement kernel set spec constant query
-    a3895dbf46a98c8a5cd58704efe85e12fa1050cd
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(native_cpu

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

@@ -0,0 +1,35 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+static const int N = 16;
+
+int main() {
+  sycl::queue q;
+
+  std::vector<int> v(N);
+
+  {
+    // We intentionally test sycl::buffer uses host ptr and trigger data write
+    // back here because in unified runtime we intercept sycl::buffer with usm,
+    // we need to cover that pattern here.
+    sycl::buffer<int, 1> buf(v.data(), v.size());
+    q.submit([&](sycl::handler &h) {
+       auto A = buf.get_access<sycl::access::mode::read_write>(h);
+       h.parallel_for<class Test>(
+           sycl::nd_range<1>(N + 1, 1),
+           [=](sycl::nd_item<1> item) { A[item.get_global_id()] *= 2; });
+     }).wait();
+    // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
+    // CHECK: {{READ of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(16, 0, 0\)}}
+    // CHECK: {{#0 .* .*buffer.cpp:}}[[@LINE-4]]
+  }
+
+  return 0;
+}

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

@@ -0,0 +1,35 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+int main() {
+  constexpr size_t size_x = 5;
+  constexpr size_t size_y = 6;
+
+  std::vector<int> v(size_x * size_y);
+
+  // We intentionally test sycl::buffer uses host ptr here because in unified
+  // runtime we intercept sycl::buffer with usm, we need to cover that pattern
+  // here.
+  sycl::buffer<int, 2> buf(v.data(), sycl::range<2>(size_x, size_y));
+
+  sycl::queue q;
+
+  q.submit([&](sycl::handler &cgh) {
+     auto accessor = buf.get_access<sycl::access::mode::read_write>(cgh);
+     cgh.parallel_for<class Test>(
+         sycl::nd_range<2>({size_x, size_y + 1}, {1, 1}),
+         [=](sycl::nd_item<2> item) {
+           accessor[item.get_global_id()] = item.get_global_linear_id();
+         });
+   }).wait();
+  // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
+  // CHECK: {{WRITE of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(6, 4, 0\)}}
+  // CHECK: {{#0 .* .*buffer_2d.cpp:}}[[@LINE-5]]
+
+  return 0;
+}

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

@@ -0,0 +1,37 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+int main() {
+  constexpr size_t size_x = 5;
+  constexpr size_t size_y = 6;
+  constexpr size_t size_z = 7;
+
+  std::vector<int> v(size_x * size_y * size_z);
+
+  // We intentionally test sycl::buffer uses host ptr here because in unified
+  // runtime we intercept sycl::buffer with usm, we need to cover that pattern
+  // here.
+  sycl::buffer<int, 3> buf(v.data(), sycl::range<3>(size_x, size_y, size_z));
+
+  sycl::queue q;
+
+  q.submit([&](sycl::handler &cgh) {
+     auto accessor = buf.get_access<sycl::access::mode::read_write>(cgh);
+
+     cgh.parallel_for<class Test>(
+         sycl::nd_range<3>({size_x, size_y, size_z + 1}, {1, 1, 1}),
+         [=](sycl::nd_item<3> item) {
+           accessor[item.get_global_id()] = item.get_global_linear_id();
+         });
+   }).wait();
+  // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
+  // CHECK: {{WRITE of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(7, 5, 4\)}}
+  // CHECK: {{#0 .* .*buffer_3d.cpp:}}[[@LINE-5]]
+
+  return 0;
+}

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

@@ -0,0 +1,46 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+#include <numeric>
+
+static const int N = 16;
+
+int main() {
+  sycl::queue q;
+
+  std::vector<int> v(N);
+  std::iota(v.begin(), v.end(), 0);
+
+  {
+    sycl::buffer<int, 1> buf(v.size());
+
+    q.submit([&](sycl::handler &h) {
+       auto A = buf.get_access<sycl::access::mode::write>(h);
+       h.copy(&v[0], A);
+     }).wait();
+
+    q.submit([&](sycl::handler &h) {
+       auto A = buf.get_access<sycl::access::mode::write>(h);
+       h.fill(A, 1);
+     }).wait();
+
+    q.submit([&](sycl::handler &h) {
+       auto A = buf.get_access<sycl::access::mode::read_write>(h);
+       h.parallel_for<class Test>(
+           sycl::nd_range<1>(N + 1, 1),
+           [=](sycl::nd_item<1> item) { A[item.get_global_id()] *= 2; });
+     }).wait();
+    // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
+    // CHECK: {{READ of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(16, 0, 0\)}}
+    // CHECK: {{#0 .* .*buffer_copy_fill.cpp:}}[[@LINE-4]]
+  }
+
+  return 0;
+}

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

@@ -0,0 +1,35 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+int main() {
+  constexpr size_t size_x = 16;
+
+  std::vector<int> v(size_x);
+  for (size_t i = 0; i < size_x; i++)
+    v[i] = i;
+
+  {
+    sycl::queue q;
+    sycl::buffer<int> buf(v.data(), v.size());
+    sycl::buffer<int> sub_buf(buf, {size_x / 2}, {size_x / 2});
+
+    q.submit([&](sycl::handler &cgh) {
+       auto accessor = sub_buf.get_access<sycl::access::mode::read_write>(cgh);
+       cgh.parallel_for<class Test>(
+           sycl::nd_range<1>(size_x / 2 + 1, 1),
+           [=](sycl::nd_item<1> item) { accessor[item.get_global_id()] *= 2; });
+     }).wait();
+    // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Memory Buffer
+    // CHECK: {{READ of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(8, 0, 0\)}}
+    // CHECK: {{#0 .* .*subbuffer.cpp:}}[[@LINE-4]]
+  }
+
+  return 0;
+}