From 2a5f9137ca2a6c1004a059ed95d3bfd79cf3ad41 Mon Sep 17 00:00:00 2001 From: Yang Zhao Date: Fri, 12 Jul 2024 03:43:01 +0800 Subject: [PATCH] [DeviceSanitizer] Support detecting misaligned access error (#14148) UR: https://github.com/oneapi-src/unified-runtime/pull/1747 --------- Co-authored-by: Kenneth Benzie (Benie) --- libdevice/sanitizer_utils.cpp | 110 ++++++++++++------ sycl/plugins/unified_runtime/CMakeLists.txt | 12 +- .../misaligned/misalign-int.cpp | 35 ++++++ .../misaligned/misalign-long.cpp | 35 ++++++ .../misaligned/misalign-short.cpp | 29 +++++ 5 files changed, 182 insertions(+), 39 deletions(-) create mode 100644 sycl/test-e2e/AddressSanitizer/misaligned/misalign-int.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/misaligned/misalign-long.cpp create mode 100644 sycl/test-e2e/AddressSanitizer/misaligned/misalign-short.cpp diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp index 651067be69851..ef7de515ade70 100644 --- a/libdevice/sanitizer_utils.cpp +++ b/libdevice/sanitizer_utils.cpp @@ -322,6 +322,9 @@ bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) { /// ASAN Save Report /// +static __SYCL_CONSTANT__ const char __mem_sanitizer_report[] = + "[kernel] SanitizerReport (ErrorType=%d, IsRecover=%d)\n"; + bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) { const int Expected = ASAN_REPORT_NONE; int Desired = ASAN_REPORT_START; @@ -339,8 +342,14 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) { if (atomicCompareAndSet(&SanitizerReport.Flag, Desired, Expected) == Expected) { SanitizerReport.ErrorType = error_type; + SanitizerReport.IsRecover = false; + // Show we've done copying atomicStore(&SanitizerReport.Flag, ASAN_REPORT_FINISH); + + if (__AsanDebug) + __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType, + SanitizerReport.IsRecover); return true; } return false; @@ -419,6 +428,10 @@ bool __asan_internal_report_save( // Show we've done copying atomicStore(&SanitizerReport.Flag, ASAN_REPORT_FINISH); + + if (__AsanDebug) + __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType, + SanitizerReport.IsRecover); return true; } return false; @@ -428,6 +441,32 @@ bool __asan_internal_report_save( /// ASAN Error Reporters /// +DeviceSanitizerMemoryType GetMemoryTypeByShadowValue(int shadow_value) { + switch (shadow_value) { + case kUsmDeviceRedzoneMagic: + case kUsmDeviceDeallocatedMagic: + return DeviceSanitizerMemoryType::USM_DEVICE; + case kUsmHostRedzoneMagic: + case kUsmHostDeallocatedMagic: + return DeviceSanitizerMemoryType::USM_HOST; + case kUsmSharedRedzoneMagic: + case kUsmSharedDeallocatedMagic: + return DeviceSanitizerMemoryType::USM_SHARED; + case kPrivateLeftRedzoneMagic: + case kPrivateMidRedzoneMagic: + case kPrivateRightRedzoneMagic: + return DeviceSanitizerMemoryType::PRIVATE; + case kMemBufferRedzoneMagic: + return DeviceSanitizerMemoryType::MEM_BUFFER; + case kSharedLocalRedzoneMagic: + return DeviceSanitizerMemoryType::LOCAL; + case kDeviceGlobalRedzoneMagic: + return DeviceSanitizerMemoryType::DEVICE_GLOBAL; + default: + return DeviceSanitizerMemoryType::UNKNOWN; + } +} + void __asan_report_access_error(uptr addr, uint32_t as, size_t size, bool is_write, uptr poisoned_addr, const char __SYCL_CONSTANT__ *file, @@ -442,54 +481,28 @@ void __asan_report_access_error(uptr addr, uint32_t as, size_t size, } // FIXME: check if shadow_address out-of-bound - DeviceSanitizerMemoryType memory_type; + DeviceSanitizerMemoryType memory_type = + GetMemoryTypeByShadowValue(shadow_value); DeviceSanitizerErrorType error_type; switch (shadow_value) { case kUsmDeviceRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::USM_DEVICE; - error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; - break; case kUsmHostRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::USM_HOST; - error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; - break; case kUsmSharedRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::USM_SHARED; - error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; - break; - case kUsmDeviceDeallocatedMagic: - memory_type = DeviceSanitizerMemoryType::USM_DEVICE; - error_type = DeviceSanitizerErrorType::USE_AFTER_FREE; - break; - case kUsmHostDeallocatedMagic: - memory_type = DeviceSanitizerMemoryType::USM_HOST; - error_type = DeviceSanitizerErrorType::USE_AFTER_FREE; - break; - case kUsmSharedDeallocatedMagic: - memory_type = DeviceSanitizerMemoryType::USM_SHARED; - error_type = DeviceSanitizerErrorType::USE_AFTER_FREE; - break; case kPrivateLeftRedzoneMagic: case kPrivateMidRedzoneMagic: case kPrivateRightRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::PRIVATE; - error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; - break; case kMemBufferRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::MEM_BUFFER; - error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; - break; case kSharedLocalRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::LOCAL; - error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; - break; case kDeviceGlobalRedzoneMagic: - memory_type = DeviceSanitizerMemoryType::DEVICE_GLOBAL; error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS; break; + case kUsmDeviceDeallocatedMagic: + case kUsmHostDeallocatedMagic: + case kUsmSharedDeallocatedMagic: + error_type = DeviceSanitizerErrorType::USE_AFTER_FREE; + break; default: - memory_type = DeviceSanitizerMemoryType::UNKNOWN; error_type = DeviceSanitizerErrorType::UNKNOWN; } @@ -497,6 +510,27 @@ void __asan_report_access_error(uptr addr, uint32_t as, size_t size, memory_type, error_type, is_recover); } +void __asan_report_misalign_error(uptr addr, uint32_t as, size_t size, + bool is_write, uptr poisoned_addr, + const char __SYCL_CONSTANT__ *file, + uint32_t line, + const char __SYCL_CONSTANT__ *func, + bool is_recover = false) { + + auto *shadow = (__SYCL_GLOBAL__ s8 *)MemToShadow(addr, as); + while (*shadow >= 0) { + ++shadow; + } + int shadow_value = *shadow; + + DeviceSanitizerErrorType error_type = DeviceSanitizerErrorType::MISALIGNED; + DeviceSanitizerMemoryType memory_type = + GetMemoryTypeByShadowValue(shadow_value); + + __asan_internal_report_save(addr, as, file, line, func, is_write, size, + memory_type, error_type, is_recover); +} + bool __asan_report_unknown_device() { return __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE); } @@ -589,6 +623,8 @@ inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) { return 0; } +constexpr size_t AlignMask(size_t n) { return n - 1; } + } // namespace /// @@ -599,6 +635,10 @@ inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) { DEVICE_EXTERN_C_NOINLINE void __asan_##type##size( \ uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file, \ uint32_t line, const char __SYCL_CONSTANT__ *func) { \ + if (addr & AlignMask(size)) { \ + __asan_report_misalign_error(addr, as, size, is_write, addr, file, line, \ + func); \ + } \ if (__asan_address_is_poisoned(addr, as, size)) { \ __asan_report_access_error(addr, as, size, is_write, addr, file, line, \ func); \ @@ -607,6 +647,10 @@ inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) { DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort( \ uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file, \ uint32_t line, const char __SYCL_CONSTANT__ *func) { \ + if (addr & AlignMask(size)) { \ + __asan_report_misalign_error(addr, as, size, is_write, addr, file, line, \ + func, true); \ + } \ if (__asan_address_is_poisoned(addr, as, size)) { \ __asan_report_access_error(addr, as, size, is_write, addr, file, line, \ func, true); \ diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 74f0fe8ce74e4..2daf0401a293b 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -100,13 +100,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) endfunction() set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 642e3437b8a0a3e702ac3d68513d7076e7618c94 - # Merge: 342bdfff 199287b7 + # commit 167ddf9320dba95c9324e373fb863baeb46e09f1 + # Merge: 76c6bf9a 0cd10f63 # Author: Kenneth Benzie (Benie) - # Date: Thu Jul 11 10:37:50 2024 +0100 - # Merge pull request #1829 from callumfare/callum/bindless_workaround - # Treat bindless images handle types as native handles - set(UNIFIED_RUNTIME_TAG 642e3437b8a0a3e702ac3d68513d7076e7618c94) + # Date: Thu Jul 11 15:12:38 2024 +0100 + # Merge pull request #1747 from AllanZyne/review/yang/misalign_access + # [DeviceSanitizer] Support detecting misaligned access error + set(UNIFIED_RUNTIME_TAG 167ddf9320dba95c9324e373fb863baeb46e09f1) fetch_adapter_source(level_zero ${UNIFIED_RUNTIME_REPO} diff --git a/sycl/test-e2e/AddressSanitizer/misaligned/misalign-int.cpp b/sycl/test-e2e/AddressSanitizer/misaligned/misalign-int.cpp new file mode 100644 index 0000000000000..92f088993d24a --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/misaligned/misalign-int.cpp @@ -0,0 +1,35 @@ +// REQUIRES: linux, cpu +// RUN: %{build} %device_asan_flags -O0 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O1 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O2 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +#include +#include + +#include + +int main() { + std::random_device rd; + std::mt19937 gen(rd()); + std::uniform_int_distribution<> distrib(1, 3); + + sycl::queue Q; + constexpr std::size_t N = 4; + auto *array = sycl::malloc_host(N, Q); + auto offset = distrib(gen); + std::cout << "offset: " << offset << std::endl; + array = (int *)((char *)array + offset); + + Q.submit([&](sycl::handler &h) { + h.parallel_for(sycl::nd_range<1>(N, 1), + [=](sycl::nd_item<1> item) { ++array[0]; }); + Q.wait(); + }); + // CHECK: ERROR: DeviceSanitizer: misaligned-access on Host USM + // CHECK: READ of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) + // CHECK: #0 {{.*}} {{.*misalign-int.cpp}}:[[@LINE-5]] + + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/misaligned/misalign-long.cpp b/sycl/test-e2e/AddressSanitizer/misaligned/misalign-long.cpp new file mode 100644 index 0000000000000..2859ca8aaf36b --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/misaligned/misalign-long.cpp @@ -0,0 +1,35 @@ +// REQUIRES: linux, cpu +// RUN: %{build} %device_asan_flags -O0 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O1 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O2 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +#include +#include + +#include + +int main() { + std::random_device rd; + std::mt19937 gen(rd()); + std::uniform_int_distribution<> distrib(1, 7); + + sycl::queue Q; + constexpr std::size_t N = 4; + auto *array = sycl::malloc_shared(N, Q); + auto offset = distrib(gen); + std::cout << "offset: " << offset << std::endl; + array = (long long *)((char *)array + offset); + + Q.submit([&](sycl::handler &h) { + h.parallel_for(sycl::nd_range<1>(N, 1), + [=](sycl::nd_item<1> item) { ++array[0]; }); + Q.wait(); + }); + // CHECK: ERROR: DeviceSanitizer: misaligned-access on Shared USM + // CHECK: READ of size 8 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) + // CHECK: #0 {{.*}} {{.*misalign-long.cpp}}:[[@LINE-5]] + + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/misaligned/misalign-short.cpp b/sycl/test-e2e/AddressSanitizer/misaligned/misalign-short.cpp new file mode 100644 index 0000000000000..acdfb6a1d0a7e --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/misaligned/misalign-short.cpp @@ -0,0 +1,29 @@ +// REQUIRES: linux, cpu +// RUN: %{build} %device_asan_flags -O0 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O1 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O2 -g -o %t +// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s +#include +#include + +#include + +int main() { + sycl::queue Q; + constexpr std::size_t N = 4; + auto *array = sycl::malloc_device(N, Q); + array = (short *)((char *)array + 1); + + Q.submit([&](sycl::handler &h) { + h.parallel_for(sycl::nd_range<1>(N, 1), + [=](sycl::nd_item<1> item) { ++array[0]; }); + Q.wait(); + }); + // CHECK: ERROR: DeviceSanitizer: misaligned-access on Device USM + // CHECK: READ of size 2 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) + // CHECK: #0 {{.*}} {{.*misalign-short.cpp}}:[[@LINE-5]] + + return 0; +}