[SYCL][DeviceSanitizer] Support GPU DG2 Device

libdevice/sanitizer_utils.cpp

@@ -83,6 +83,10 @@ static __SYCL_CONSTANT__ const char __generic_to[] =
 static __SYCL_CONSTANT__ const char __generic_to_fail[] =
     "[kernel] %p(4) - unknown address space\n";
 
+static __SYCL_CONSTANT__ const char __mem_launch_info[] =
+    "[kernel] launch_info: %p (local_shadow=%p~%p, numLocalArgs=%d, "
+    "localArgs=%p)\n";
+
 #define ASAN_REPORT_NONE 0
 #define ASAN_REPORT_START 1
 #define ASAN_REPORT_FINISH 2
@@ -111,56 +115,120 @@ __SYCL_PRIVATE__ void *ToPrivate(void *ptr) {
   return __spirv_GenericCastToPtrExplicit_ToPrivate(ptr, 7);
 }
 
+inline bool ConvertGenericPointer(uptr &addr, uint32_t &as) {
+  auto old = addr;
+  if ((addr = (uptr)ToPrivate((void *)old))) {
+    as = ADDRESS_SPACE_PRIVATE;
+  } else if ((addr = (uptr)ToLocal((void *)old))) {
+    as = ADDRESS_SPACE_LOCAL;
+  } else if ((addr = (uptr)ToGlobal((void *)old))) {
+    as = ADDRESS_SPACE_GLOBAL;
+  } else {
+    if (__AsanDebug)
+      __spirv_ocl_printf(__generic_to_fail, old);
+    return false;
+  }
+  if (__AsanDebug)
+    __spirv_ocl_printf(__generic_to, old, addr, as);
+  return true;
+}
+
 inline uptr MemToShadow_CPU(uptr addr) {
   return __AsanShadowMemoryGlobalStart + (addr >> ASAN_SHADOW_SCALE);
 }
 
 inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
-  uptr shadow_ptr = 0;
-  if (addr & (~0xffffffffffff)) {
-    shadow_ptr = (((addr & 0xffffffffffff) >> ASAN_SHADOW_SCALE) +
-                  __AsanShadowMemoryGlobalStart) |
-                 (~0xffffffffffff);
-  } else {
-    shadow_ptr = (addr >> ASAN_SHADOW_SCALE) + __AsanShadowMemoryGlobalStart;
+  if (as == ADDRESS_SPACE_GENERIC) {
+    if (!ConvertGenericPointer(addr, as)) {
+      return 0;
+    }
   }
 
-  if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
-    if (__asan_report_out_of_shadow_bounds()) {
-      __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr);
+  if (as == ADDRESS_SPACE_GLOBAL) {     // global
+    if (addr & 0xFFFF000000000000ULL) { // Device USM
+      return __AsanShadowMemoryGlobalStart + 0x80000000000ULL +
+             ((addr & 0x7FFFFFFFFFFFULL) >> ASAN_SHADOW_SCALE);
+    } else { // Host/Shared USM
+      return __AsanShadowMemoryGlobalStart + (addr >> ASAN_SHADOW_SCALE);
     }
-  }
+  } else if (as == ADDRESS_SPACE_LOCAL) { // local
+    // The size of SLM is 64KB on DG2
+    constexpr unsigned slm_size = 64 * 1024;
+    const auto wg_lid =
+        __spirv_BuiltInWorkgroupId.x * __spirv_BuiltInNumWorkgroups.y *
+            __spirv_BuiltInNumWorkgroups.z +
+        __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z +
+        __spirv_BuiltInWorkgroupId.z;
 
-  return shadow_ptr;
-}
+    auto launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
+    const auto shadow_offset = launch_info->LocalShadowOffset;
+    const auto shadow_offset_end = launch_info->LocalShadowOffsetEnd;
 
-static __SYCL_CONSTANT__ const char __mem_launch_info[] =
-    "[kernel] launch_info: %p (local_shadow=%p~%p, numLocalArgs=%d, "
-    "localArgs=%p)\n";
+    if (shadow_offset == 0) {
+      return 0;
+    }
 
-static __SYCL_CONSTANT__ const char __generic_to[] =
-    "[kernel] %p(4) - %p(%d)\n";
+    if (__AsanDebug)
+      __spirv_ocl_printf(__mem_launch_info, launch_info,
+                         launch_info->LocalShadowOffset,
+                         launch_info->LocalShadowOffsetEnd,
+                         launch_info->NumLocalArgs, launch_info->LocalArgs);
 
-static __SYCL_CONSTANT__ const char __generic_to_fail[] =
-    "[kernel] %p(4) - unknown address space\n";
+    auto shadow_ptr = shadow_offset +
+                      ((wg_lid * slm_size) >> ASAN_SHADOW_SCALE) +
+                      ((addr & (slm_size - 1)) >> ASAN_SHADOW_SCALE);
 
-inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
+    if (shadow_ptr > shadow_offset_end) {
+      if (__asan_report_out_of_shadow_bounds()) {
+        __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
+                           wg_lid, (uptr)shadow_offset);
+      }
+      return 0;
+    }
+    return shadow_ptr;
+  } else if (as == ADDRESS_SPACE_PRIVATE) { // private
+    // work-group linear id
+    const auto WG_LID =
+        __spirv_BuiltInWorkgroupId.x * __spirv_BuiltInNumWorkgroups.y *
+            __spirv_BuiltInNumWorkgroups.z +
+        __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z +
+        __spirv_BuiltInWorkgroupId.z;
 
-  if (as == ADDRESS_SPACE_GENERIC) {
-    auto old = addr;
-    if ((addr = (uptr)ToPrivate((void *)old))) {
-      as = ADDRESS_SPACE_PRIVATE;
-    } else if ((addr = (uptr)ToLocal((void *)old))) {
-      as = ADDRESS_SPACE_LOCAL;
-    } else if ((addr = (uptr)ToGlobal((void *)old))) {
-      as = ADDRESS_SPACE_GLOBAL;
-    } else {
-      if (__AsanDebug)
-        __spirv_ocl_printf(__generic_to_fail, old);
+    auto launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
+    const auto shadow_offset = launch_info->PrivateShadowOffset;
+    const auto shadow_offset_end = launch_info->LocalShadowOffsetEnd;
+
+    if (shadow_offset == 0) {
       return 0;
     }
+
     if (__AsanDebug)
-      __spirv_ocl_printf(__generic_to, old, addr, as);
+      __spirv_ocl_printf(__mem_launch_info, launch_info,
+                         launch_info->PrivateShadowOffset, 0,
+                         launch_info->NumLocalArgs, launch_info->LocalArgs);
+
+    uptr shadow_ptr = shadow_offset +
+                      ((WG_LID * ASAN_PRIVATE_SIZE) >> ASAN_SHADOW_SCALE) +
+                      ((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE);
+
+    if (shadow_ptr > shadow_offset_end) {
+      if (__asan_report_out_of_shadow_bounds()) {
+        __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr,
+                           WG_LID, (uptr)shadow_offset);
+      }
+      return 0;
+    }
+    return shadow_ptr;
+  }
+
+  return 0;
+}
+
+inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
+  if (as == ADDRESS_SPACE_GENERIC) {
+    if (!ConvertGenericPointer(addr, as)) {
+      return 0;
+    }
   }
 
   if (as == ADDRESS_SPACE_GLOBAL) { // global
@@ -262,6 +330,8 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
     shadow_ptr = MemToShadow_CPU(addr);
   } else if (__DeviceType == DeviceType::GPU_PVC) {
     shadow_ptr = MemToShadow_PVC(addr, as);
+  } else if (__DeviceType == DeviceType::GPU_DG2) {
+    shadow_ptr = MemToShadow_DG2(addr, as);
   } else {
     if (__asan_report_unknown_device() && __AsanDebug) {
       __spirv_ocl_printf(__asan_print_unsupport_device_type, (int)__DeviceType);

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -142,8 +142,8 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
 
   message(STATUS "Will fetch Unified Runtime from ${UNIFIED_RUNTIME_REPO}")
   FetchContent_Declare(unified-runtime
-    GIT_REPOSITORY    ${UNIFIED_RUNTIME_REPO}
-    GIT_TAG           ${UNIFIED_RUNTIME_TAG}
+    GIT_REPOSITORY    "https://github.com/AllanZyne/unified-runtime.git"
+    GIT_TAG           "review/yang/dg2"
   )
 
   FetchContent_GetProperties(unified-runtime)

sycl/test-e2e/AddressSanitizer/bad-free/bad-free-host.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
 // RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck %s
 #include <sycl/usm.hpp>

sycl/test-e2e/AddressSanitizer/bad-free/bad-free-minus1.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
 // RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_HOST -O0 -g -o %t

sycl/test-e2e/AddressSanitizer/bad-free/bad-free-plus1.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
 // RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_HOST -O0 -g -o %t

sycl/test-e2e/AddressSanitizer/common/config-red-zone-size.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DUNSAFE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS=redzone:64 %{run} not %t 2>&1 | FileCheck %s
 // RUN: %{build} %device_asan_flags -DSAFE -O0 -g -o %t

sycl/test-e2e/AddressSanitizer/common/demangle-kernel-name.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O2 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t &> %t.txt ; FileCheck --input-file %t.txt %s
 #include <sycl/detail/core.hpp>

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

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O2 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS=debug:1 %{run} %t 2>&1 | FileCheck --check-prefixes CHECK-DEBUG %s
 // RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS=debug:0 %{run} %t 2>&1 | FileCheck %s

sycl/test-e2e/AddressSanitizer/double-free/double-free.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
 // RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_HOST -O0 -g -o %t

sycl/test-e2e/AddressSanitizer/multiple-reports/multiple_kernels.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -Xarch_device -fsanitize-recover=address -O2 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} %t 2>&1 | FileCheck %s
 

sycl/test-e2e/AddressSanitizer/multiple-reports/one_kernel.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -Xarch_device -fsanitize-recover=address -O2 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} %t 2>&1 | FileCheck %s
 

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O2 -g -DUSER_CODE_1 -c -o %t1.o
 // RUN: %{build} %device_asan_flags -O2 -g -DUSER_CODE_2 -c -o %t2.o
 // RUN: %clangxx -fsycl %device_asan_flags -O2 -g %t1.o %t2.o -o %t.out

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu, aspect-fp64
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux, aspect-fp64
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t
@@ -16,7 +17,7 @@
 
 int main() {
   sycl::queue Q;
-  constexpr std::size_t N = 1234567;
+  constexpr std::size_t N = 512;
 #if defined(MALLOC_HOST)
   auto *array = sycl::malloc_host<int>(N, Q);
 #elif defined(MALLOC_SHARED)
@@ -34,7 +35,7 @@ int main() {
   // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
   // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
   // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
-  // CHECK: {{READ of size 4 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(1234567, 0, 0\)}}
+  // CHECK: {{READ of size 4 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(512, 0, 0\)}}
   // CHECK: {{  #0 .* .*parallel_for_int.cpp:}}[[@LINE-7]]
 
   sycl::free(array, Q);

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t
@@ -16,7 +17,7 @@
 
 int main() {
   sycl::queue Q;
-  constexpr std::size_t N = 123456789;
+  constexpr std::size_t N = 1024;
 #if defined(MALLOC_HOST)
   auto *array = sycl::malloc_host<short>(N, Q);
 #elif defined(MALLOC_SHARED)
@@ -34,7 +35,7 @@ int main() {
   // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
   // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
   // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
-  // CHECK: {{READ of size 2 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456789, 0, 0\)}}
+  // CHECK: {{READ of size 2 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(1024, 0, 0\)}}
   // CHECK: {{  #0 .* .*parallel_for_short.cpp:}}[[@LINE-7]]
 
   sycl::free(array, Q);

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_no_local_size.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/unaligned_shadow_memory.cpp

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // RUN: %{build} %device_asan_flags -DTEST1 -O0 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK1 %s
 // RUN: %{build} %device_asan_flags -DTEST2 -O0 -g -o %t

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

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

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

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

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

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

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

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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

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

@@ -1,4 +1,5 @@
-// REQUIRES: linux, cpu
+// UNSUPPORTED: cuda || hip
+// REQUIRES: linux
 // 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