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

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -94,14 +94,8 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 8cdd099ae3d1a34d3bcd7cbed7f5745c3dc8e112
-  # Merge: fc9bb61b c893a3c4
-  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Mon May 20 15:50:02 2024 +0100
-  #     Merge pull request #954 from jchlanda/jakub/rqwgs_hip
-  #     [HIP] Handle required wg size attribute in HIP
-  set(UNIFIED_RUNTIME_TAG 8cdd099ae3d1a34d3bcd7cbed7f5745c3dc8e112)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/AllanZyne/unified-runtime.git")
+  set(UNIFIED_RUNTIME_TAG sanitizer-buffer)
 
   fetch_adapter_source(level_zero
     ${UNIFIED_RUNTIME_REPO}

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

@@ -0,0 +1,34 @@
+// 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);
+  for (int i = 0; i < N; i++)
+    v[i] = i;
+
+  {
+    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,34 @@
+// 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);
+
+  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_id(0) * item.get_global_range(1) +
+               item.get_global_id(1);
+         });
+   }).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-7]]
+
+  return 0;
+}

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

@@ -0,0 +1,38 @@
+// 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);
+
+  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_id(0) * item.get_global_range(1) *
+                   item.get_global_range(2) +
+               item.get_global_id(1) * item.get_global_range(2) +
+               item.get_global_id(2);
+         });
+   }).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-9]]
+
+  return 0;
+}

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

@@ -0,0 +1,40 @@
+// 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);
+  for (int i = 0; i < N; i++)
+    v[i] = i;
+
+  {
+    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::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.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;
+}