diff --git a/libdevice/include/asan_libdevice.hpp b/libdevice/include/asan_libdevice.hpp
new file mode 100644
index 0000000000000..2ac312252fcd9
--- /dev/null
+++ b/libdevice/include/asan_libdevice.hpp
@@ -0,0 +1,155 @@
+//===---- asan_libdevice.hpp - Structure and declaration for sanitizer ----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include <cinttypes>
+
+// NOTE This file should be sync with
+// unified-runtime/source/loader/layers/sanitizer/device_sanitizer_report.hpp
+
+enum class DeviceSanitizerErrorType : int32_t {
+  UNKNOWN,
+  OUT_OF_BOUNDS,
+  MISALIGNED,
+  USE_AFTER_FREE,
+  OUT_OF_SHADOW_BOUNDS,
+  UNKNOWN_DEVICE,
+  NULL_POINTER,
+};
+
+enum class DeviceSanitizerMemoryType : int32_t {
+  UNKNOWN,
+  USM_DEVICE,
+  USM_HOST,
+  USM_SHARED,
+  LOCAL,
+  PRIVATE,
+  MEM_BUFFER,
+  DEVICE_GLOBAL,
+};
+
+struct DeviceSanitizerReport {
+  int Flag = 0;
+
+  char File[256 + 1] = {};
+  char Func[256 + 1] = {};
+
+  int32_t Line = 0;
+
+  uint64_t GID0 = 0;
+  uint64_t GID1 = 0;
+  uint64_t GID2 = 0;
+
+  uint64_t LID0 = 0;
+  uint64_t LID1 = 0;
+  uint64_t LID2 = 0;
+
+  uintptr_t Address = 0;
+  bool IsWrite = false;
+  uint32_t AccessSize = 0;
+  DeviceSanitizerMemoryType MemoryType = DeviceSanitizerMemoryType::UNKNOWN;
+  DeviceSanitizerErrorType ErrorType = DeviceSanitizerErrorType::UNKNOWN;
+
+  bool IsRecover = false;
+};
+
+struct LocalArgsInfo {
+  uint32_t ArgIndex = 0;
+  uint64_t Size = 0;
+  uint64_t SizeWithRedZone = 0;
+};
+
+struct LaunchInfo {
+  uintptr_t PrivateShadowOffset =
+      0; // don't move this field, we use it in AddressSanitizerPass
+
+  uintptr_t LocalShadowOffset = 0;
+  uintptr_t LocalShadowOffsetEnd = 0;
+  DeviceSanitizerReport SanitizerReport;
+
+  uint32_t NumLocalArgs = 0;
+  LocalArgsInfo *LocalArgs = nullptr; // ordered by ArgIndex
+};
+
+constexpr unsigned ASAN_SHADOW_SCALE = 3;
+constexpr unsigned ASAN_SHADOW_GRANULARITY = 1ULL << ASAN_SHADOW_SCALE;
+
+// Based on the observation, only the last 24 bits of the address of the private
+// variable have changed, we use 31 bits(2G) to be safe.
+constexpr std::size_t ASAN_PRIVATE_SIZE = 0x7fffffffULL + 1;
+
+// These magic values are written to shadow for better error
+// reporting.
+constexpr int kUsmDeviceRedzoneMagic = (char)0x81;
+constexpr int kUsmHostRedzoneMagic = (char)0x82;
+constexpr int kUsmSharedRedzoneMagic = (char)0x83;
+constexpr int kMemBufferRedzoneMagic = (char)0x84;
+constexpr int kDeviceGlobalRedzoneMagic = (char)0x85;
+constexpr int kNullPointerRedzoneMagic = (char)0x86;
+
+constexpr int kUsmDeviceDeallocatedMagic = (char)0x91;
+constexpr int kUsmHostDeallocatedMagic = (char)0x92;
+constexpr int kUsmSharedDeallocatedMagic = (char)0x93;
+constexpr int kMemBufferDeallocatedMagic = (char)0x93;
+
+constexpr int kSharedLocalRedzoneMagic = (char)0xa1;
+
+// Same with host ASan stack
+const int kPrivateLeftRedzoneMagic = (char)0xf1;
+const int kPrivateMidRedzoneMagic = (char)0xf2;
+const int kPrivateRightRedzoneMagic = (char)0xf3;
+
+constexpr auto kSPIR_AsanShadowMemoryGlobalStart =
+    "__AsanShadowMemoryGlobalStart";
+constexpr auto kSPIR_AsanShadowMemoryGlobalEnd = "__AsanShadowMemoryGlobalEnd";
+
+constexpr auto kSPIR_DeviceType = "__DeviceType";
+constexpr auto kSPIR_AsanDebug = "__AsanDebug";
+
+constexpr auto kSPIR_AsanDeviceGlobalCount = "__AsanDeviceGlobalCount";
+constexpr auto kSPIR_AsanDeviceGlobalMetadata = "__AsanDeviceGlobalMetadata";
+
+inline const char *ToString(DeviceSanitizerMemoryType MemoryType) {
+  switch (MemoryType) {
+  case DeviceSanitizerMemoryType::USM_DEVICE:
+    return "Device USM";
+  case DeviceSanitizerMemoryType::USM_HOST:
+    return "Host USM";
+  case DeviceSanitizerMemoryType::USM_SHARED:
+    return "Shared USM";
+  case DeviceSanitizerMemoryType::LOCAL:
+    return "Local Memory";
+  case DeviceSanitizerMemoryType::PRIVATE:
+    return "Private Memory";
+  case DeviceSanitizerMemoryType::MEM_BUFFER:
+    return "Memory Buffer";
+  case DeviceSanitizerMemoryType::DEVICE_GLOBAL:
+    return "Device Global";
+  default:
+    return "Unknown Memory";
+  }
+}
+
+inline const char *ToString(DeviceSanitizerErrorType ErrorType) {
+  switch (ErrorType) {
+  case DeviceSanitizerErrorType::OUT_OF_BOUNDS:
+    return "out-of-bounds-access";
+  case DeviceSanitizerErrorType::MISALIGNED:
+    return "misaligned-access";
+  case DeviceSanitizerErrorType::USE_AFTER_FREE:
+    return "use-after-free";
+  case DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS:
+    return "out-of-shadow-bounds-access";
+  case DeviceSanitizerErrorType::UNKNOWN_DEVICE:
+    return "unknown-device";
+  case DeviceSanitizerErrorType::NULL_POINTER:
+    return "null-pointer-access";
+  default:
+    return "unknown-error";
+  }
+}
diff --git a/libdevice/include/device-sanitizer-report.hpp b/libdevice/include/device-sanitizer-report.hpp
deleted file mode 100644
index bc4f286ce9525..0000000000000
--- a/libdevice/include/device-sanitizer-report.hpp
+++ /dev/null
@@ -1,56 +0,0 @@
-//==-- device-sanitizer-report.hpp - Structure and declaration for assert
-// support --==//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-#pragma once
-
-#include <cinttypes>
-
-enum class DeviceSanitizerErrorType : int32_t {
-  UNKNOWN,
-  OUT_OF_BOUND,
-  MISALIGNED,
-  USE_AFTER_FREE,
-  OUT_OF_SHADOW_BOUND,
-};
-
-enum class DeviceSanitizerMemoryType : int32_t {
-  UNKNOWN,
-  USM_DEVICE,
-  USM_HOST,
-  USM_SHARED,
-  LOCAL,
-  PRIVATE,
-  MEM_BUFFER,
-  DEVICE_GLOBAL,
-};
-
-// NOTE Layout of this structure should be aligned with the one in
-// sycl/include/sycl/detail/device_sanitizer_report.hpp
-struct DeviceSanitizerReport {
-  int Flag = 0;
-
-  char File[256 + 1] = "";
-  char Func[256 + 1] = "";
-
-  int32_t Line = 0;
-
-  uint64_t GID0 = 0;
-  uint64_t GID1 = 0;
-  uint64_t GID2 = 0;
-
-  uint64_t LID0 = 0;
-  uint64_t LID1 = 0;
-  uint64_t LID2 = 0;
-
-  bool IsWrite = false;
-  uint32_t AccessSize = 0;
-  DeviceSanitizerMemoryType MemoryType = DeviceSanitizerMemoryType::UNKNOWN;
-  DeviceSanitizerErrorType ErrorType = DeviceSanitizerErrorType::UNKNOWN;
-
-  bool IsRecover = false;
-};
diff --git a/libdevice/include/sanitizer_device_utils.hpp b/libdevice/include/sanitizer_device_utils.hpp
index 0cb26e7cf3cf8..9d496d56513f2 100644
--- a/libdevice/include/sanitizer_device_utils.hpp
+++ b/libdevice/include/sanitizer_device_utils.hpp
@@ -40,4 +40,4 @@ class
   T val;
 };
 
-enum DeviceType : uintptr_t { UNKNOWN, CPU, GPU_PVC, GPU_DG2 };
+enum DeviceType : uint64_t { UNKNOWN, CPU, GPU_PVC, GPU_DG2 };
diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp
index c560d1d731bdc..3112fb7e33517 100644
--- a/libdevice/sanitizer_utils.cpp
+++ b/libdevice/sanitizer_utils.cpp
@@ -10,9 +10,8 @@
 #include "device.h"
 #include "spirv_vars.h"
 
-#include "include/device-sanitizer-report.hpp"
+#include "include/asan_libdevice.hpp"
 #include "include/sanitizer_device_utils.hpp"
-
 #include <cstddef>
 #include <cstdint>
 
@@ -22,25 +21,29 @@ using u8 = unsigned char;
 using s16 = short;
 using u16 = unsigned short;
 
-#define ASAN_SHADOW_SCALE 3
-#define ASAN_SHADOW_GRANULARITY (1ULL << ASAN_SHADOW_SCALE)
-
 DeviceGlobal<uptr> __AsanShadowMemoryGlobalStart;
 DeviceGlobal<uptr> __AsanShadowMemoryGlobalEnd;
 DeviceGlobal<uptr> __AsanShadowMemoryLocalStart;
 DeviceGlobal<uptr> __AsanShadowMemoryLocalEnd;
-
-DeviceGlobal<DeviceSanitizerReport> __DeviceSanitizerReportMem;
-
 DeviceGlobal<DeviceType> __DeviceType;
+DeviceGlobal<uint64_t> __AsanDebug;
+DeviceGlobal<DeviceSanitizerReport> __DeviceSanitizerReportMem;
 
 #if defined(__SPIR__) || defined(__SPIRV__)
 
-#ifdef __SYCL_DEVICE_ONLY__
+#if defined(__SYCL_DEVICE_ONLY__)
+
 #define __USE_SPIR_BUILTIN__ 1
-#else
+
+#ifndef SYCL_EXTERNAL
+#define SYCL_EXTERNAL
+#endif // SYCL_EXTERNAL
+
+#else // __SYCL_DEVICE_ONLY__
+
 #define __USE_SPIR_BUILTIN__
-#endif
+
+#endif // __SYCL_DEVICE_ONLY__
 
 #if __USE_SPIR_BUILTIN__
 extern SYCL_EXTERNAL int
@@ -52,54 +55,52 @@ extern SYCL_EXTERNAL __SYCL_LOCAL__ void *
 __spirv_GenericCastToPtrExplicit_ToLocal(void *, int);
 extern SYCL_EXTERNAL __SYCL_PRIVATE__ void *
 __spirv_GenericCastToPtrExplicit_ToPrivate(void *, int);
-#endif
-
-// These magic values are written to shadow for better error
-// reporting.
-const int kUsmDeviceRedzoneMagic = (char)0x81;
-const int kUsmHostRedzoneMagic = (char)0x82;
-const int kUsmSharedRedzoneMagic = (char)0x83;
-const int kMemBufferRedzoneMagic = (char)0x84;
-const int kDeviceGlobalRedZoneMagic = (char)0x85;
-
-const int kUsmDeviceDeallocatedMagic = (char)0x91;
-const int kUsmHostDeallocatedMagic = (char)0x92;
-const int kUsmSharedDeallocatedMagic = (char)0x93;
-
-const int kSharedLocalRedzoneMagic = (char)0xa1;
-
-// Same with Asan Stack
-const int kPrivateLeftRedzoneMagic = (char)0xf1;
-const int kPrivateMidRedzoneMagic = (char)0xf2;
-const int kPrivateRightRedzoneMagic = (char)0xf3;
+#endif // __USE_SPIR_BUILTIN__
 
 static const __SYCL_CONSTANT__ char __asan_shadow_value_start[] =
-    "%p(%d) -> %p:";
+    "[kernel] %p(%d) -> %p:";
 static const __SYCL_CONSTANT__ char __asan_shadow_value[] = " %02X";
 static const __SYCL_CONSTANT__ char __asan_current_shadow_value[] = ">%02X";
 static const __SYCL_CONSTANT__ char __newline[] = "\n";
 
 static const __SYCL_CONSTANT__ char __global_shadow_out_of_bound[] =
-    "ERROR: Global shadow memory out-of-bound (ptr: %p -> %p, base: %p)\n";
+    "[kernel] Global shadow memory out-of-bound (ptr: %p -> %p, base: %p)\n";
 static const __SYCL_CONSTANT__ char __local_shadow_out_of_bound[] =
-    "ERROR: Local shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: "
+    "[kernel] Local shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: "
     "%p)\n";
 
-static const __SYCL_CONSTANT__ char __unsupport_device_type[] =
-    "ERROR: Unsupport device type: %d\n";
+static const __SYCL_CONSTANT__ char __asan_print_unsupport_device_type[] =
+    "[kernel] Unsupport device type: %d\n";
+
+static const __SYCL_CONSTANT__ char __asan_print_shadow_value1[] =
+    "[kernel] %p(%d) -> %p: %02X\n";
+static const __SYCL_CONSTANT__ char __asan_print_shadow_value2[] =
+    "[kernel] %p(%d) -> %p: --\n";
+
+static __SYCL_CONSTANT__ const char __generic_to[] =
+    "[kernel] %p(4) - %p(%d)\n";
+
+static __SYCL_CONSTANT__ const char __generic_to_fail[] =
+    "[kernel] %p(4) - unknown address space\n";
 
 #define ASAN_REPORT_NONE 0
 #define ASAN_REPORT_START 1
 #define ASAN_REPORT_FINISH 2
 
-#define AS_PRIVATE 0
-#define AS_GLOBAL 1
-#define AS_CONSTANT 2
-#define AS_LOCAL 3
-#define AS_GENERIC 4
+enum ADDRESS_SPACE : uint32_t {
+  ADDRESS_SPACE_PRIVATE = 0,
+  ADDRESS_SPACE_GLOBAL = 1,
+  ADDRESS_SPACE_CONSTANT = 2,
+  ADDRESS_SPACE_LOCAL = 3,
+  ADDRESS_SPACE_GENERIC = 4,
+};
 
 namespace {
 
+bool __asan_report_unknown_device();
+bool __asan_report_out_of_shadow_bounds();
+void __asan_print_shadow_memory(uptr addr, uptr shadow_address, uint32_t as);
+
 __SYCL_GLOBAL__ void *ToGlobal(void *ptr) {
   return __spirv_GenericCastToPtrExplicit_ToGlobal(ptr, 5);
 }
@@ -110,11 +111,11 @@ __SYCL_PRIVATE__ void *ToPrivate(void *ptr) {
   return __spirv_GenericCastToPtrExplicit_ToPrivate(ptr, 7);
 }
 
-inline uptr MemToShadow_CPU(uptr addr, int32_t as) {
+inline uptr MemToShadow_CPU(uptr addr) {
   return __AsanShadowMemoryGlobalStart + (addr >> 3);
 }
 
-inline uptr MemToShadow_DG2(uptr addr, int32_t as) {
+inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
   uptr shadow_ptr = 0;
   if (addr & (~0xffffffffffff)) {
     shadow_ptr =
@@ -125,77 +126,107 @@ inline uptr MemToShadow_DG2(uptr addr, int32_t as) {
   }
 
   if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
-    __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr);
+    if (__asan_report_out_of_shadow_bounds() && __AsanDebug) {
+      __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr);
+    }
   }
 
   return shadow_ptr;
 }
 
-inline uptr MemToShadow_PVC(uptr addr, int32_t as) {
-  uptr shadow_ptr = 0;
+inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
 
-  if (as == AS_GENERIC) {
-    if ((shadow_ptr = (uptr)ToGlobal((void *)addr))) {
-      as = AS_GLOBAL;
-    } else if ((shadow_ptr = (uptr)ToPrivate((void *)addr))) {
-      as = AS_PRIVATE;
-    } else if ((shadow_ptr = (uptr)ToLocal((void *)addr))) {
-      as = AS_LOCAL;
+  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);
       return 0;
     }
+    if (__AsanDebug)
+      __spirv_ocl_printf(__generic_to, old, addr, as);
   }
 
-  if (as == AS_PRIVATE) {            // private
-  } else if (as == AS_GLOBAL) {      // global
+  if (as == ADDRESS_SPACE_GLOBAL) { // global
+    uptr shadow_ptr;
     if (addr & 0xFF00000000000000) { // Device USM
       shadow_ptr = __AsanShadowMemoryGlobalStart + 0x200000000000 +
                    ((addr & 0xFFFFFFFFFFFF) >> 3);
     } else { // Only consider 47bit VA
-      shadow_ptr =
-          __AsanShadowMemoryGlobalStart + ((addr & 0x7FFFFFFFFFFF) >> 3);
+      shadow_ptr = __AsanShadowMemoryGlobalStart +
+                   ((addr & 0x7FFFFFFFFFFF) >> ASAN_SHADOW_SCALE);
     }
 
     if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
-      __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
-                         (uptr)__AsanShadowMemoryGlobalStart);
-      shadow_ptr = 0;
+      if (__asan_report_out_of_shadow_bounds() && __AsanDebug) {
+        __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
+                           (uptr)__AsanShadowMemoryGlobalStart);
+      }
+      return 0;
+    }
+    return shadow_ptr;
+  } else if (as == ADDRESS_SPACE_LOCAL) { // local
+    if (__AsanShadowMemoryLocalStart == 0) {
+      return 0;
     }
-  } else if (as == AS_CONSTANT) { // constant
-  } else if (as == AS_LOCAL) {    // local
     // The size of SLM is 128KB on PVC
-    constexpr unsigned slm_size = 128 * 1024;
+    constexpr unsigned SLM_SIZE = 128 * 1024;
+    // 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;
 
-    shadow_ptr = __AsanShadowMemoryLocalStart + ((wg_lid * slm_size) >> 3) +
-                 ((addr & (slm_size - 1)) >> 3);
+    uptr shadow_ptr = __AsanShadowMemoryLocalStart +
+                      ((wg_lid * SLM_SIZE) >> ASAN_SHADOW_SCALE) +
+                      ((addr & (SLM_SIZE - 1)) >> 3);
 
     if (shadow_ptr > __AsanShadowMemoryLocalEnd) {
-      __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr, wg_lid,
-                         (uptr)__AsanShadowMemoryLocalStart);
-      shadow_ptr = 0;
+      if (__asan_report_out_of_shadow_bounds() && __AsanDebug) {
+        __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
+                           wg_lid, (uptr)__AsanShadowMemoryLocalStart);
+      }
+      return 0;
     }
+    return shadow_ptr;
   }
 
-  return shadow_ptr;
+  return 0;
 }
 
-inline uptr MemToShadow(uptr addr, int32_t as) {
+inline uptr MemToShadow(uptr addr, uint32_t as) {
   uptr shadow_ptr = 0;
 
   if (__DeviceType == DeviceType::CPU) {
-    shadow_ptr = MemToShadow_CPU(addr, as);
+    shadow_ptr = MemToShadow_CPU(addr);
   } else if (__DeviceType == DeviceType::GPU_PVC) {
     shadow_ptr = MemToShadow_PVC(addr, as);
   } else {
-    __spirv_ocl_printf(__unsupport_device_type, (int)__DeviceType);
+    if (__asan_report_unknown_device() && __AsanDebug) {
+      __spirv_ocl_printf(__asan_print_unsupport_device_type, (int)__DeviceType);
+    }
     return shadow_ptr;
   }
 
+  if (__AsanDebug) {
+    if (shadow_ptr) {
+      if (as == ADDRESS_SPACE_PRIVATE)
+        __asan_print_shadow_memory(addr, shadow_ptr, as);
+      else
+        __spirv_ocl_printf(__asan_print_shadow_value1, addr, as, shadow_ptr,
+                           *(u8 *)shadow_ptr);
+    } else {
+      __spirv_ocl_printf(__asan_print_shadow_value2, addr, as, shadow_ptr);
+    }
+  }
+
   return shadow_ptr;
 }
 
@@ -207,64 +238,54 @@ inline constexpr uptr RoundDownTo(uptr x, uptr boundary) {
   return x & ~(boundary - 1);
 }
 
-bool MemIsZero(const char *beg, uptr size) {
-  const char *end = beg + size;
-  uptr *aligned_beg = (uptr *)RoundUpTo((uptr)beg, sizeof(uptr));
-  uptr *aligned_end = (uptr *)RoundDownTo((uptr)end, sizeof(uptr));
+bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) {
+  __SYCL_GLOBAL__ const char *end = beg + size;
+  auto *aligned_beg =
+      (__SYCL_GLOBAL__ uptr *)RoundUpTo((uptr)beg, sizeof(uptr));
+  auto *aligned_end =
+      (__SYCL_GLOBAL__ uptr *)RoundDownTo((uptr)end, sizeof(uptr));
   uptr all = 0;
   // Prologue.
-  for (const char *mem = beg; mem < (char *)aligned_beg && mem < end; mem++)
+  for (__SYCL_GLOBAL__ const char *mem = beg;
+       mem < (__SYCL_GLOBAL__ char *)aligned_beg && mem < end; mem++)
     all |= *mem;
   // Aligned loop.
   for (; aligned_beg < aligned_end; aligned_beg++)
     all |= *aligned_beg;
   // Epilogue.
-  if ((char *)aligned_end >= beg) {
-    for (const char *mem = (char *)aligned_end; mem < end; mem++)
+  if ((__SYCL_GLOBAL__ char *)aligned_end >= beg) {
+    for (__SYCL_GLOBAL__ const char *mem = (__SYCL_GLOBAL__ char *)aligned_end;
+         mem < end; mem++)
       all |= *mem;
   }
   return all == 0;
 }
 
-void print_shadow_memory(uptr addr, int32_t as) {
-  uptr shadow_address = MemToShadow(addr, as);
-  uptr p = shadow_address & (~0xf);
-  __spirv_ocl_printf(__asan_shadow_value_start, addr, as, p);
-  for (int i = 0; i < 0xf; ++i) {
-    u8 shadow_value = *(u8 *)(p + i);
-    if (p + i == shadow_address) {
-      __spirv_ocl_printf(__asan_current_shadow_value, shadow_value);
-    } else {
-      __spirv_ocl_printf(__asan_shadow_value, shadow_value);
-    }
-  }
-  __spirv_ocl_printf(__newline);
-}
-
-} // namespace
+///
+/// ASAN Save Report
+///
 
-bool __asan_region_is_value(uptr addr, int32_t as, std::size_t size,
-                            char value) {
-  if (size == 0)
+bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
+  const int Expected = ASAN_REPORT_NONE;
+  int Desired = ASAN_REPORT_START;
+  if (atomicCompareAndSet(&__DeviceSanitizerReportMem.get().Flag, Desired,
+                          Expected) == Expected) {
+    __DeviceSanitizerReportMem.get().ErrorType = error_type;
+    // Show we've done copying
+    atomicStore(&__DeviceSanitizerReportMem.get().Flag, ASAN_REPORT_FINISH);
     return true;
-  while (size--) {
-    char *shadow = (char *)MemToShadow(addr, as);
-    if (*shadow != value) {
-      return false;
-    }
-    ++addr;
   }
-  return true;
+  return false;
 }
 
 #ifdef __SYCL_DEVICE_ONLY__
 #define __DEVICE_SANITIZER_REPORT_ACCESSOR __DeviceSanitizerReportMem.get()
-#else
+#else // __SYCL_DEVICE_ONLY__
 #define __DEVICE_SANITIZER_REPORT_ACCESSOR
-#endif
+#endif // __SYCL_DEVICE_ONLY__
 
-static void __asan_internal_report_save(
-    uptr ptr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line,
+bool __asan_internal_report_save(
+    uptr ptr, uint32_t as, const char __SYCL_CONSTANT__ *file, uint32_t line,
     const char __SYCL_CONSTANT__ *func, bool is_write, uint32_t access_size,
     DeviceSanitizerMemoryType memory_type, DeviceSanitizerErrorType error_type,
     bool is_recover = false) {
@@ -308,6 +329,7 @@ static void __asan_internal_report_save(
     __DEVICE_SANITIZER_REPORT_ACCESSOR.LID1 = __spirv_LocalInvocationId_y();
     __DEVICE_SANITIZER_REPORT_ACCESSOR.LID2 = __spirv_LocalInvocationId_z();
 
+    __DEVICE_SANITIZER_REPORT_ACCESSOR.Address = ptr;
     __DEVICE_SANITIZER_REPORT_ACCESSOR.IsWrite = is_write;
     __DEVICE_SANITIZER_REPORT_ACCESSOR.AccessSize = access_size;
     __DEVICE_SANITIZER_REPORT_ACCESSOR.ErrorType = error_type;
@@ -317,20 +339,21 @@ static void __asan_internal_report_save(
     // Show we've done copying
     atomicStore(&__DEVICE_SANITIZER_REPORT_ACCESSOR.Flag, ASAN_REPORT_FINISH);
   }
+  return false;
 }
 
 ///
 /// ASAN Error Reporters
 ///
 
-void __asan_report_access_error(uptr addr, int32_t as, size_t size,
+void __asan_report_access_error(uptr addr, uint32_t as, size_t size,
                                 bool is_write, uptr poisoned_addr,
                                 const char __SYCL_CONSTANT__ *file,
-                                int32_t line,
+                                uint32_t line,
                                 const char __SYCL_CONSTANT__ *func,
                                 bool is_recover = false) {
   // Check Error Type
-  s8 *shadow_address = (s8 *)MemToShadow(poisoned_addr, as);
+  auto *shadow_address = (__SYCL_GLOBAL__ s8 *)MemToShadow(poisoned_addr, as);
   int shadow_value = *shadow_address;
   if (shadow_value > 0) {
     shadow_value = *(shadow_address + 1);
@@ -343,15 +366,15 @@ void __asan_report_access_error(uptr addr, int32_t as, size_t size,
   switch (shadow_value) {
   case kUsmDeviceRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::USM_DEVICE;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
   case kUsmHostRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::USM_HOST;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
   case kUsmSharedRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::USM_SHARED;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
   case kUsmDeviceDeallocatedMagic:
     memory_type = DeviceSanitizerMemoryType::USM_DEVICE;
@@ -369,19 +392,19 @@ void __asan_report_access_error(uptr addr, int32_t as, size_t size,
   case kPrivateMidRedzoneMagic:
   case kPrivateRightRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::PRIVATE;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
   case kMemBufferRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::MEM_BUFFER;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
   case kSharedLocalRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::LOCAL;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
-  case kDeviceGlobalRedZoneMagic:
+  case kDeviceGlobalRedzoneMagic:
     memory_type = DeviceSanitizerMemoryType::DEVICE_GLOBAL;
-    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
     break;
   default:
     memory_type = DeviceSanitizerMemoryType::UNKNOWN;
@@ -392,13 +415,50 @@ void __asan_report_access_error(uptr addr, int32_t as, size_t size,
                               memory_type, error_type, is_recover);
 }
 
+bool __asan_report_unknown_device() {
+  return __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE);
+}
+
+bool __asan_report_out_of_shadow_bounds() {
+  return __asan_internal_report_save(
+      DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS);
+}
+
 ///
-/// Check if memory is poisoned
+/// ASan utils
 ///
 
+void __asan_print_shadow_memory(uptr addr, uptr shadow_address, uint32_t as) {
+  uptr p = shadow_address & (~0xf);
+  __spirv_ocl_printf(__asan_shadow_value_start, addr, as, p);
+  for (int i = 0; i < 0xf; ++i) {
+    u8 shadow_value = *(u8 *)(p + i);
+    if (p + i == shadow_address) {
+      __spirv_ocl_printf(__asan_current_shadow_value, shadow_value);
+    } else {
+      __spirv_ocl_printf(__asan_shadow_value, shadow_value);
+    }
+  }
+  __spirv_ocl_printf(__newline);
+}
+
+bool __asan_region_is_value(uptr addr, uint32_t as, std::size_t size,
+                            char value) {
+  if (size == 0)
+    return true;
+  while (size--) {
+    auto *shadow = (__SYCL_GLOBAL__ char *)MemToShadow(addr, as);
+    if (*shadow != value) {
+      return false;
+    }
+    ++addr;
+  }
+  return true;
+}
+
 // NOTE: size < 8
-inline int __asan_address_is_poisoned(uptr a, int32_t as, size_t size) {
-  auto *shadow_address = (s8 *)MemToShadow(a, as);
+inline int __asan_address_is_poisoned(uptr a, uint32_t as, size_t size) {
+  auto *shadow_address = (__SYCL_GLOBAL__ s8 *)MemToShadow(a, as);
   if (shadow_address) {
     auto shadow_value = *shadow_address;
     if (shadow_value) {
@@ -410,11 +470,11 @@ inline int __asan_address_is_poisoned(uptr a, int32_t as, size_t size) {
 }
 
 // NOTE: size = 1
-inline int __asan_address_is_poisoned(uptr a, int32_t as) {
+inline int __asan_address_is_poisoned(uptr a, uint32_t as) {
   return __asan_address_is_poisoned(a, as, 1);
 }
 
-inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) {
+inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) {
   if (!size)
     return 0;
 
@@ -437,7 +497,8 @@ inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) {
   if (!__asan_address_is_poisoned(beg, as) &&
       !__asan_address_is_poisoned(end - 1, as) &&
       (shadow_end <= shadow_beg ||
-       MemIsZero((const char *)shadow_beg, shadow_end - shadow_beg)))
+       MemIsZero((__SYCL_GLOBAL__ const char *)shadow_beg,
+                 shadow_end - shadow_beg)))
     return 0;
 
   // The fast check failed, so we have a poisoned byte somewhere.
@@ -449,22 +510,24 @@ inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) {
   return 0;
 }
 
+} // namespace
+
 ///
 /// ASAN Load/Store Report Built-ins
 ///
 
 #define ASAN_REPORT_ERROR(type, is_write, size)                                \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size(                           \
-      uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \
-      const char __SYCL_CONSTANT__ *func) {                                    \
+      uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
     if (__asan_address_is_poisoned(addr, as, size)) {                          \
       __asan_report_access_error(addr, as, size, is_write, addr, file, line,   \
                                  func);                                        \
     }                                                                          \
   }                                                                            \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort(                 \
-      uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \
-      const char __SYCL_CONSTANT__ *func) {                                    \
+      uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
     if (__asan_address_is_poisoned(addr, as, size)) {                          \
       __asan_report_access_error(addr, as, size, is_write, addr, file, line,   \
                                  func, true);                                  \
@@ -480,8 +543,8 @@ ASAN_REPORT_ERROR(store, true, 4)
 
 #define ASAN_REPORT_ERROR_BYTE(type, is_write, size)                           \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size(                           \
-      uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \
-      const char __SYCL_CONSTANT__ *func) {                                    \
+      uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
     u##size *shadow_address = (u##size *)MemToShadow(addr, as);                \
     if (shadow_address && *shadow_address) {                                   \
       __asan_report_access_error(addr, as, size, is_write, addr, file, line,   \
@@ -489,8 +552,8 @@ ASAN_REPORT_ERROR(store, true, 4)
     }                                                                          \
   }                                                                            \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort(                 \
-      uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \
-      const char __SYCL_CONSTANT__ *func) {                                    \
+      uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
     u##size *shadow_address = (u##size *)MemToShadow(addr, as);                \
     if (shadow_address && *shadow_address) {                                   \
       __asan_report_access_error(addr, as, size, is_write, addr, file, line,   \
@@ -505,16 +568,16 @@ ASAN_REPORT_ERROR_BYTE(store, true, 16)
 
 #define ASAN_REPORT_ERROR_N(type, is_write)                                    \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##N(                              \
-      uptr addr, size_t size, int32_t as, const char __SYCL_CONSTANT__ *file,  \
-      int32_t line, const char __SYCL_CONSTANT__ *func) {                      \
+      uptr addr, size_t size, uint32_t as, const char __SYCL_CONSTANT__ *file, \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
     if (auto poisoned_addr = __asan_region_is_poisoned(addr, as, size)) {      \
       __asan_report_access_error(addr, as, size, is_write, poisoned_addr,      \
                                  file, line, func);                            \
     }                                                                          \
   }                                                                            \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##N_noabort(                      \
-      uptr addr, size_t size, int32_t as, const char __SYCL_CONSTANT__ *file,  \
-      int32_t line, const char __SYCL_CONSTANT__ *func) {                      \
+      uptr addr, size_t size, uint32_t as, const char __SYCL_CONSTANT__ *file, \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
     if (auto poisoned_addr = __asan_region_is_poisoned(addr, as, size)) {      \
       __asan_report_access_error(addr, as, size, is_write, poisoned_addr,      \
                                  file, line, func, true);                      \
@@ -524,24 +587,55 @@ ASAN_REPORT_ERROR_BYTE(store, true, 16)
 ASAN_REPORT_ERROR_N(load, false)
 ASAN_REPORT_ERROR_N(store, true)
 
+///
+/// ASAN initialize shdadow memory of local memory
+///
+
+static __SYCL_CONSTANT__ const char __mem_set_shadow_local[] =
+    "[kernel] set_shadow_local(beg=%p, end=%p, val:%02X)\n";
+
 DEVICE_EXTERN_C_NOINLINE void
 __asan_set_shadow_local_memory(uptr ptr, size_t size,
                                size_t size_with_redzone) {
+  // Since ptr is aligned to ASAN_SHADOW_GRANULARITY,
+  // if size != aligned_size, then the buffer tail of ptr is not aligned
   uptr aligned_size = RoundUpTo(size, ASAN_SHADOW_GRANULARITY);
 
+  // Set user zone to zero
+  {
+    auto shadow_begin = MemToShadow(ptr, ADDRESS_SPACE_LOCAL);
+    auto shadow_end = MemToShadow(ptr + size, ADDRESS_SPACE_LOCAL);
+    if (__AsanDebug)
+      __spirv_ocl_printf(__mem_set_shadow_local, shadow_begin, shadow_end, 0);
+    while (shadow_begin <= shadow_end) {
+      *((__SYCL_GLOBAL__ u8 *)shadow_begin) = 0;
+      ++shadow_begin;
+    }
+  }
+
+  // Set left red zone
   {
-    auto shadow_address = MemToShadow(ptr + aligned_size, AS_LOCAL);
-    auto count = (size_with_redzone - aligned_size) / ASAN_SHADOW_GRANULARITY;
+    auto shadow_address = MemToShadow(ptr + aligned_size, ADDRESS_SPACE_LOCAL);
+    auto count = (size_with_redzone - aligned_size) >> ASAN_SHADOW_SCALE;
+    if (__AsanDebug)
+      __spirv_ocl_printf(__mem_set_shadow_local, shadow_address,
+                         shadow_address + count,
+                         (unsigned char)kSharedLocalRedzoneMagic);
     for (size_t i = 0; i < count; ++i) {
-      ((u8 *)shadow_address)[i] = kSharedLocalRedzoneMagic;
+      ((__SYCL_GLOBAL__ u8 *)shadow_address)[i] = kSharedLocalRedzoneMagic;
     }
   }
 
+  // Set unaligned tail
   if (size != aligned_size) {
-    auto user_end = ptr + size - 1;
-    auto *shadow_end = (s8 *)MemToShadow(user_end, AS_LOCAL);
-    *shadow_end = user_end - RoundDownTo(user_end, ASAN_SHADOW_GRANULARITY);
+    auto user_end = ptr + size;
+    auto *shadow_end =
+        (__SYCL_GLOBAL__ s8 *)MemToShadow(user_end, ADDRESS_SPACE_LOCAL);
+    auto value = user_end - RoundDownTo(user_end, ASAN_SHADOW_GRANULARITY) + 1;
+    if (__AsanDebug)
+      __spirv_ocl_printf(__mem_set_shadow_local, shadow_end, shadow_end, value);
+    *shadow_end = value;
   }
 }
 
-#endif
+#endif // __SPIR__ || __SPIRV__
diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
index c0b50579ff546..c895556eb6d01 100644
--- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
@@ -69,6 +69,7 @@
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/MathExtras.h"
+#include "llvm/Support/Path.h"
 #include "llvm/Support/raw_ostream.h"
 #include "llvm/TargetParser/Triple.h"
 #include "llvm/Transforms/Instrumentation.h"
@@ -1320,11 +1321,14 @@ static bool isUnsupportedAMDGPUAddrspace(Value *Addr) {
 }
 
 static bool isUnsupportedSPIRAccess(Value *Addr, Function *Func) {
+  std::ignore = Func;
   Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
   // Private address space: skip kernel arguments
   if (PtrTy->getPointerAddressSpace() == 0) {
-    return Func->getCallingConv() == CallingConv::SPIR_KERNEL &&
-           isa<Argument>(Addr);
+    // FIXME: Currently we don't suppport all private variables
+    // return Func->getCallingConv() == CallingConv::SPIR_KERNEL &&
+    //        isa<Argument>(Addr);
+    return true;
   }
 
   // All the rest address spaces: skip SPIR-V built-in varibles
@@ -1366,9 +1370,10 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore,
 
   // File & Line
   if (Loc) {
-    StringRef FileName = Loc->getFilename();
+    llvm::SmallString<128> Source = Loc->getDirectory();
+    sys::path::append(Source, Loc->getFilename());
     auto *FileNameGV =
-        GetOrCreateGlobalString(*M, "__asan_file", FileName, ConstantAS);
+        GetOrCreateGlobalString(*M, "__asan_file", Source, ConstantAS);
     Args.push_back(ConstantExpr::getPointerCast(FileNameGV, ConstASPtrTy));
     Args.push_back(ConstantInt::get(Type::getInt32Ty(C), Loc.getLine()));
   } else {
@@ -2911,6 +2916,14 @@ bool ModuleAddressSanitizer::instrumentModule(Module &M) {
     }
   }
 
+  if (TargetTriple.isSPIR()) {
+    // Add module metadata "device.sanitizer" for sycl-post-link
+    LLVMContext &Ctx = M.getContext();
+    auto *MD = M.getOrInsertNamedMetadata("device.sanitizer");
+    Metadata *MDVals[] = {MDString::get(Ctx, "asan")};
+    MD->addOperand(MDNode::get(Ctx, MDVals));
+  }
+
   const uint64_t Priority = GetCtorAndDtorPriority(TargetTriple);
 
   // Put the constructor and destructor in comdat if both
@@ -2959,8 +2972,13 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T
         }
       }
 
-      // Extend __asan_load/store arguments: unsigned int address_space, char*
-      // file, unsigned int line, char* func
+      // __asan_loadX/__asan_storeX(
+      //   ...
+      //   int32_t as, // Address Space
+      //   char* file,
+      //   unsigned int line,
+      //   char* func
+      // )
       if (TargetTriple.isSPIR()) {
         constexpr unsigned ConstantAS = 2;
         auto *Int8PtrTy = Type::getInt8Ty(*C)->getPointerTo(ConstantAS);
diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
new file mode 100644
index 0000000000000..002b14076dec2
--- /dev/null
+++ b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
@@ -0,0 +1,76 @@
+; This test checks that the post-link tool properly generates "asanUsed=1"
+; in [SYCL/misc properties]
+
+; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table
+; RUN: FileCheck %s -input-file=%t_0.prop
+; CHECK: [SYCL/misc properties]
+; CHECK: asanUsed=1
+
+; ModuleID = 'parallel_for_int.cpp'
+source_filename = "parallel_for_int.cpp"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E11MyKernelR_4 = comdat any
+
+@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write)
+declare void @llvm.assume(i1 noundef) #0
+
+; Function Attrs: mustprogress norecurse nounwind sanitize_address uwtable
+define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E11MyKernelR_4(ptr addrspace(1) noundef align 4 %_arg_array, i64 %__asan_launch) local_unnamed_addr #1 comdat !srcloc !7 !kernel_arg_buffer_location !8 !sycl_fixed_targets !9 {
+entry:
+  call spir_func void @__itt_offload_wi_start_wrapper()
+  %0 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32, !noalias !10
+  %cmp.i = icmp ult i64 %0, 2147483648
+  tail call void @llvm.assume(i1 %cmp.i)
+  %arrayidx.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_array, i64 %0
+  %1 = ptrtoint ptr addrspace(1) %arrayidx.i to i64
+  call void @__asan_load4(i64 %1, i32 1, i64 %__asan_launch, ptr addrspace(2) null, i32 0, ptr addrspace(2) @__asan_func)
+  %2 = load i32, ptr addrspace(1) %arrayidx.i, align 4, !tbaa !17
+  %inc.i = add nsw i32 %2, 1
+  store i32 %inc.i, ptr addrspace(1) %arrayidx.i, align 4, !tbaa !17
+  call spir_func void @__itt_offload_wi_finish_wrapper()
+  ret void
+}
+
+declare void @__asan_load4(i64, i32, i64, ptr addrspace(2), i32, ptr addrspace(2))
+
+declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(ptr addrspace(2), ...)
+
+declare spir_func void @__itt_offload_wi_start_wrapper()
+
+declare spir_func void @__itt_offload_wi_finish_wrapper()
+
+attributes #0 = { mustprogress nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
+attributes #1 = { mustprogress norecurse nounwind sanitize_address uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="parallel_for_int.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
+
+!llvm.module.flags = !{!0, !1, !2}
+!opencl.spir.version = !{!3}
+!spirv.Source = !{!4}
+!llvm.ident = !{!5}
+!device.sanitizer = !{!6}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"uwtable", i32 2}
+!2 = !{i32 7, !"frame-pointer", i32 2}
+!3 = !{i32 1, i32 2}
+!4 = !{i32 4, i32 100000}
+!5 = !{!"clang version 19.0.0git (https://github.com/intel/llvm f8eada76c08c6a5e6c5842842ac5b98fa72669be)"}
+!6 = !{!"asan"}
+!7 = !{i32 1536}
+!8 = !{i32 -1, i32 -1}
+!9 = !{}
+!10 = !{!11, !13, !15}
+!11 = distinct !{!11, !12, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv: %agg.result"}
+!12 = distinct !{!12, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv"}
+!13 = distinct !{!13, !14, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v: %agg.result"}
+!14 = distinct !{!14, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v"}
+!15 = distinct !{!15, !16, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv: %agg.result"}
+!16 = distinct !{!16, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv"}
+!17 = !{!18, !18, i64 0}
+!18 = !{!"int", !19, i64 0}
+!19 = !{!"omnipotent char", !20, i64 0}
+!20 = !{!"Simple C++ TBAA"}
diff --git a/llvm/test/tools/sycl-post-link/sycl-sanitize/asan.ll b/llvm/test/tools/sycl-post-link/sycl-sanitize/asan.ll
deleted file mode 100644
index 2e06b77131ff0..0000000000000
--- a/llvm/test/tools/sycl-post-link/sycl-sanitize/asan.ll
+++ /dev/null
@@ -1,175 +0,0 @@
-; This test checks that the post-link tool properly generates "asanUsed=1"
-; in [SYCL/misc properties]
-
-; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table
-; RUN: FileCheck %s -input-file=%t_0.prop
-; CHECK: [SYCL/misc properties]
-; CHECK: asanUsed=1
-
-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
-target triple = "spir64-unknown-unknown"
-
-%class.DeviceGlobal = type { i64 }
-
-$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E4Test = comdat any
-
-@__spirv_BuiltInGlobalLinearId = external dso_local local_unnamed_addr addrspace(1) constant i64, align 8
-@__spirv_BuiltInWorkgroupSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
-@__DeviceType = dso_local local_unnamed_addr addrspace(1) global %class.DeviceGlobal zeroinitializer, align 8 #0
-@__AsanShadowMemoryGlobalStart = dso_local local_unnamed_addr addrspace(1) global %class.DeviceGlobal zeroinitializer, align 8 #1
-@__AsanShadowMemoryGlobalEnd = dso_local local_unnamed_addr addrspace(1) global %class.DeviceGlobal zeroinitializer, align 8 #2
-@__spirv_BuiltInWorkgroupId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
-@__AsanShadowMemoryLocalStart = dso_local local_unnamed_addr addrspace(1) global %class.DeviceGlobal zeroinitializer, align 8 #3
-@__AsanShadowMemoryLocalEnd = dso_local local_unnamed_addr addrspace(1) global %class.DeviceGlobal zeroinitializer, align 8 #4
-@__DeviceSanitizerReportMem = dso_local addrspace(1) global { { i32, [257 x i8], [257 x i8], i32, i64, i64, i64, i64, i64, i64, i8, i32, i32, i32, i8 } } zeroinitializer, align 8 #5
-
-; Function Attrs: mustprogress norecurse nounwind sanitize_address uwtable
-define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E4Test() local_unnamed_addr #6 comdat !srcloc !7 !kernel_arg_buffer_location !8 !sycl_fixed_targets !8 !sycl_kernel_omit_args !8 {
-entry:
-  call spir_func void @__itt_offload_wi_start_wrapper()
-  call spir_func void @__itt_offload_wi_finish_wrapper()
-  ret void
-}
-
-; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
-declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #7
-
-; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
-declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #7
-
-; Function Attrs: convergent nounwind
-declare dso_local spir_func signext i8 @__spirv_SpecConstant(i32 noundef, i8 noundef signext) local_unnamed_addr #8
-
-; Function Attrs: alwaysinline convergent mustprogress norecurse nounwind
-define dso_local spir_func void @__itt_offload_wi_start_wrapper() #9 !srcloc !9 {
-entry:
-  %GroupID = alloca [3 x i64], align 8
-  %call.i = tail call spir_func signext i8 @__spirv_SpecConstant(i32 noundef -9145239, i8 noundef signext 0) #11
-  %cmp.i.not = icmp eq i8 %call.i, 0
-  br i1 %cmp.i.not, label %return, label %if.end
-
-if.end:                                           ; preds = %entry
-  %GroupID.ascast = addrspacecast ptr %GroupID to ptr addrspace(4)
-  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %GroupID) #12
-  %0 = load <3 x i64>, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
-  %1 = extractelement <3 x i64> %0, i64 0
-  store i64 %1, ptr %GroupID, align 8, !tbaa !10
-  %arrayinit.element = getelementptr inbounds i8, ptr %GroupID, i64 8
-  %2 = extractelement <3 x i64> %0, i64 1
-  store i64 %2, ptr %arrayinit.element, align 8, !tbaa !10
-  %arrayinit.element1 = getelementptr inbounds i8, ptr %GroupID, i64 16
-  %3 = extractelement <3 x i64> %0, i64 2
-  store i64 %3, ptr %arrayinit.element1, align 8, !tbaa !10
-  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalLinearId, align 8, !tbaa !10
-  %5 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32
-  %6 = load i64, ptr addrspace(1) getelementptr inbounds (i8, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, i64 8), align 8
-  %mul = mul i64 %5, %6
-  %7 = load i64, ptr addrspace(1) getelementptr inbounds (i8, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, i64 16), align 16
-  %mul2 = mul i64 %mul, %7
-  %conv = trunc i64 %mul2 to i32
-  call spir_func void @__itt_offload_wi_start_stub(ptr addrspace(4) noundef %GroupID.ascast, i64 noundef %4, i32 noundef %conv) #11
-  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %GroupID) #12
-  br label %return
-
-return:                                           ; preds = %if.end, %entry
-  ret void
-}
-
-; Function Attrs: alwaysinline convergent mustprogress norecurse nounwind
-define dso_local spir_func void @__itt_offload_wi_finish_wrapper() #9 !srcloc !14 {
-entry:
-  %GroupID = alloca [3 x i64], align 8
-  %call.i = tail call spir_func signext i8 @__spirv_SpecConstant(i32 noundef -9145239, i8 noundef signext 0) #11
-  %cmp.i.not = icmp eq i8 %call.i, 0
-  br i1 %cmp.i.not, label %return, label %if.end
-
-if.end:                                           ; preds = %entry
-  %GroupID.ascast = addrspacecast ptr %GroupID to ptr addrspace(4)
-  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %GroupID) #12
-  %0 = load <3 x i64>, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
-  %1 = extractelement <3 x i64> %0, i64 0
-  store i64 %1, ptr %GroupID, align 8, !tbaa !10
-  %arrayinit.element = getelementptr inbounds i8, ptr %GroupID, i64 8
-  %2 = extractelement <3 x i64> %0, i64 1
-  store i64 %2, ptr %arrayinit.element, align 8, !tbaa !10
-  %arrayinit.element1 = getelementptr inbounds i8, ptr %GroupID, i64 16
-  %3 = extractelement <3 x i64> %0, i64 2
-  store i64 %3, ptr %arrayinit.element1, align 8, !tbaa !10
-  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalLinearId, align 8, !tbaa !10
-  call spir_func void @__itt_offload_wi_finish_stub(ptr addrspace(4) noundef %GroupID.ascast, i64 noundef %4) #11
-  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %GroupID) #12
-  br label %return
-
-return:                                           ; preds = %if.end, %entry
-  ret void
-}
-
-; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
-define dso_local spir_func void @__itt_offload_wi_start_stub(ptr addrspace(4) noundef %group_id, i64 noundef %wi_id, i32 noundef %wg_size) local_unnamed_addr #10 !srcloc !15 {
-entry:
-  %group_id.addr = alloca ptr addrspace(4), align 8
-  %wi_id.addr = alloca i64, align 8
-  %wg_size.addr = alloca i32, align 4
-  %group_id.addr.ascast = addrspacecast ptr %group_id.addr to ptr addrspace(4)
-  %wi_id.addr.ascast = addrspacecast ptr %wi_id.addr to ptr addrspace(4)
-  %wg_size.addr.ascast = addrspacecast ptr %wg_size.addr to ptr addrspace(4)
-  store ptr addrspace(4) %group_id, ptr addrspace(4) %group_id.addr.ascast, align 8, !tbaa !16
-  store i64 %wi_id, ptr addrspace(4) %wi_id.addr.ascast, align 8, !tbaa !10
-  store i32 %wg_size, ptr addrspace(4) %wg_size.addr.ascast, align 4, !tbaa !18
-  ret void
-}
-
-; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
-define dso_local spir_func void @__itt_offload_wi_finish_stub(ptr addrspace(4) noundef %group_id, i64 noundef %wi_id) local_unnamed_addr #10 !srcloc !20 {
-entry:
-  %group_id.addr = alloca ptr addrspace(4), align 8
-  %wi_id.addr = alloca i64, align 8
-  %group_id.addr.ascast = addrspacecast ptr %group_id.addr to ptr addrspace(4)
-  %wi_id.addr.ascast = addrspacecast ptr %wi_id.addr to ptr addrspace(4)
-  store ptr addrspace(4) %group_id, ptr addrspace(4) %group_id.addr.ascast, align 8, !tbaa !16
-  store i64 %wi_id, ptr addrspace(4) %wi_id.addr.ascast, align 8, !tbaa !10
-  ret void
-}
-
-attributes #0 = { "sycl-unique-id"="_Z12__DeviceType" }
-attributes #1 = { "sycl-unique-id"="_Z29__AsanShadowMemoryGlobalStart" }
-attributes #2 = { "sycl-unique-id"="_Z27__AsanShadowMemoryGlobalEnd" }
-attributes #3 = { "sycl-unique-id"="_Z28__AsanShadowMemoryLocalStart" }
-attributes #4 = { "sycl-unique-id"="_Z26__AsanShadowMemoryLocalEnd" }
-attributes #5 = { "sycl-unique-id"="_Z26__DeviceSanitizerReportMem" }
-attributes #6 = { mustprogress norecurse nounwind sanitize_address uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test.cpp" "sycl-optlevel"="2" "sycl-single-task" "uniform-work-group-size"="true" }
-attributes #7 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
-attributes #8 = { convergent nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-attributes #9 = { alwaysinline convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="/export/users/maosuzha/ics_workspace/syclos/libdevice/itt_compiler_wrappers.cpp" "sycl-optlevel"="2" }
-attributes #10 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="/export/users/maosuzha/ics_workspace/syclos/libdevice/itt_stubs.cpp" "sycl-optlevel"="2" }
-attributes #11 = { convergent nounwind }
-attributes #12 = { nounwind }
-
-!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0}
-!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1}
-!llvm.ident = !{!2, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !2}
-!llvm.module.flags = !{!4, !5, !6}
-!sycl.specialization-constants = !{}
-!sycl.specialization-constants-default-values = !{}
-
-!0 = !{i32 1, i32 2}
-!1 = !{i32 4, i32 100000}
-!2 = !{!"clang version 19.0.0git (https://github.com/intel/llvm.git c4308cc8751c15934d154a9a9d5cac8c31a7743a)"}
-!3 = !{!"clang version 19.0.0git (https://github.com/intel/llvm.git 55c6d2b3751e3b59e9aaf3972f375c33dc0b9d8b)"}
-!4 = !{i32 1, !"wchar_size", i32 4}
-!5 = !{i32 7, !"uwtable", i32 2}
-!6 = !{i32 7, !"frame-pointer", i32 2}
-!7 = !{i32 5141640}
-!8 = !{}
-!9 = !{i32 442}
-!10 = !{!11, !11, i64 0}
-!11 = !{!"long", !12, i64 0}
-!12 = !{!"omnipotent char", !13, i64 0}
-!13 = !{!"Simple C++ TBAA"}
-!14 = !{i32 1030}
-!15 = !{i32 462}
-!16 = !{!17, !17, i64 0}
-!17 = !{!"any pointer", !12, i64 0}
-!18 = !{!19, !19, i64 0}
-!19 = !{!"int", !12, i64 0}
-!20 = !{i32 592}
diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp
index 6d1941c58958b..cfc77534edc16 100644
--- a/llvm/tools/sycl-post-link/sycl-post-link.cpp
+++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp
@@ -339,7 +339,12 @@ std::vector<StringRef> getKernelNamesUsingAssert(const Module &M) {
 }
 
 bool isModuleUsingAsan(const Module &M) {
-  return nullptr != M.getNamedGlobal("__DeviceSanitizerReportMem");
+  NamedMDNode *MD = M.getNamedMetadata("device.sanitizer");
+  if (MD == nullptr)
+    return false;
+  assert(MD->getNumOperands() != 0);
+  auto *MDVal = cast<MDString>(MD->getOperand(0)->getOperand(0));
+  return MDVal->getString() == "asan";
 }
 
 // Gets reqd_work_group_size information for function Func.
diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt
index 7d57dfa01f220..0aaa4d5c762eb 100644
--- a/sycl/plugins/unified_runtime/CMakeLists.txt
+++ b/sycl/plugins/unified_runtime/CMakeLists.txt
@@ -95,13 +95,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit fe9a05e528992cd1db7b05e2857fb17879442e86
-  # Merge: ee2feb22 9222315f
+  # commit 003d4da84fd6eda5240c9b90217f8901559a28a4
+  # Merge: e38e79e6 7a5c1ad1
   # Author: aarongreig <aaron.greig@codeplay.com>
-  # Date:   Tue Apr 16 10:10:10 2024 +0100
-  #    Merge pull request #1507 from nrspruit/fix_p2p_properties_init
-  #    [L0] Fix to p2p properties init for pNext and stype
-  set(UNIFIED_RUNTIME_TAG fe9a05e528992cd1db7b05e2857fb17879442e86)
+  # Date:   Wed Apr 17 17:52:21 2024 +0100
+  #    Merge pull request #1402 from AllanZyne/user-after-free
+  #    [DeviceSanitizer] Checking "sycl::free" related errors
+  set(UNIFIED_RUNTIME_TAG 003d4da84fd6eda5240c9b90217f8901559a28a4)
 
   fetch_adapter_source(level_zero
     ${UNIFIED_RUNTIME_REPO}
@@ -110,46 +110,22 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
 
   fetch_adapter_source(opencl
     ${UNIFIED_RUNTIME_REPO}
-    # commit 0d2a972c71ba4dd5935478c7b7124a372a1eeca0
-    # Merge: ac89abfe 44aef877
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Thu Apr 11 10:24:19 2024 +0100
-    #     Merge pull request #1440 from fabiomestre/fabio/opencl_remove_queued_hack
-    #     [OPENCL] Remove EVENT_STATUS_QUEUED workaround
-    0d2a972c71ba4dd5935478c7b7124a372a1eeca0
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(cuda
     ${UNIFIED_RUNTIME_REPO}
-    # commit 3f5f5688471a23a8d49ad6f47333df92e9f2e5c6
-    # Merge: fe9a05e5 6027c6bc
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Tue Apr 16 13:59:32 2024 +0100
-    #     Merge pull request #1510 from kbenzie/benie/cuda-fix-wrapper-escape
-    #     Fix Coverity wrapper escape issue
-    3f5f5688471a23a8d49ad6f47333df92e9f2e5c6
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(hip
     ${UNIFIED_RUNTIME_REPO}
-    # commit 55cf3ba8139f8ffea9cf5cf42ea1f79c049b3520
-    # Merge: 15233fd2 35a6899d
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Mon Apr 15 09:50:52 2024 +0100
-    #     Merge pull request #1437 from hdelan/event-wait-with-good-context
-    #     [HIP] Use context to get active device
-    55cf3ba8139f8ffea9cf5cf42ea1f79c049b3520
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(native_cpu
     ${UNIFIED_RUNTIME_REPO}
-    # commit 15233fd2521f9e9b35e3a24037be99ceef334a8e
-    # Merge: 68e525a4 a04b062e
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Fri Apr 12 15:46:49 2024 +0100
-    #     Merge pull request #1489 from konradkusiak97/nativeCPUqueueFill
-    #     [NATIVECPU] Extended usm fill to bigger patterns than 1 byte
-    15233fd2521f9e9b35e3a24037be99ceef334a8e
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
diff --git a/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-host.cpp b/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-host.cpp
new file mode 100644
index 0000000000000..08307de42f075
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-host.cpp
@@ -0,0 +1,15 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck %s
+#include <sycl/sycl.hpp>
+
+constexpr size_t N = 64;
+
+int main() {
+  sycl::queue Q;
+  auto *data = new int[N];
+  sycl::free(data, Q);
+  return 0;
+}
+// CHECK: ERROR: DeviceSanitizer: bad-free on address [[ADDR:0x.*]]
+// CHECK: [[ADDR]] may be allocated on Host Memory
diff --git a/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-minus1.cpp b/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-minus1.cpp
new file mode 100644
index 0000000000000..58f2c9e781c7e
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-minus1.cpp
@@ -0,0 +1,29 @@
+// REQUIRES: linux, cpu
+// 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
+// RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O0 -g -o %t
+// RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-SHARED %s
+#include <sycl/sycl.hpp>
+
+constexpr size_t N = 64;
+
+int main() {
+  sycl::queue Q;
+
+#if defined(MALLOC_HOST)
+  auto *data = sycl::malloc_host<int>(N, Q);
+#elif defined(MALLOC_SHARED)
+  auto *data = sycl::malloc_shared<int>(N, Q);
+#else
+  auto *data = sycl::malloc_device<int>(N, Q);
+#endif
+
+  sycl::free(data - 1, Q);
+  return 0;
+}
+// CHECK: ERROR: DeviceSanitizer: bad-free on address [[ADDR:0x.*]]
+// CHECK-HOST:   [[ADDR]] is located inside of Host USM region {{\[0x.*, 0x.*\)}}
+// CHECK-SHARED: [[ADDR]] is located inside of Shared USM region {{\[0x.*, 0x.*\)}}
+// CHECK-DEVICE: [[ADDR]] is located inside of Device USM region {{\[0x.*, 0x.*\)}}
diff --git a/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-plus1.cpp b/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-plus1.cpp
new file mode 100644
index 0000000000000..dab9302cca85d
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/bad-free/bad-free-plus1.cpp
@@ -0,0 +1,29 @@
+// REQUIRES: linux, cpu
+// 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
+// RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O0 -g -o %t
+// RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-SHARED %s
+#include <sycl/sycl.hpp>
+
+constexpr size_t N = 64;
+
+int main() {
+  sycl::queue Q;
+
+#if defined(MALLOC_HOST)
+  auto *data = sycl::malloc_host<int>(N, Q);
+#elif defined(MALLOC_SHARED)
+  auto *data = sycl::malloc_shared<int>(N, Q);
+#else
+  auto *data = sycl::malloc_device<int>(N, Q);
+#endif
+
+  sycl::free(data + 1, Q);
+  // CHECK: ERROR: DeviceSanitizer: bad-free on address [[ADDR:0x.*]]
+  // CHECK-HOST:   [[ADDR]] is located inside of Host USM region {{\[0x.*, 0x.*\)}}
+  // CHECK-SHARED: [[ADDR]] is located inside of Shared USM region {{\[0x.*, 0x.*\)}}
+  // CHECK-DEVICE: [[ADDR]] is located inside of Device USM region {{\[0x.*, 0x.*\)}}
+  return 0;
+}
diff --git a/sycl/test-e2e/AddressSanitizer/common/demangle-kernel-name.cpp b/sycl/test-e2e/AddressSanitizer/common/demangle-kernel-name.cpp
index bfb3efa89d81c..4b97c8d7f3672 100644
--- a/sycl/test-e2e/AddressSanitizer/common/demangle-kernel-name.cpp
+++ b/sycl/test-e2e/AddressSanitizer/common/demangle-kernel-name.cpp
@@ -1,5 +1,5 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t
+// 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>
 
diff --git a/sycl/test-e2e/AddressSanitizer/common/kernel-debug.cpp b/sycl/test-e2e/AddressSanitizer/common/kernel-debug.cpp
new file mode 100644
index 0000000000000..1527a6efe3291
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/common/kernel-debug.cpp
@@ -0,0 +1,20 @@
+// REQUIRES: linux, cpu
+// 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
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+  int *array = sycl::malloc_device<int>(1, Q);
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { *array = 0; });
+  });
+  Q.wait();
+  // CHECK-DEBUG: [kernel]
+  // CHECK-NOT: [kernel]
+
+  std::cout << "PASS" << std::endl;
+  return 0;
+}
diff --git a/sycl/test-e2e/AddressSanitizer/double-free/double-free.cpp b/sycl/test-e2e/AddressSanitizer/double-free/double-free.cpp
new file mode 100644
index 0000000000000..91b2aa0213647
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/double-free/double-free.cpp
@@ -0,0 +1,33 @@
+// REQUIRES: linux, cpu
+// 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
+// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -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-SHARED %s
+#include <sycl/sycl.hpp>
+
+constexpr size_t N = 64;
+
+int main() {
+  sycl::queue Q;
+
+#if defined(MALLOC_HOST)
+  auto *data = sycl::malloc_host<int>(N, Q);
+#elif defined(MALLOC_SHARED)
+  auto *data = sycl::malloc_shared<int>(N, Q);
+#else
+  auto *data = sycl::malloc_device<int>(N, Q);
+#endif
+
+  sycl::free(data, Q);
+  sycl::free(data, Q);
+
+  return 0;
+}
+// CHECK: ERROR: DeviceSanitizer: double-free on address [[ADDR:0x.*]]
+// CHECK-HOST:   [[ADDR]] is located inside of Host USM region {{\[0x.*, 0x.*\)}}
+// CHECK-SHARED: [[ADDR]] is located inside of Shared USM region {{\[0x.*, 0x.*\)}}
+// CHECK-DEVICE: [[ADDR]] is located inside of Device USM region {{\[0x.*, 0x.*\)}}
+// CHECK: freed here
+// CHECK: previously allocated here
diff --git a/sycl/test-e2e/AddressSanitizer/lit.local.cfg b/sycl/test-e2e/AddressSanitizer/lit.local.cfg
new file mode 100644
index 0000000000000..63fd4ac5bd18a
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/lit.local.cfg
@@ -0,0 +1,7 @@
+config.substitutions.append(
+    ("%device_asan_flags", "-Xarch_device -fsanitize=address")
+)
+
+config.substitutions.append(
+    ("%force_device_asan_rt", "env SYCL_PREFER_UR=1 UR_ENABLE_LAYERS=UR_LAYER_ASAN")
+)
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp
index 143529c9f9891..dde453a659a12 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp
@@ -1,6 +1,10 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t.out
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+// 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>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp
index 4af96a37bbc1b..4836d367bc14d 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp
@@ -1,6 +1,10 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t.out
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+// 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>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp
index bdeef5a65fd61..088408c8820e8 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp
@@ -1,6 +1,10 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t.out
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+// 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>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp
index fb602705d239f..e1d46dee3c10c 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp
@@ -1,8 +1,8 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -O2 -g -DUSER_CODE_1 -c -o %t1.o
-// RUN: %{build} %device_sanitizer_flags -O2 -g -DUSER_CODE_2 -c -o %t2.o
-// RUN: %clangxx -fsycl %device_sanitizer_flags -O2 -g %t1.o %t2.o -o %t.out
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+// 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
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp
index 898442449b243..7ab4eea9b9123 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp
@@ -1,12 +1,14 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+// 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
+// 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 -O2 -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_HOST -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp
index d72d7677433dc..ac8d17fd5528d 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp
@@ -1,12 +1,14 @@
 // REQUIRES: linux, cpu, aspect-fp64
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+// 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
+// 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 -O2 -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_HOST -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp
index 470d17758748e..bd9dbe7fca999 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp
@@ -1,12 +1,14 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+// 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
+// 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 -O2 -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_HOST -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp
index 9b72df1d9cb48..139de679ae6c5 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp
@@ -1,12 +1,14 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+// 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
+// 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 -O2 -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_HOST -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp
index 87e7b495fff28..a0ae55b6674e1 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp
@@ -1,12 +1,14 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
-// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+// 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
+// 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 -O2 -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_HOST -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp
index b781b4840a751..fc3afac179ee6 100644
--- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp
+++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp
@@ -1,6 +1,6 @@
 // REQUIRES: linux, cpu
-// RUN: %{build} %device_sanitizer_flags -g -o %t.out
-// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t.out 2>&1 | FileCheck %s
 
 #include <sycl/detail/core.hpp>
 
diff --git a/sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-free.cpp b/sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-free.cpp
new file mode 100644
index 0000000000000..6fdbb6a782f3c
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-free.cpp
@@ -0,0 +1,50 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:5 UR_LOG_SANITIZER=level:info %{run} %t 2>&1 | FileCheck %s
+#include <sycl/sycl.hpp>
+
+/// Quarantine Cache Test
+///
+/// The "sycl::free"d buffer are not freed immediately, but enqueued into
+/// quarantine cache.
+/// The maximum size of quarantine cache (per device) is configured by
+/// "quarantine_size_mb" on env "UR_LAYER_ASAN_OPTIONS".
+/// If the total size of enqueued buffers is larger than "quarantine_size_mb",
+/// then the enqueued buffers will be freed by FIFO.
+///
+/// In this test, the maximum size of quarantine cache is 5MB (5242880 bytes).
+
+constexpr size_t N = 1024 * 1024;
+
+int main() {
+  sycl::queue Q;
+  auto *array =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672 <= 5242880
+  // 1. allocated size: {currently the size of all allocated memory} <= {maximum
+  // size of quarantine cache}"
+  // 2. 1052672 = 1024*1024 + 4096, 4096 is the size of red zone
+  // CHECK: Alloc={{\[}}[[ADDR1:0x[0-9a-f]+]]
+  sycl::free(array, Q);
+
+  auto *temp =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672*2 <= 5242880
+  // CHECK: Alloc={{\[}}[[ADDR2:0x[0-9a-f]+]]
+  sycl::free(temp, Q);
+
+  temp =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672*3 <= 5242880
+  // CHECK: Alloc={{\[}}[[ADDR3:0x[0-9a-f]+]]
+  sycl::free(temp, Q);
+
+  temp =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672*4 <= 5242880
+  // CHECK: Alloc={{\[}}[[ADDR4:0x[0-9a-f]+]]
+  sycl::free(temp, Q);
+
+  temp = sycl::malloc_device<char>(N, Q); // allocated size: 1052672*5 > 5242880
+  // CHECK: Alloc={{\[}}[[ADDR5:0x[0-9a-f]+]]
+  sycl::free(temp, Q);
+  // CHECK: Quarantine Free: [[ADDR1]]
+
+  return 0;
+}
diff --git a/sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-no-free.cpp b/sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-no-free.cpp
new file mode 100644
index 0000000000000..ad5a8e36119cc
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-no-free.cpp
@@ -0,0 +1,52 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:5 UR_LOG_SANITIZER=level:info %{run} not %t 2>&1 | FileCheck %s
+#include <sycl/sycl.hpp>
+
+/// Quarantine Cache Test
+///
+/// The "sycl::free"d buffer are not freed immediately, but enqueued into
+/// quarantine cache.
+/// The maximum size of quarantine cache (per device) is configured by
+/// "quarantine_size_mb" on env "UR_LAYER_ASAN_OPTIONS".
+/// If the total size of enqueued buffers is larger than "quarantine_size_mb",
+/// then the enqueued buffers will be freed by FIFO.
+///
+/// In this test, the maximum size of quarantine cache is 5MB (5242880 bytes).
+
+constexpr size_t N = 1024 * 1024;
+
+int main() {
+  sycl::queue Q;
+  auto *array =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672 <= 5242880
+  // 1. allocated size: {currently the size of all allocated memory} <= {maximum
+  // size of quarantine cache}"
+  // 2. 1052672 = 1024*1024 + 4096, 4096 is the size of red zone
+  sycl::free(array, Q);
+
+  auto *temp =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672*2 <= 5242880
+  sycl::free(temp, Q);
+  temp =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672*3 <= 5242880
+  sycl::free(temp, Q);
+  temp =
+      sycl::malloc_device<char>(N, Q); // allocated size: 1052672*4 <= 5242880
+  sycl::free(temp, Q);
+  // Make sure the first allocated buffer is not freed
+  // CHECK-NOT: <SANITIZER>[INFO]: Quarantine Free
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { array[0] = 0; });
+  });
+  Q.wait();
+  // CHECK: ERROR: DeviceSanitizer: use-after-free on address [[ADDR:0x.*]]
+  // CHECK: WRITE of size 1 at kernel <{{.*MyKernel}}>
+  // CHECK:   #0 {{.*}} {{.*quarantine-no-free.cpp}}:[[@LINE-5]]
+  // CHECK: [[ADDR]] is located inside of Device USM region [{{0x.*}}, {{0x.*}})
+  // CHECK: allocated here:
+  // CHECK: released here:
+
+  return 0;
+}
diff --git a/sycl/test-e2e/AddressSanitizer/use-after-free/use-after-free.cpp b/sycl/test-e2e/AddressSanitizer/use-after-free/use-after-free.cpp
new file mode 100644
index 0000000000000..b14b092722b9a
--- /dev/null
+++ b/sycl/test-e2e/AddressSanitizer/use-after-free/use-after-free.cpp
@@ -0,0 +1,27 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck %s
+#include <sycl/sycl.hpp>
+
+constexpr size_t N = 1024;
+
+int main() {
+  sycl::queue Q;
+  auto *array = sycl::malloc_device<char>(N, Q);
+  sycl::free(array, Q);
+
+  Q.submit([&](sycl::handler &h) {
+    h.parallel_for<class MyKernel>(
+        sycl::nd_range<1>(N, 1),
+        [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
+  });
+  Q.wait();
+  // CHECK: ERROR: DeviceSanitizer: use-after-free on address [[ADDR:0x.*]]
+  // CHECK: READ of size 1 at kernel <{{.*MyKernel}}>
+  // CHECK:   #0 {{.*}} {{.*use-after-free.cpp:}}[[@LINE-5]]
+  // CHECK: [[ADDR]] is located inside of Device USM region [{{0x.*}}, {{0x.*}})
+  // CHECK: allocated here:
+  // CHECK: released here:
+
+  return 0;
+}
diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py
index 0da52813e1fa1..4637f2743f510 100644
--- a/sycl/test-e2e/lit.cfg.py
+++ b/sycl/test-e2e/lit.cfg.py
@@ -701,7 +701,3 @@
     lit_config.maxIndividualTestTime = 600
 except ImportError:
     pass
-
-config.substitutions.append(
-    ("%device_sanitizer_flags", "-Xsycl-target-frontend -fsanitize=address")
-)