enhance kernel-fitler.cpp test

Author: AllanZyne

sycl/test-e2e/AddressSanitizer/common/ignorelist.txt

@@ -1 +1 @@
-fun:*MyKernel*
+fun:*NoSanitized*

sycl/test-e2e/AddressSanitizer/common/kernel-filter.cpp

@@ -1,27 +1,51 @@
-// REQUIRES: linux
-// RUN: %{build} %device_asan_flags -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t
-// RUN: %{run} %t 2>&1 | FileCheck %s
-// RUN: %{build} %device_asan_flags %if cpu %{ -fsycl-targets=spir64_x86_64 %} %if gpu %{ -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %} -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t2
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_asan_flags -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1
+// RUN: %{run} %t1 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_aot_flags -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t2
 // RUN: %{run} %t2 2>&1 | FileCheck %s
 
 #include <sycl/detail/core.hpp>
 #include <sycl/usm.hpp>
 
 int main() {
+  constexpr std::size_t N = 8;
+  constexpr std::size_t group_size = 4;
+
   sycl::queue Q;
-  constexpr std::size_t N = 16;
+
   auto *array = sycl::malloc_device<int>(N, Q);
 
+  std::vector<int> v(N);
+  sycl::buffer<int, 1> buf(v.data(), v.size());
+
   Q.submit([&](sycl::handler &h) {
-    h.parallel_for<class MyKernel>(
-        sycl::nd_range<1>(N + 1, 1),
-        [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
+    auto buf_acc = buf.get_access<sycl::access::mode::read_write>(h);
+    auto loc_acc = sycl::local_accessor<int>(group_size, h);
+    h.parallel_for<class NoSanitized>(
+        sycl::nd_range<1>(N, group_size), [=](sycl::nd_item<1> item) {
+          auto gid = item.get_global_id(0);
+          auto lid = item.get_local_id(0);
+          array[gid] = buf_acc[gid] + loc_acc[lid];
+        });
+  });
+  Q.wait();
+  // CHECK-NOT: ERROR: DeviceSanitizer: out-of-bounds-access
+
+  Q.submit([&](sycl::handler &h) {
+    auto buf_acc = buf.get_access<sycl::access::mode::read_write>(h);
+    auto loc_acc = sycl::local_accessor<int>(group_size, h);
+    h.parallel_for<class Sanitized>(sycl::nd_range<1>(N, group_size),
+                                    [=](sycl::nd_item<1> item) {
+                                      auto gid = item.get_global_id(0);
+                                      auto lid = item.get_local_id(0);
+                                      array[gid] = buf_acc[gid] + loc_acc[lid];
+                                    });
   });
   Q.wait();
 
   sycl::free(array, Q);
   std::cout << "PASS" << std::endl;
+  // CHECK: PASS
+
   return 0;
 }
-
-// CHECK: PASS

sycl/test-e2e/AddressSanitizer/lit.local.cfg

@@ -1,14 +1,12 @@
-# This test assumes it can only run in CPU/PVC/DG2 devices, which support usm_device_allocations aspect
-
 config.substitutions.append(
     ("%device_asan_flags", "-Xarch_device -fsanitize=address")
 )
-
+config.substitutions.append(
+    ("%device_asan_aot_flags", "-Xarch_device -fsanitize=address %if cpu %{ -fsycl-targets=spir64_x86_64 %} %if gpu %{ -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %}")
+)
 config.substitutions.append(
     ("%force_device_asan_rt", "env UR_ENABLE_LAYERS=UR_LAYER_ASAN")
 )
 
-config.unsupported_features += ['cuda', 'hip']
-
 # https://github.com/intel/llvm/issues/15953
 config.unsupported_features += ['gpu-intel-gen12']