From fd5126c67d2376d1049e87f1bffe893e4c1eb1b5 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 15 Sep 2025 04:39:46 +0200 Subject: [PATCH 1/9] init --- libdevice/sanitizer/asan_rtl.cpp | 37 ++++++++++++++++++- .../out-of-bounds/USM/arbitary_access.cpp | 27 ++++++++++++++ .../layers/sanitizer/asan/asan_libdevice.hpp | 3 ++ .../layers/sanitizer/asan/asan_shadow.cpp | 5 +++ .../layers/sanitizer/asan/asan_shadow.hpp | 4 +- .../sanitizer_common/sanitizer_libdevice.hpp | 19 ++++++---- 6 files changed, 85 insertions(+), 10 deletions(-) create mode 100644 sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index 076cd61df578f..d80c7f8ab284e 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/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) { + __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) { + __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; }; diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp new file mode 100644 index 0000000000000..c8c5ed786e1f8 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp @@ -0,0 +1,27 @@ +// REQUIRES: linux, cpu || (gpu && level_zero) +// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t1.out +// RUN: %{run} not %t1.out 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t2.out +// RUN: %{run} not %t2.out 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O2 -g -o %t3.out +// RUN: %{run} not %t3.out 2>&1 | FileCheck %s + +#include +#include + +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.cpp}}:[[@LINE-4]] + +int main() { + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { out_of_bounds_function(); }); + }); + Q.wait(); + + return 0; +} diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_libdevice.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_libdevice.hpp index c3e05479915ca..308be288f1d1a 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_libdevice.hpp +++ b/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; diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.cpp index 209752d718458..b784eb228a975 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.cpp +++ b/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, diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.hpp index 7ab897dc280c7..e504043ae3965 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.hpp +++ b/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 { diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp index b0f447711cb5f..e027f10e9c1e1 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp @@ -68,16 +68,19 @@ inline const char *ToString(ErrorType ErrorType) { } } +// clang-format off enum class MemoryType : int32_t { - UNKNOWN, - USM_DEVICE, - USM_HOST, - USM_SHARED, - LOCAL, - PRIVATE, - MEM_BUFFER, - DEVICE_GLOBAL, + 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, }; +// clang-format on inline const char *ToString(MemoryType MemoryType) { switch (MemoryType) { From 199451d195bc9f8d2557d2f6f3c2d14f4ad2f924 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 15 Sep 2025 07:49:34 +0200 Subject: [PATCH 2/9] wip --- libdevice/sanitizer/asan_rtl.cpp | 12 ++--- .../out-of-bounds/USM/arbitary_access.cpp | 10 ++--- .../sanitizer/asan/asan_interceptor.cpp | 44 ++++++++++++++----- .../sanitizer_common/sanitizer_libdevice.hpp | 4 +- 4 files changed, 45 insertions(+), 25 deletions(-) diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index d80c7f8ab284e..9f7c1a1bbd44d 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/libdevice/sanitizer/asan_rtl.cpp @@ -122,9 +122,9 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as, if (shadow_ptr < launch_info->GlobalShadowLowerBound || shadow_ptr > launch_info->GlobalShadowUpperBound) { - __spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr, - launch_info->GlobalShadowLowerBound, - 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; } @@ -215,9 +215,9 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as, if (shadow_ptr < launch_info->GlobalShadowLowerBound || shadow_ptr > launch_info->GlobalShadowUpperBound) { - __spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr, - launch_info->GlobalShadowLowerBound, - 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; } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp index c8c5ed786e1f8..63d5a21819a9a 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp @@ -1,9 +1,7 @@ -// REQUIRES: linux, cpu || (gpu && level_zero) -// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t1.out +// 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 -DMALLOC_DEVICE -O1 -g -o %t2.out -// RUN: %{run} not %t2.out 2>&1 | FileCheck %s -// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O2 -g -o %t3.out +// RUN: %{build} %device_asan_flags -O2 -g -o %t3.out // RUN: %{run} not %t3.out 2>&1 | FileCheck %s #include @@ -13,7 +11,7 @@ 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.cpp}}:[[@LINE-4]] +// CHECK: {{.*arbitary_access.cpp}}:[[@LINE-4]] int main() { sycl::queue Q; diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp index f31642369cac3..3fc370f9a6c5b 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp +++ b/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 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; } diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp index e027f10e9c1e1..e15d189fa5663 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp @@ -69,7 +69,7 @@ inline const char *ToString(ErrorType ErrorType) { } // clang-format off -enum class MemoryType : int32_t { +enum MemoryType : uint32_t { UNKNOWN = 0x000000'00, GLOBAL = 0x000001'00, USM_DEVICE = 0x000001'01, @@ -79,6 +79,8 @@ enum class MemoryType : int32_t { DEVICE_GLOBAL = 0x000001'05, LOCAL = 0x000002'00, PRIVATE = 0x000004'00, + CONSTANT = 0x000008'00, + GENERIC = 0x000010'00, }; // clang-format on From 499b4a10e424f3ffab22d0766d4ef184865c2ed6 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 15 Sep 2025 07:59:36 +0200 Subject: [PATCH 3/9] fix spell --- .../USM/{arbitary_access.cpp => arbitrary_access.cpp} | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) rename sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/{arbitary_access.cpp => arbitrary_access.cpp} (92%) diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp similarity index 92% rename from sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp rename to sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp index 63d5a21819a9a..44addda7dcc2f 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp @@ -11,7 +11,7 @@ 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]] +// CHECK: {{.*arbitrary_access.cpp}}:[[@LINE-4]] int main() { sycl::queue Q; From 14b373d96fbb372306b8752b406d809aa6405ed0 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 16 Sep 2025 10:09:30 +0200 Subject: [PATCH 4/9] wip --- libdevice/sanitizer/asan_rtl.cpp | 14 ++-- .../SPIRVSanitizerCommonUtils.h | 1 + .../Instrumentation/AddressSanitizer.cpp | 66 ++++++++++++++----- .../sanitizer/asan/asan_interceptor.cpp | 52 +++++++++++---- .../sanitizer/asan/asan_interceptor.hpp | 19 +++++- .../sanitizer_common/sanitizer_common.hpp | 1 + .../sanitizer_common/sanitizer_options.cpp | 16 +++++ .../sanitizer_common/sanitizer_options.hpp | 2 - 8 files changed, 131 insertions(+), 40 deletions(-) diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index 9f7c1a1bbd44d..ef668ff2348ae 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/libdevice/sanitizer/asan_rtl.cpp @@ -9,6 +9,8 @@ #include "include/asan_rtl.hpp" #include "asan/asan_libdevice.hpp" +extern "C" __attribute__((weak)) const int __asan_check_shadow_bounds; + // Save the pointer to LaunchInfo __SYCL_GLOBAL__ uptr *__SYCL_LOCAL__ __AsanLaunchInfo; @@ -68,6 +70,8 @@ struct DebugInfo { uint32_t line; }; +inline bool IsCheckShadowBounds() { return __asan_check_shadow_bounds; } + 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, @@ -120,8 +124,9 @@ 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) { + if (IsCheckShadowBounds() && + (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)); @@ -213,8 +218,9 @@ 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) { + if (IsCheckShadowBounds() && + (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)); diff --git a/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h b/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h index d6a206a40bdd6..ceea4b5cd5179 100644 --- a/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h +++ b/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h @@ -50,6 +50,7 @@ enum SanitizedKernelFlags : uint32_t { CHECK_PRIVATES = 1U << 3, CHECK_GENERICS = 1U << 4, MSAN_TRACK_ORIGINS = 1U << 5, + ASAN_CHECK_SHADOW_BOUNDS = 1U << 6, }; } // namespace llvm diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 61068979bc0c8..df800cd510bce 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -441,27 +441,35 @@ static cl::opt ClOverrideDestructorKind( cl::init(AsanDtorKind::Invalid), cl::Hidden); // SYCL flags +static cl::opt ClSpirOffloadPrivates( + "asan-spir-privates", + cl::desc("Instrument private pointer on SPIR-V target"), cl::Hidden, + cl::init(true)); + static cl::opt - ClSpirOffloadPrivates("asan-spir-privates", - cl::desc("instrument private pointer"), cl::Hidden, - cl::init(true)); + ClSpirOffloadGlobals("asan-spir-globals", + cl::desc("Instrument global pointer on SPIR-V target"), + cl::Hidden, cl::init(true)); -static cl::opt ClSpirOffloadGlobals("asan-spir-globals", - cl::desc("instrument global pointer"), - cl::Hidden, cl::init(true)); +static cl::opt + ClSpirOffloadLocals("asan-spir-locals", + cl::desc("Instrument local pointer on SPIR-V target"), + cl::Hidden, cl::init(true)); -static cl::opt ClSpirOffloadLocals("asan-spir-locals", - cl::desc("instrument local pointer"), - cl::Hidden, cl::init(true)); +static cl::opt ClSpirOffloadGenerics( + "asan-spir-generics", + cl::desc("Instrument generic pointer on SPIR-V target"), cl::Hidden, + cl::init(true)); static cl::opt - ClSpirOffloadGenerics("asan-spir-generics", - cl::desc("instrument generic pointer"), cl::Hidden, - cl::init(true)); + ClDeviceGlobals("asan-device-globals", + cl::desc("Instrument device globals on SPIR-V target"), + cl::Hidden, cl::init(true)); -static cl::opt ClDeviceGlobals("asan-device-globals", - cl::desc("instrument device globals"), - cl::Hidden, cl::init(true)); +static cl::opt ClSpirCheckShadowBounds( + "asan-spir-shadow-bounds", + cl::desc("Enable checking shadow bounds on SPIR-V target"), cl::Hidden, + cl::init(false)); // Debug flags. @@ -1411,7 +1419,8 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, // following structure: // uptr unmangled_kernel_name // uptr unmangled_kernel_name_size - StructType *StructTy = StructType::get(IntptrTy, IntptrTy); + // uptr sanitized_flags + StructType *StructTy = StructType::get(IntptrTy, IntptrTy, IntptrTy); if (!HasESIMD) for (Function &F : M) { @@ -1442,9 +1451,21 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, KernelNamesBytes.append(KernelName.begin(), KernelName.end()); auto *KernelNameGV = GetOrCreateGlobalString( M, "__asan_kernel", KernelName, kSpirOffloadConstantAS); + + uintptr_t SanitizerFlags = 0; + SanitizerFlags |= ClSpirOffloadLocals ? SanitizedKernelFlags::CHECK_LOCALS + : SanitizedKernelFlags::NO_CHECK; + SanitizerFlags |= ClSpirOffloadPrivates + ? SanitizedKernelFlags::CHECK_PRIVATES + : SanitizedKernelFlags::NO_CHECK; + SanitizerFlags |= ClSpirCheckShadowBounds != 0 + ? SanitizedKernelFlags::ASAN_CHECK_SHADOW_BOUNDS + : SanitizedKernelFlags::NO_CHECK; + SpirKernelsMetadata.emplace_back(ConstantStruct::get( StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy), - ConstantInt::get(IntptrTy, KernelName.size()))); + ConstantInt::get(IntptrTy, KernelName.size()), + ConstantInt::get(IntptrTy, SanitizerFlags))); } // Create global variable to record spirv kernels' information @@ -1632,6 +1653,17 @@ PreservedAnalyses AddressSanitizerPass::run(Module &M, ExtendSpirKernelArgs(M, FAM, HasESIMD); Modified = true; + { + IRBuilder<> IRB(M.getContext()); + M.getOrInsertGlobal("__asan_check_shadow_bounds", IRB.getInt32Ty(), [&] { + return new GlobalVariable( + M, IRB.getInt32Ty(), true, GlobalValue::WeakODRLinkage, + ConstantInt::get(IRB.getInt32Ty(), ClSpirCheckShadowBounds), + "__asan_check_shadow_bounds", nullptr, + llvm::GlobalValue::NotThreadLocal, kSpirOffloadGlobalAS); + }); + } + if (HasESIMD) { GlobalStringMap.clear(); return PreservedAnalyses::none(); diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp index 3fc370f9a6c5b..a02abc8a8b74e 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp @@ -465,7 +465,7 @@ ur_result_t AsanInterceptor::unregisterProgram(ur_program_handle_t Program) { } ProgramInfo->AllocInfoForGlobals.clear(); - ProgramInfo->InstrumentedKernels.clear(); + ProgramInfo->KernelMetadataMap.clear(); return UR_RESULT_SUCCESS; } @@ -520,14 +520,21 @@ ur_result_t AsanInterceptor::registerSpirKernels(ur_program_handle_t Program) { std::string KernelName = std::string(KernelNameV.begin(), KernelNameV.end()); + bool CheckLocals = SKI.Flags & SanitizedKernelFlags::CHECK_LOCALS; + bool CheckPrivates = SKI.Flags & SanitizedKernelFlags::CHECK_PRIVATES; + bool CheckShadowBounds = + SKI.Flags & SanitizedKernelFlags::ASAN_CHECK_SHADOW_BOUNDS; UR_LOG_L(getContext()->logger, INFO, - "SpirKernel(name='{}', isInstrumented={})", KernelName, true); + "SpirKernel(name='{}', isInstrumented={}, " + "checkLocals={}, checkPrivates={}, checkShadowBounds={})", + KernelName, true, CheckLocals, CheckPrivates, CheckShadowBounds); - PI->InstrumentedKernels.insert(std::move(KernelName)); + PI->KernelMetadataMap[KernelName] = ProgramInfo::KernelMetada{ + CheckLocals, CheckPrivates, CheckShadowBounds}; } UR_LOG_L(getContext()->logger, INFO, "Number of sanitized kernel: {}", - PI->InstrumentedKernels.size()); + PI->KernelMetadataMap.size()); } return UR_RESULT_SUCCESS; @@ -691,11 +698,18 @@ KernelInfo &AsanInterceptor::getOrCreateKernelInfo(ur_kernel_handle_t Kernel) { auto Program = GetProgram(Kernel); auto PI = getProgramInfo(Program); assert(PI != nullptr && "unregistered program!"); - bool IsInstrumented = PI->isKernelInstrumented(Kernel); + + auto KI = std::make_unique(Kernel); + KI->IsInstrumented = PI->isKernelInstrumented(Kernel); + if (KI->IsInstrumented) { + auto &KM = PI->getKernelMetadata(Kernel); + KI->IsCheckLocals = KM.CheckLocals; + KI->IsCheckPrivates = KM.CheckPrivates; + KI->IsCheckShadowBounds = KM.CheckShadowBounds; + } std::scoped_lock Guard(m_KernelMapMutex); - m_KernelMap.emplace(Kernel, - std::make_unique(Kernel, IsInstrumented)); + m_KernelMap.emplace(Kernel, std::move(KI)); return *m_KernelMap[Kernel].get(); } @@ -837,14 +851,17 @@ 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; + if (KernelInfo.IsCheckShadowBounds) { + LaunchInfo.Data.Host.GlobalShadowLowerBound = + DeviceInfo->Shadow->ShadowLowerBound; + LaunchInfo.Data.Host.GlobalShadowUpperBound = + DeviceInfo->Shadow->ShadowUpperBound; + } + // Write shadow memory offset for local memory - if (getContext()->Options.DetectLocals) { + if (KernelInfo.IsCheckLocals) { if (DeviceInfo->Shadow->AllocLocalShadow( Queue, NumWG, LaunchInfo.Data.Host.LocalShadowOffset, LaunchInfo.Data.Host.LocalShadowOffsetEnd) != UR_RESULT_SUCCESS) { @@ -864,7 +881,7 @@ ur_result_t AsanInterceptor::prepareLaunch( } // Write shadow memory offset for private memory - if (getContext()->Options.DetectPrivates) { + if (KernelInfo.IsCheckPrivates) { if (DeviceInfo->Shadow->AllocPrivateShadow( Queue, NumSG, LaunchInfo.Data.Host.PrivateBase, LaunchInfo.Data.Host.PrivateShadowOffset, @@ -950,7 +967,14 @@ AsanInterceptor::findAllocInfoByContext(ur_context_handle_t Context) { bool ProgramInfo::isKernelInstrumented(ur_kernel_handle_t Kernel) const { const auto Name = GetKernelName(Kernel); - return InstrumentedKernels.find(Name) != InstrumentedKernels.end(); + return KernelMetadataMap.find(Name) != KernelMetadataMap.end(); +} + +const ProgramInfo::KernelMetada & +ProgramInfo::getKernelMetadata(ur_kernel_handle_t Kernel) const { + const auto Name = GetKernelName(Kernel); + assert(KernelMetadataMap.find(Name) != KernelMetadataMap.end()); + return KernelMetadataMap.at(Name); } ContextInfo::~ContextInfo() { diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp index a0c69fe038807..aee8f9209e8e7 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp @@ -86,6 +86,12 @@ struct KernelInfo { // sanitized kernel bool IsInstrumented = false; + // check local memory + bool IsCheckLocals = true; + // check private memory + bool IsCheckPrivates = true; + // check shadow bounds + bool IsCheckShadowBounds = false; // lock this mutex if following fields are accessed ur_shared_mutex Mutex; @@ -95,8 +101,7 @@ struct KernelInfo { // Need preserve the order of local arguments std::map LocalArgs; - explicit KernelInfo(ur_kernel_handle_t Kernel, bool IsInstrumented) - : Handle(Kernel), IsInstrumented(IsInstrumented) { + explicit KernelInfo(ur_kernel_handle_t Kernel) : Handle(Kernel) { [[maybe_unused]] auto Result = getContext()->urDdiTable.Kernel.pfnRetain(Kernel); assert(Result == UR_RESULT_SUCCESS); @@ -113,9 +118,15 @@ struct ProgramInfo { ur_program_handle_t Handle; std::atomic RefCount = 1; + struct KernelMetada { + bool CheckLocals; + bool CheckPrivates; + bool CheckShadowBounds; + }; + // Program is built only once, so we don't need to lock it std::unordered_set> AllocInfoForGlobals; - std::unordered_set InstrumentedKernels; + std::unordered_map KernelMetadataMap; explicit ProgramInfo(ur_program_handle_t Program) : Handle(Program) { [[maybe_unused]] auto Result = @@ -130,6 +141,7 @@ struct ProgramInfo { } bool isKernelInstrumented(ur_kernel_handle_t Kernel) const; + const KernelMetada &getKernelMetadata(ur_kernel_handle_t Kernel) const; }; struct ContextInfo { @@ -277,6 +289,7 @@ struct DeviceGlobalInfo { struct SpirKernelInfo { uptr KernelName; uptr Size; + uptr Flags; }; class AsanInterceptor { diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_common.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_common.hpp index c4854d96ad2e2..c122557b000c9 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_common.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_common.hpp @@ -157,6 +157,7 @@ enum SanitizedKernelFlags : uint32_t { CHECK_PRIVATES = 1U << 3, CHECK_GENERICS = 1U << 4, MSAN_TRACK_ORIGINS = 1U << 5, + ASAN_CHECK_SHADOW_BOUNDS = 1U << 6, }; } // namespace ur_sanitizer_layer diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp index 7ecca4e62387e..14774c792d122 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp @@ -41,8 +41,24 @@ void SanitizerOptions::Init(const std::string &EnvName, Parser.ParseBool("debug", Debug); Parser.ParseBool("detect_kernel_arguments", DetectKernelArguments); + + bool DetectLocals = true; Parser.ParseBool("detect_locals", DetectLocals); + if (!DetectLocals) { + UR_LOG_L(Logger, WARN, + "Option `detect_locals` has been deprecated. Just using compiler " + "flag is enough."); + } + + bool DetectPrivates = true; Parser.ParseBool("detect_privates", DetectPrivates); + if (!DetectPrivates) { + UR_LOG_L( + Logger, WARN, + "Option `detect_privates` has been deprecated. Just using compiler " + "flag is enough."); + } + Parser.ParseBool("print_stats", PrintStats); Parser.ParseBool("detect_leaks", DetectLeaks); Parser.ParseBool("halt_on_error", HaltOnError); diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp index 328bdc05b2171..d1a58798b7687 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp @@ -25,8 +25,6 @@ struct SanitizerOptions { bool Debug = false; uint64_t MinRZSize = 16; uint64_t MaxQuarantineSizeMB = 8; - bool DetectLocals = true; - bool DetectPrivates = true; bool PrintStats = false; bool DetectKernelArguments = true; bool DetectLeaks = true; From 4f9c73270562375ddbac312ac04c4a7b0c051550 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 16 Sep 2025 10:31:31 +0200 Subject: [PATCH 5/9] wip --- libdevice/sanitizer/asan_rtl.cpp | 12 ------------ .../sanitizer_common/sanitizer_libdevice.hpp | 3 +++ 2 files changed, 3 insertions(+), 12 deletions(-) diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index ef668ff2348ae..2230d867d3762 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/libdevice/sanitizer/asan_rtl.cpp @@ -139,8 +139,6 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as, 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; }); @@ -162,8 +160,6 @@ 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; @@ -190,8 +186,6 @@ 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; } @@ -233,8 +227,6 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as, 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; @@ -255,8 +247,6 @@ 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; @@ -283,8 +273,6 @@ 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; }; diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp index e15d189fa5663..5ecfe351099de 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp @@ -69,6 +69,9 @@ inline const char *ToString(ErrorType ErrorType) { } // clang-format off +// We treat "GLOBAL/LOCAL/PRIVATE/CONSTANT/GENERIC" as address space mask as well, +// So it's easy to check that USM_XXX is also in global memory, and we can also +// mark an address is a generic & global & usm_device address enum MemoryType : uint32_t { UNKNOWN = 0x000000'00, GLOBAL = 0x000001'00, From 64b2b45b3fd0f6f9525de7078a2fac15ec586e82 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 17 Sep 2025 04:55:09 +0200 Subject: [PATCH 6/9] add comp flag asan-spir-shadow-bounds --- .../Instrumentation/AddressSanitizer.cpp | 16 +++++++++------- .../out-of-bounds/USM/arbitrary_access.cpp | 4 ++-- 2 files changed, 11 insertions(+), 9 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index df800cd510bce..9c9ba40a77745 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -1742,19 +1742,21 @@ static bool isUnsupportedAMDGPUAddrspace(Value *Addr) { } static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { - // Non image scope device globals are implemented by device USM, and the - // out-of-bounds check for them will be done by sanitizer USM part. So we - // exclude them here. - if (!G->hasAttribute("sycl-device-image-scope")) - return true; - // Skip instrumenting on "__AsanKernelMetadata" etc. - if (G->getName().starts_with("__Asan")) + if (G->getName().starts_with("__Asan") || G->getName().starts_with("__asan")) return true; if (G->getAddressSpace() == kSpirOffloadLocalAS) return !ClSpirOffloadLocals; + // When shadow bounds check is disabled, we need to instrument all global + // variables that user code can access + if (ClSpirCheckShadowBounds) + return false; + + // Non image scope device globals are implemented by device USM, and the + // out-of-bounds check for them will be done by sanitizer USM part. So we + // exclude them here. Attribute Attr = G->getAttribute("sycl-device-image-scope"); return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false"); } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp index 44addda7dcc2f..02c233a179158 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitrary_access.cpp @@ -1,7 +1,7 @@ // REQUIRES: linux, gpu && level_zero -// RUN: %{build} %device_asan_flags -O0 -g -o %t1.out +// RUN: %{build} %device_asan_flags -Xarch_device -mllvm=-asan-spir-shadow-bounds=1 -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: %{build} %device_asan_flags -Xarch_device -mllvm=-asan-spir-shadow-bounds=1 -O2 -g -o %t3.out // RUN: %{run} not %t3.out 2>&1 | FileCheck %s #include From f6635677080525ea937a6b51339163c6bfde61d3 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 17 Sep 2025 05:17:28 +0200 Subject: [PATCH 7/9] fix copilot comments --- libdevice/sanitizer/asan_rtl.cpp | 2 +- llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp | 2 +- .../loader/layers/sanitizer/asan/asan_interceptor.cpp | 4 ++-- .../loader/layers/sanitizer/asan/asan_interceptor.hpp | 6 +++--- .../loader/layers/sanitizer/msan/msan_interceptor.cpp | 4 ++-- .../loader/layers/sanitizer/msan/msan_interceptor.hpp | 6 +++--- 6 files changed, 12 insertions(+), 12 deletions(-) diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index 2230d867d3762..1525c021ea395 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/libdevice/sanitizer/asan_rtl.cpp @@ -43,7 +43,7 @@ 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"; + "[kernel] addr: %p, shadow: %p, lower: %p, upper: %p\n"; #define ASAN_REPORT_NONE 0 #define ASAN_REPORT_START 1 diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 9c9ba40a77745..d3506c3bc508e 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -1749,7 +1749,7 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { if (G->getAddressSpace() == kSpirOffloadLocalAS) return !ClSpirOffloadLocals; - // When shadow bounds check is disabled, we need to instrument all global + // When shadow bounds check is enabled, we need to instrument all global // variables that user code can access if (ClSpirCheckShadowBounds) return false; diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp index a02abc8a8b74e..bce0a911260a8 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp @@ -530,7 +530,7 @@ ur_result_t AsanInterceptor::registerSpirKernels(ur_program_handle_t Program) { "checkLocals={}, checkPrivates={}, checkShadowBounds={})", KernelName, true, CheckLocals, CheckPrivates, CheckShadowBounds); - PI->KernelMetadataMap[KernelName] = ProgramInfo::KernelMetada{ + PI->KernelMetadataMap[KernelName] = ProgramInfo::KernelMetadata{ CheckLocals, CheckPrivates, CheckShadowBounds}; } UR_LOG_L(getContext()->logger, INFO, "Number of sanitized kernel: {}", @@ -970,7 +970,7 @@ bool ProgramInfo::isKernelInstrumented(ur_kernel_handle_t Kernel) const { return KernelMetadataMap.find(Name) != KernelMetadataMap.end(); } -const ProgramInfo::KernelMetada & +const ProgramInfo::KernelMetadata & ProgramInfo::getKernelMetadata(ur_kernel_handle_t Kernel) const { const auto Name = GetKernelName(Kernel); assert(KernelMetadataMap.find(Name) != KernelMetadataMap.end()); diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp index aee8f9209e8e7..97a478b4dc694 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp @@ -118,7 +118,7 @@ struct ProgramInfo { ur_program_handle_t Handle; std::atomic RefCount = 1; - struct KernelMetada { + struct KernelMetadata { bool CheckLocals; bool CheckPrivates; bool CheckShadowBounds; @@ -126,7 +126,7 @@ struct ProgramInfo { // Program is built only once, so we don't need to lock it std::unordered_set> AllocInfoForGlobals; - std::unordered_map KernelMetadataMap; + std::unordered_map KernelMetadataMap; explicit ProgramInfo(ur_program_handle_t Program) : Handle(Program) { [[maybe_unused]] auto Result = @@ -141,7 +141,7 @@ struct ProgramInfo { } bool isKernelInstrumented(ur_kernel_handle_t Kernel) const; - const KernelMetada &getKernelMetadata(ur_kernel_handle_t Kernel) const; + const KernelMetadata &getKernelMetadata(ur_kernel_handle_t Kernel) const; }; struct ContextInfo { diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp index 4838a1a58e22e..0b3022d0554cd 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp @@ -256,7 +256,7 @@ ur_result_t MsanInterceptor::registerSpirKernels(ur_program_handle_t Program) { KernelName, true, CheckLocals, CheckPrivates, TrackOrigins); PI->KernelMetadataMap[KernelName] = - ProgramInfo::KernelMetada{CheckLocals, CheckPrivates, TrackOrigins}; + ProgramInfo::KernelMetadata{CheckLocals, CheckPrivates, TrackOrigins}; } UR_LOG_L(getContext()->logger, INFO, "Number of sanitized kernel: {}", PI->KernelMetadataMap.size()); @@ -650,7 +650,7 @@ bool ProgramInfo::isKernelInstrumented(ur_kernel_handle_t Kernel) const { return KernelMetadataMap.find(Name) != KernelMetadataMap.end(); } -const ProgramInfo::KernelMetada & +const ProgramInfo::KernelMetadata & ProgramInfo::getKernelMetadata(ur_kernel_handle_t Kernel) const { const auto Name = GetKernelName(Kernel); assert(KernelMetadataMap.find(Name) != KernelMetadataMap.end()); diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp index 0098783dfe8d6..6c020974a0004 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp @@ -110,14 +110,14 @@ struct ProgramInfo { ur_program_handle_t Handle; std::atomic RefCount = 1; - struct KernelMetada { + struct KernelMetadata { bool CheckLocals; bool CheckPrivates; bool TrackOrigins; }; // Program is built only once, so we don't need to lock it - std::unordered_map KernelMetadataMap; + std::unordered_map KernelMetadataMap; explicit ProgramInfo(ur_program_handle_t Program) : Handle(Program) { [[maybe_unused]] auto Result = @@ -132,7 +132,7 @@ struct ProgramInfo { } bool isKernelInstrumented(ur_kernel_handle_t Kernel) const; - const KernelMetada &getKernelMetadata(ur_kernel_handle_t Kernel) const; + const KernelMetadata &getKernelMetadata(ur_kernel_handle_t Kernel) const; }; struct ContextInfo { From 4e3ca14be1c7912a28e1e74c5919679166e9336b Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Thu, 18 Sep 2025 07:19:59 +0200 Subject: [PATCH 8/9] fix lit failure --- .../AddressSanitizer/SPIRV/extend_launch_info_arg.ll | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll index 1590d1f8997fb..7f844a10ed49f 100644 --- a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll +++ b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll @@ -26,4 +26,4 @@ entry: attributes #0 = { sanitize_address } ;; sycl-device-global-size = 16 * 2 ;; sycl-host-access = 0 read-only -; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" } +; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="48" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" } From cf2358cd558216e1bb25a77ac382d6dd7818f01a Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 23 Sep 2025 04:47:00 +0200 Subject: [PATCH 9/9] revert changes related to local/private --- .../layers/sanitizer/asan/asan_interceptor.cpp | 17 ++++++----------- .../layers/sanitizer/asan/asan_interceptor.hpp | 6 ------ .../sanitizer_common/sanitizer_options.cpp | 16 ---------------- .../sanitizer_common/sanitizer_options.hpp | 2 ++ 4 files changed, 8 insertions(+), 33 deletions(-) diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp index 805a2b98fa43b..45eed900f9150 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp @@ -520,18 +520,15 @@ ur_result_t AsanInterceptor::registerSpirKernels(ur_program_handle_t Program) { std::string KernelName = std::string(KernelNameV.begin(), KernelNameV.end()); - bool CheckLocals = SKI.Flags & SanitizedKernelFlags::CHECK_LOCALS; - bool CheckPrivates = SKI.Flags & SanitizedKernelFlags::CHECK_PRIVATES; bool CheckShadowBounds = SKI.Flags & SanitizedKernelFlags::ASAN_CHECK_SHADOW_BOUNDS; UR_LOG_L(getContext()->logger, INFO, - "SpirKernel(name='{}', isInstrumented={}, " - "checkLocals={}, checkPrivates={}, checkShadowBounds={})", - KernelName, true, CheckLocals, CheckPrivates, CheckShadowBounds); + "SpirKernel(name='{}', isInstrumented={}, checkShadowBounds={})", + KernelName, true, CheckShadowBounds); - PI->KernelMetadataMap[KernelName] = ProgramInfo::KernelMetadata{ - CheckLocals, CheckPrivates, CheckShadowBounds}; + PI->KernelMetadataMap[KernelName] = + ProgramInfo::KernelMetadata{CheckShadowBounds}; } UR_LOG_L(getContext()->logger, INFO, "Number of sanitized kernel: {}", PI->KernelMetadataMap.size()); @@ -690,8 +687,6 @@ KernelInfo &AsanInterceptor::getOrCreateKernelInfo(ur_kernel_handle_t Kernel) { KI->IsInstrumented = PI->isKernelInstrumented(Kernel); if (KI->IsInstrumented) { auto &KM = PI->getKernelMetadata(Kernel); - KI->IsCheckLocals = KM.CheckLocals; - KI->IsCheckPrivates = KM.CheckPrivates; KI->IsCheckShadowBounds = KM.CheckShadowBounds; } @@ -849,7 +844,7 @@ ur_result_t AsanInterceptor::prepareLaunch( } // Write shadow memory offset for local memory - if (KernelInfo.IsCheckLocals) { + if (getContext()->Options.DetectLocals) { if (DeviceInfo->Shadow->AllocLocalShadow( Queue, NumWG, LaunchInfo.Data.Host.LocalShadowOffset, LaunchInfo.Data.Host.LocalShadowOffsetEnd) != UR_RESULT_SUCCESS) { @@ -869,7 +864,7 @@ ur_result_t AsanInterceptor::prepareLaunch( } // Write shadow memory offset for private memory - if (KernelInfo.IsCheckPrivates) { + if (getContext()->Options.DetectPrivates) { if (DeviceInfo->Shadow->AllocPrivateShadow( Queue, NumSG, LaunchInfo.Data.Host.PrivateBase, LaunchInfo.Data.Host.PrivateShadowOffset, diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp index 97a478b4dc694..7e4bc8edee7e3 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp @@ -86,10 +86,6 @@ struct KernelInfo { // sanitized kernel bool IsInstrumented = false; - // check local memory - bool IsCheckLocals = true; - // check private memory - bool IsCheckPrivates = true; // check shadow bounds bool IsCheckShadowBounds = false; @@ -119,8 +115,6 @@ struct ProgramInfo { std::atomic RefCount = 1; struct KernelMetadata { - bool CheckLocals; - bool CheckPrivates; bool CheckShadowBounds; }; diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp index 14774c792d122..7ecca4e62387e 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp @@ -41,24 +41,8 @@ void SanitizerOptions::Init(const std::string &EnvName, Parser.ParseBool("debug", Debug); Parser.ParseBool("detect_kernel_arguments", DetectKernelArguments); - - bool DetectLocals = true; Parser.ParseBool("detect_locals", DetectLocals); - if (!DetectLocals) { - UR_LOG_L(Logger, WARN, - "Option `detect_locals` has been deprecated. Just using compiler " - "flag is enough."); - } - - bool DetectPrivates = true; Parser.ParseBool("detect_privates", DetectPrivates); - if (!DetectPrivates) { - UR_LOG_L( - Logger, WARN, - "Option `detect_privates` has been deprecated. Just using compiler " - "flag is enough."); - } - Parser.ParseBool("print_stats", PrintStats); Parser.ParseBool("detect_leaks", DetectLeaks); Parser.ParseBool("halt_on_error", HaltOnError); diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp index d1a58798b7687..328bdc05b2171 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.hpp @@ -25,6 +25,8 @@ struct SanitizerOptions { bool Debug = false; uint64_t MinRZSize = 16; uint64_t MaxQuarantineSizeMB = 8; + bool DetectLocals = true; + bool DetectPrivates = true; bool PrintStats = false; bool DetectKernelArguments = true; bool DetectLeaks = true;