From 4723efc481cc18160cfa2f76d89378a84c43df64 Mon Sep 17 00:00:00 2001 From: Yang Zhao Date: Fri, 19 Apr 2024 17:30:23 +0800 Subject: [PATCH] [SYCL][DeviceSanitizer] Checking "sycl::free" related errors (#12882) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit UR: https://github.com/oneapi-src/unified-runtime/pull/1402 This PR added supports for checking the following types of error in "UR_LAYER_ASAN": - bad-free: the memory address to be freed is not allocated by UR - bad-context: the memory address to be freed uses a wrong "context" - double-free: the memory address to be freed is already freed - use-after-free: the freed memory is used in kernel I added the environment variable "UR_LAYER_ASAN_OPTIONS" to have additional control over "UR_LAYER_ASAN", which is similar to "ASAN_OPTIONS" in [ASan](https://github.com/google/sanitizers/wiki/AddressSanitizerFlags). Currently, it supports: - "quarantine_size_mb" (default = 0) - Size (in MB) of quarantine per device. The pointers passed to urUSMFree are not freed immediately, but saved into QuarantineCache (per device cache), and when the cached chunk size (only counts the size of USM buffer, not shadow memory) is more than "quarantine_size_mb", the first enqueued chunk will be freed (aka., FIFO). Lower value may reduce memory usage but increase the chance of false negatives - This option must be enabled for checking "double-free" and "use-after-free" - "debug" (default = 0) - Print extra debug messages in kernel ("__AsanDebug” in “libdevice/sanitizer_utils.cpp”), which is helpful for DeviceSanitizer developers. For example, to enable "use-after-free" with 5MB quarantine cache and debug message in kernel, you need to ```bash UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:5;debug:1" ./sycl_app ``` --------- Co-authored-by: Maosu Zhao Co-authored-by: Aaron Greig --- libdevice/include/asan_libdevice.hpp | 155 +++++++ libdevice/include/device-sanitizer-report.hpp | 56 --- libdevice/include/sanitizer_device_utils.hpp | 2 +- libdevice/sanitizer_utils.cpp | 384 +++++++++++------- .../Instrumentation/AddressSanitizer.cpp | 30 +- .../sycl-post-link/device-sanitizer/asan.ll | 76 ++++ .../sycl-post-link/sycl-sanitize/asan.ll | 175 -------- llvm/tools/sycl-post-link/sycl-post-link.cpp | 7 +- sycl/plugins/unified_runtime/CMakeLists.txt | 44 +- .../bad-free/bad-free-host.cpp | 15 + .../bad-free/bad-free-minus1.cpp | 29 ++ .../bad-free/bad-free-plus1.cpp | 29 ++ .../common/demangle-kernel-name.cpp | 2 +- .../AddressSanitizer/common/kernel-debug.cpp | 20 + .../double-free/double-free.cpp | 33 ++ sycl/test-e2e/AddressSanitizer/lit.local.cfg | 7 + .../DeviceGlobal/device_global.cpp | 8 +- .../device_global_image_scope.cpp | 8 +- .../device_global_image_scope_unaligned.cpp | 8 +- .../DeviceGlobal/multi_device_images.cpp | 8 +- .../out-of-bounds/USM/parallel_for_char.cpp | 18 +- .../out-of-bounds/USM/parallel_for_double.cpp | 18 +- .../out-of-bounds/USM/parallel_for_func.cpp | 18 +- .../out-of-bounds/USM/parallel_for_int.cpp | 18 +- .../out-of-bounds/USM/parallel_for_short.cpp | 18 +- .../out-of-bounds/local/local-overflow-1.cpp | 4 +- .../use-after-free/quarantine-free.cpp | 50 +++ .../use-after-free/quarantine-no-free.cpp | 52 +++ .../use-after-free/use-after-free.cpp | 27 ++ sycl/test-e2e/lit.cfg.py | 4 - 30 files changed, 848 insertions(+), 475 deletions(-) create mode 100644 libdevice/include/asan_libdevice.hpp delete mode 100644 libdevice/include/device-sanitizer-report.hpp create mode 100644 llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll delete mode 100644 llvm/test/tools/sycl-post-link/sycl-sanitize/asan.ll create mode 100644 sycl/test-e2e/AddressSanitizer/bad-free/bad-free-host.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/bad-free/bad-free-minus1.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/bad-free/bad-free-plus1.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/common/kernel-debug.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/double-free/double-free.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/lit.local.cfg create mode 100644 sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-free.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-no-free.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/use-after-free/use-after-free.cpp 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 + +// 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 - -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 #include @@ -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 __AsanShadowMemoryGlobalStart; DeviceGlobal __AsanShadowMemoryGlobalEnd; DeviceGlobal __AsanShadowMemoryLocalStart; DeviceGlobal __AsanShadowMemoryLocalEnd; - -DeviceGlobal __DeviceSanitizerReportMem; - DeviceGlobal __DeviceType; +DeviceGlobal __AsanDebug; +DeviceGlobal __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(Addr->getType()->getScalarType()); // Private address space: skip kernel arguments if (PtrTy->getPointerAddressSpace() == 0) { - return Func->getCallingConv() == CallingConv::SPIR_KERNEL && - isa(Addr); + // FIXME: Currently we don't suppport all private variables + // return Func->getCallingConv() == CallingConv::SPIR_KERNEL && + // isa(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 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(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 - # 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) - # 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) - # 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) - # 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) - # 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 + +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 + +constexpr size_t N = 64; + +int main() { + sycl::queue Q; + +#if defined(MALLOC_HOST) + auto *data = sycl::malloc_host(N, Q); +#elif defined(MALLOC_SHARED) + auto *data = sycl::malloc_shared(N, Q); +#else + auto *data = sycl::malloc_device(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 + +constexpr size_t N = 64; + +int main() { + sycl::queue Q; + +#if defined(MALLOC_HOST) + auto *data = sycl::malloc_host(N, Q); +#elif defined(MALLOC_SHARED) + auto *data = sycl::malloc_shared(N, Q); +#else + auto *data = sycl::malloc_device(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 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 + +int main() { + sycl::queue Q; + int *array = sycl::malloc_device(1, Q); + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { *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 + +constexpr size_t N = 64; + +int main() { + sycl::queue Q; + +#if defined(MALLOC_HOST) + auto *data = sycl::malloc_host(N, Q); +#elif defined(MALLOC_SHARED) + auto *data = sycl::malloc_shared(N, Q); +#else + auto *data = sycl::malloc_device(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 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 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 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 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 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 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 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 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 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 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 + +/// 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(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(N, Q); // allocated size: 1052672*2 <= 5242880 + // CHECK: Alloc={{\[}}[[ADDR2:0x[0-9a-f]+]] + sycl::free(temp, Q); + + temp = + sycl::malloc_device(N, Q); // allocated size: 1052672*3 <= 5242880 + // CHECK: Alloc={{\[}}[[ADDR3:0x[0-9a-f]+]] + sycl::free(temp, Q); + + temp = + sycl::malloc_device(N, Q); // allocated size: 1052672*4 <= 5242880 + // CHECK: Alloc={{\[}}[[ADDR4:0x[0-9a-f]+]] + sycl::free(temp, Q); + + temp = sycl::malloc_device(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 + +/// 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(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(N, Q); // allocated size: 1052672*2 <= 5242880 + sycl::free(temp, Q); + temp = + sycl::malloc_device(N, Q); // allocated size: 1052672*3 <= 5242880 + sycl::free(temp, Q); + temp = + sycl::malloc_device(N, Q); // allocated size: 1052672*4 <= 5242880 + sycl::free(temp, Q); + // Make sure the first allocated buffer is not freed + // CHECK-NOT: [INFO]: Quarantine Free + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { 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 + +constexpr size_t N = 1024; + +int main() { + sycl::queue Q; + auto *array = sycl::malloc_device(N, Q); + sycl::free(array, Q); + + Q.submit([&](sycl::handler &h) { + h.parallel_for( + 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") -)