[DeviceASAN] Fix invalid test-e2e "kernel-filter.cpp" (#17328)

The old "kernel-filter.cpp" is invalid to test kernel filter feature.

Author: AllanZyne

sycl/test-e2e/AddressSanitizer/aot/gpu.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: linux, ocloc, gpu && level_zero
+// REQUIRES: linux, gpu && level_zero
 // REQUIRES: (arch-intel_gpu_pvc || gpu-intel-dg2)
 
 // RUN: %{run-aux} %{build} %device_asan_aot_flags -O0 -g %S/Inputs/host-usm-oob.cpp -o %t.out

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

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

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

@@ -1,4 +1,6 @@
 // REQUIRES: linux, opencl-aot, cpu
 
-// RUN: %{run-aux} %{build} %device_asan_aot_flags %S/kernel-filter.cpp -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1
-// RUN: %{run} %t1 2>&1 | FileCheck %S/kernel-filter.cpp
+// RUN: %{run-aux} %{build} %device_asan_aot_flags %S/kernel-filter.cpp -g -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1
+// RUN: %{run} %t1 2>&1 | FileCheck %S/kernel-filter.cpp --check-prefixes CHECK-IGNORE
+// RUN: %{run-aux} %{build} %device_asan_aot_flags %S/kernel-filter.cpp -g -O2 -o %t2
+// RUN: %{run} not %t2 2>&1 | FileCheck %S/kernel-filter.cpp

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

@@ -1,5 +1,7 @@
-// REQUIRES: linux, ocloc, gpu && level_zero
+// REQUIRES: linux, gpu && level_zero
 // REQUIRES: (arch-intel_gpu_pvc || gpu-intel-dg2)
 
-// RUN: %{run-aux} %{build} %device_asan_aot_flags %S/kernel-filter.cpp -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1
-// RUN: %{run} %t1 2>&1 | FileCheck %S/kernel-filter.cpp
+// RUN: %{run-aux} %{build} %device_asan_aot_flags %S/kernel-filter.cpp -g -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1
+// RUN: %{run} %t1 2>&1 | FileCheck %S/kernel-filter.cpp --check-prefixes CHECK-IGNORE
+// RUN: %{run-aux} %{build} %device_asan_aot_flags %S/kernel-filter.cpp -g -O2 -o %t2
+// RUN: %{run} not %t2 2>&1 | FileCheck %S/kernel-filter.cpp

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

@@ -1,6 +1,8 @@
 // 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_flags -g -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1
+// RUN: %{run} %t1 2>&1 | FileCheck %s --check-prefixes CHECK-IGNORE
+// RUN: %{build} %device_asan_flags -g -O2 -o %t2
+// RUN: %{run} not %t2 2>&1 | FileCheck %s
 
 #include <sycl/detail/core.hpp>
 #include <sycl/usm.hpp>
@@ -19,31 +21,17 @@ int main() {
   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 NoSanitized>(
+    h.parallel_for<class MyKernel>(
         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];
+          array[gid + 1] = 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();
+  // CHECK-IGNORE-NOT: ERROR: DeviceSanitizer: out-of-bounds-access
+  // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access
 
   sycl::free(array, Q);
-  std::cout << "PASS" << std::endl;
-  // CHECK: PASS
-
   return 0;
 }