[DeviceASAN] Enhance out-of-bound by checking shadow bounds on global memory

libdevice/sanitizer/asan_rtl.cpp

@@ -40,6 +40,9 @@ static const __SYCL_CONSTANT__ char __asan_print_shadow_value2[] =
 static __SYCL_CONSTANT__ const char __generic_to[] =
     "[kernel] %p(4) - %p(%d)\n";
 
+static __SYCL_CONSTANT__ const char __asan_print_shadow_bound[] =
+    "[kernel] addr: %p, shadow: %p, lower: %p, uppper: %p\n";
+
 #define ASAN_REPORT_NONE 0
 #define ASAN_REPORT_START 1
 #define ASAN_REPORT_FINISH 2
@@ -67,6 +70,8 @@ struct DebugInfo {
 
 void ReportUnknownDevice(const DebugInfo *debug);
 void PrintShadowMemory(uptr addr, uptr shadow_address, uint32_t as);
+void SaveReport(ErrorType error_type, MemoryType memory_type, bool is_recover,
+                const DebugInfo *debug);
 
 __SYCL_GLOBAL__ void *ToGlobal(void *ptr) {
   return __spirv_GenericCastToPtrExplicit_ToGlobal(ptr, 5);
@@ -115,11 +120,22 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
           launch_info->GlobalShadowOffset + (addr >> ASAN_SHADOW_SCALE);
     }
 
+    if (shadow_ptr < launch_info->GlobalShadowLowerBound ||
+        shadow_ptr > launch_info->GlobalShadowUpperBound) {
+      ASAN_DEBUG(__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
+                                    launch_info->GlobalShadowLowerBound,
+                                    launch_info->GlobalShadowUpperBound));
+      SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
+      return 0;
+    }
+
     ASAN_DEBUG(
         const auto shadow_offset_end = launch_info->GlobalShadowOffsetEnd;
         if (shadow_ptr > shadow_offset_end) {
           __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
                              (uptr)launch_info->GlobalShadowOffset);
+          SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::GLOBAL, false,
+                     debug);
           return 0;
         });
 
@@ -141,6 +157,8 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
                if (shadow_ptr > shadow_offset_end) {
                  __spirv_ocl_printf(__local_shadow_out_of_bound, addr,
                                     shadow_ptr, wid, (uptr)shadow_offset);
+                 SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::LOCAL,
+                            false, debug);
                  return 0;
                });
     return shadow_ptr;
@@ -167,8 +185,10 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
     if (shadow_ptr > shadow_offset_end) {
       __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, sid,
                          private_base);
+      SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::PRIVATE, false,
+                 debug);
       return 0;
-    };
+    }
 
     return shadow_ptr;
   }
@@ -193,11 +213,22 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
                    ((addr & 0x7FFFFFFFFFFF) >> ASAN_SHADOW_SCALE);
     }
 
+    if (shadow_ptr < launch_info->GlobalShadowLowerBound ||
+        shadow_ptr > launch_info->GlobalShadowUpperBound) {
+      ASAN_DEBUG(__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
+                                    launch_info->GlobalShadowLowerBound,
+                                    launch_info->GlobalShadowUpperBound));
+      SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
+      return 0;
+    }
+
     ASAN_DEBUG(
         const auto shadow_offset_end = launch_info->GlobalShadowOffsetEnd;
         if (shadow_ptr > shadow_offset_end) {
           __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
                              (uptr)launch_info->GlobalShadowOffset);
+          SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::GLOBAL, false,
+                     debug);
           return 0;
         });
     return shadow_ptr;
@@ -218,6 +249,8 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
                if (shadow_ptr > shadow_offset_end) {
                  __spirv_ocl_printf(__local_shadow_out_of_bound, addr,
                                     shadow_ptr, wid, (uptr)shadow_offset);
+                 SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::LOCAL,
+                            false, debug);
                  return 0;
                });
     return shadow_ptr;
@@ -244,6 +277,8 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
     if (shadow_ptr > shadow_offset_end) {
       __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, sid,
                          private_base);
+      SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::PRIVATE, false,
+                 debug);
       return 0;
     };
 

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

@@ -0,0 +1,25 @@
+// REQUIRES: linux, gpu && level_zero
+// RUN: %{build} %device_asan_flags -O0 -g -o %t1.out
+// RUN: %{run} not %t1.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t3.out
+// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+void out_of_bounds_function() { *(int *)0xdeadbeef = 42; }
+// CHECK: out-of-bounds-access
+// CHECK-SAME: 0xdeadbeef
+// CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}}
+// CHECK: {{.*arbitary_access.cpp}}:[[@LINE-4]]
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { out_of_bounds_function(); });
+  });
+  Q.wait();
+
+  return 0;
+}

unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp

@@ -267,11 +267,23 @@ ur_result_t AsanInterceptor::preLaunchKernel(ur_kernel_handle_t Kernel,
 
   ur_queue_handle_t InternalQueue = ContextInfo->getInternalQueue(Device);
 
+  // To get right shadow boundary, shadow memory should be updated before
+  // prepareLaunch
+  {
+    // Force to allocate membuffer before prepareLaunch
+    auto &KernelInfo = getOrCreateKernelInfo(Kernel);
+    std::shared_lock<ur_shared_mutex> Guard(KernelInfo.Mutex);
+    for (const auto &[ArgIndex, MemBuffer] : KernelInfo.BufferArgs) {
+      char *ArgPointer = nullptr;
+      UR_CALL(MemBuffer->getHandle(DeviceInfo->Handle, ArgPointer));
+      (void)ArgPointer;
+    }
+  }
+  UR_CALL(updateShadowMemory(ContextInfo, DeviceInfo, InternalQueue));
+
   UR_CALL(prepareLaunch(ContextInfo, DeviceInfo, InternalQueue, Kernel,
                         LaunchInfo));
 
-  UR_CALL(updateShadowMemory(ContextInfo, DeviceInfo, InternalQueue));
-
   UR_CALL(getContext()->urDdiTable.Queue.pfnFinish(InternalQueue));
 
   return UR_RESULT_SUCCESS;
@@ -825,6 +837,10 @@ ur_result_t AsanInterceptor::prepareLaunch(
   // Prepare asan runtime data
   LaunchInfo.Data.Host.GlobalShadowOffset = DeviceInfo->Shadow->ShadowBegin;
   LaunchInfo.Data.Host.GlobalShadowOffsetEnd = DeviceInfo->Shadow->ShadowEnd;
+  LaunchInfo.Data.Host.GlobalShadowLowerBound =
+      DeviceInfo->Shadow->ShadowLowerBound;
+  LaunchInfo.Data.Host.GlobalShadowUpperBound =
+      DeviceInfo->Shadow->ShadowUpperBound;
   LaunchInfo.Data.Host.Debug = getContext()->Options.Debug ? 1 : 0;
 
   // Write shadow memory offset for local memory
@@ -884,16 +900,20 @@ ur_result_t AsanInterceptor::prepareLaunch(
   // sync asan runtime data to device side
   UR_CALL(LaunchInfo.Data.syncToDevice(Queue));
 
-  UR_LOG_L(getContext()->logger, INFO,
-           "LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateBase={}, "
-           "PrivateShadow={}, LocalArgs={}, NumLocalArgs={}, Debug={})",
-           (void *)LaunchInfo.Data.getDevicePtr(),
-           (void *)LaunchInfo.Data.Host.GlobalShadowOffset,
-           (void *)LaunchInfo.Data.Host.LocalShadowOffset,
-           (void *)LaunchInfo.Data.Host.PrivateBase,
-           (void *)LaunchInfo.Data.Host.PrivateShadowOffset,
-           (void *)LaunchInfo.Data.Host.LocalArgs,
-           LaunchInfo.Data.Host.NumLocalArgs, LaunchInfo.Data.Host.Debug);
+  UR_LOG_L(
+      getContext()->logger, INFO,
+      "LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateBase={}, "
+      "PrivateShadow={}, GlobalShadowLowerBound={}, GlobalShadowUpperBound={}, "
+      "LocalArgs={}, NumLocalArgs={}, Debug={})",
+      (void *)LaunchInfo.Data.getDevicePtr(),
+      (void *)LaunchInfo.Data.Host.GlobalShadowOffset,
+      (void *)LaunchInfo.Data.Host.LocalShadowOffset,
+      (void *)LaunchInfo.Data.Host.PrivateBase,
+      (void *)LaunchInfo.Data.Host.PrivateShadowOffset,
+      (void *)LaunchInfo.Data.Host.GlobalShadowLowerBound,
+      (void *)LaunchInfo.Data.Host.GlobalShadowUpperBound,
+      (void *)LaunchInfo.Data.Host.LocalArgs, LaunchInfo.Data.Host.NumLocalArgs,
+      LaunchInfo.Data.Host.Debug);
 
   return UR_RESULT_SUCCESS;
 }

unified-runtime/source/loader/layers/sanitizer/asan/asan_libdevice.hpp

@@ -59,6 +59,9 @@ struct AsanRuntimeData {
   uintptr_t GlobalShadowOffset = 0;
   uintptr_t GlobalShadowOffsetEnd = 0;
 
+  uintptr_t GlobalShadowLowerBound = 0;
+  uintptr_t GlobalShadowUpperBound = 0;
+
   uintptr_t *PrivateBase = nullptr;
   uintptr_t PrivateShadowOffset = 0;
   uintptr_t PrivateShadowOffsetEnd = 0;

unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.cpp

@@ -226,6 +226,11 @@ ur_result_t ShadowMemoryGPU::EnqueuePoisonShadow(ur_queue_handle_t Queue,
         VirtualMemMaps[MappedPtr] = PhysicalMem;
       }
     }
+
+    ShadowLowerBound =
+        std::min(ShadowLowerBound, RoundDownTo(ShadowBegin, PageSize));
+    ShadowUpperBound =
+        std::max(ShadowUpperBound, RoundUpTo(ShadowEnd, PageSize));
   }
 
   auto URes = EnqueueUSMSet(Queue, (void *)ShadowBegin, Value,

unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.hpp

@@ -64,8 +64,10 @@ struct ShadowMemory {
   ur_device_handle_t Device{};
 
   uptr ShadowBegin = 0;
-
   uptr ShadowEnd = 0;
+
+  uptr ShadowLowerBound = 0xffff'ffff'ffff'ffff;
+  uptr ShadowUpperBound = 0;
 };
 
 struct ShadowMemoryCPU final : public ShadowMemory {

unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp

@@ -68,16 +68,21 @@ inline const char *ToString(ErrorType ErrorType) {
   }
 }
 
-enum class MemoryType : int32_t {
-  UNKNOWN,
-  USM_DEVICE,
-  USM_HOST,
-  USM_SHARED,
-  LOCAL,
-  PRIVATE,
-  MEM_BUFFER,
-  DEVICE_GLOBAL,
+// clang-format off
+enum MemoryType : uint32_t {
+  UNKNOWN       = 0x000000'00,
+  GLOBAL        = 0x000001'00,
+  USM_DEVICE    = 0x000001'01,
+  USM_HOST      = 0x000001'02,
+  USM_SHARED    = 0x000001'03,
+  MEM_BUFFER    = 0x000001'04,
+  DEVICE_GLOBAL = 0x000001'05,
+  LOCAL         = 0x000002'00,
+  PRIVATE       = 0x000004'00,
+  CONSTANT      = 0x000008'00,
+  GENERIC       = 0x000010'00,
 };
+// clang-format on
 
 inline const char *ToString(MemoryType MemoryType) {
   switch (MemoryType) {