[SYCL][DeviceSanitizer] Checking "sycl::free" related errors

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";
+ }
+}

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 };