From dd2202d7216bc9493fd2d8ca11027592dde0ef99 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 31 Jul 2024 22:49:06 -0700 Subject: [PATCH 1/8] add nullpointer detection --- libdevice/sanitizer_utils.cpp | 7 +++++ .../nullpointer/nullpointer.cpp | 26 +++++++++++++++++++ 2 files changed, 33 insertions(+) create mode 100644 sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp index f59bc17bb948e..51eec8dae1f95 100644 --- a/libdevice/sanitizer_utils.cpp +++ b/libdevice/sanitizer_utils.cpp @@ -52,6 +52,8 @@ extern SYCL_EXTERNAL __SYCL_LOCAL__ void * __spirv_GenericCastToPtrExplicit_ToLocal(void *, int); extern SYCL_EXTERNAL __SYCL_PRIVATE__ void * __spirv_GenericCastToPtrExplicit_ToPrivate(void *, int); + +extern "C" SYCL_EXTERNAL void __devicelib_exit(); #endif // __USE_SPIR_BUILTIN__ static const __SYCL_CONSTANT__ char __asan_shadow_value_start[] = @@ -422,6 +424,7 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) { SanitizerReport.IsRecover); return true; } + __devicelib_exit(); return false; } @@ -504,6 +507,7 @@ bool __asan_internal_report_save( SanitizerReport.IsRecover); return true; } + __devicelib_exit(); return false; } @@ -572,6 +576,9 @@ void __asan_report_access_error(uptr addr, uint32_t as, size_t size, case kUsmSharedDeallocatedMagic: error_type = DeviceSanitizerErrorType::USE_AFTER_FREE; break; + case kNullPointerRedzoneMagic: + error_type = DeviceSanitizerErrorType::NULL_POINTER; + break; default: error_type = DeviceSanitizerErrorType::UNKNOWN; } diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp new file mode 100644 index 0000000000000..73e37e0981ee0 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp @@ -0,0 +1,26 @@ +// REQUIRES: linux +// RUN-NOT: %{build} %device_asan_flags -O0 -g -o %t +// RUN-NOT: %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O1 -g -o %t +// RUN: %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O2 -g -o %t +// RUN: %{run} not %t 2>&1 | FileCheck %s + +#include + +int main() { + sycl::queue Q; + constexpr std::size_t N = 4; + int *array = 0; + + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { array[0] = 0; }); + Q.wait(); + }); + // CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory + // CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) + // CHECK: #0 {{.*}} {{.*null-pointer.cpp}}:[[@LINE-5]] + + return 0; +} From b22141993ce6c19d15bc72ba653cd1072a1635b1 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 31 Jul 2024 22:52:36 -0700 Subject: [PATCH 2/8] change ur repo --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 35a9142059418..e303ee0db00d5 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -115,14 +115,14 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/AllanZyne/unified-runtime.git") # commit a985a81dc9ba8adfcc8b54e35ad287e97766fb3e # Merge: b7b0c8b3 f772f907 # Author: Piotr Balcer # Date: Mon Jul 29 09:11:29 2024 +0200 # Merge pull request #1905 from igchor/umf_hwloc_disable # Bump UMF version to allow disabling hwloc - set(UNIFIED_RUNTIME_TAG a985a81dc9ba8adfcc8b54e35ad287e97766fb3e) + set(UNIFIED_RUNTIME_TAG review/yang/dsan_nullpointer) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") # Due to the use of dependentloadflag and no installer for UMF and hwloc we need From 5a9a04a3ac0aa34918f3e8ed936aa090e5b70658 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 31 Jul 2024 23:47:32 -0700 Subject: [PATCH 3/8] fix libdevice --- libdevice/sanitizer_utils.cpp | 44 +++++++++++++++++------------------ 1 file changed, 21 insertions(+), 23 deletions(-) diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp index b606af36ae20b..b2819e5032a0f 100644 --- a/libdevice/sanitizer_utils.cpp +++ b/libdevice/sanitizer_utils.cpp @@ -103,8 +103,8 @@ enum ADDRESS_SPACE : uint32_t { namespace { -bool __asan_report_unknown_device(); -bool __asan_report_out_of_shadow_bounds(); +void __asan_report_unknown_device(); +void __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) { @@ -181,10 +181,11 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) { ((addr & (slm_size - 1)) >> ASAN_SHADOW_SCALE); if (shadow_ptr > shadow_offset_end) { - if (__asan_report_out_of_shadow_bounds()) { + if (__AsanDebug) { __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr, wg_lid, (uptr)shadow_offset); } + __asan_report_out_of_shadow_bounds(); return 0; } return shadow_ptr; @@ -214,10 +215,11 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) { ((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE); if (shadow_ptr > shadow_offset_end) { - if (__asan_report_out_of_shadow_bounds()) { + if (__AsanDebug) { __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, WG_LID, (uptr)shadow_offset); } + __asan_report_out_of_shadow_bounds(); return 0; } return shadow_ptr; @@ -244,10 +246,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) { } if (shadow_ptr > __AsanShadowMemoryGlobalEnd) { - if (__asan_report_out_of_shadow_bounds()) { + if (__AsanDebug) { __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr, (uptr)__AsanShadowMemoryGlobalStart); } + __asan_report_out_of_shadow_bounds(); return 0; } return shadow_ptr; @@ -280,10 +283,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) { ((addr & (SLM_SIZE - 1)) >> ASAN_SHADOW_SCALE); if (shadow_ptr > shadow_offset_end) { - if (__asan_report_out_of_shadow_bounds()) { + if (__AsanDebug) { __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr, wg_lid, (uptr)shadow_offset); } + __asan_report_out_of_shadow_bounds(); return 0; } return shadow_ptr; @@ -313,10 +317,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) { ((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE); if (shadow_ptr > shadow_offset_end) { - if (__asan_report_out_of_shadow_bounds()) { + if (__AsanDebug) { __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, WG_LID, (uptr)shadow_offset); } + __asan_report_out_of_shadow_bounds(); return 0; } return shadow_ptr; @@ -335,14 +340,13 @@ inline uptr MemToShadow(uptr addr, uint32_t as) { } else if (__DeviceType == DeviceType::GPU_DG2) { shadow_ptr = MemToShadow_DG2(addr, as); } else { - if (__asan_report_unknown_device() && __AsanDebug) { + if (__AsanDebug) { __spirv_ocl_printf(__asan_print_unsupport_device_type, (int)__DeviceType); } - return shadow_ptr; + __asan_report_unknown_device(); + return 0; } -// FIXME: OCL "O2" optimizer doesn't work well with following code -#if 0 if (__AsanDebug) { if (shadow_ptr) { if (as == ADDRESS_SPACE_PRIVATE) @@ -354,7 +358,6 @@ inline uptr MemToShadow(uptr addr, uint32_t as) { __spirv_ocl_printf(__asan_print_shadow_value2, addr, as, shadow_ptr); } } -#endif return shadow_ptr; } @@ -397,7 +400,7 @@ bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) { static __SYCL_CONSTANT__ const char __mem_sanitizer_report[] = "[kernel] SanitizerReport (ErrorType=%d, IsRecover=%d)\n"; -bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) { +void __asan_internal_report_save(DeviceSanitizerErrorType error_type) { const int Expected = ASAN_REPORT_NONE; int Desired = ASAN_REPORT_START; @@ -422,13 +425,11 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) { if (__AsanDebug) __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType, SanitizerReport.IsRecover); - return true; } __devicelib_exit(); - return false; } -bool __asan_internal_report_save( +void __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, @@ -505,10 +506,8 @@ bool __asan_internal_report_save( if (__AsanDebug) __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType, SanitizerReport.IsRecover); - return true; } __devicelib_exit(); - return false; } /// @@ -608,13 +607,12 @@ void __asan_report_misalign_error(uptr addr, uint32_t as, size_t size, memory_type, error_type, is_recover); } -bool __asan_report_unknown_device() { - return __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE); +void __asan_report_unknown_device() { + __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE); } -bool __asan_report_out_of_shadow_bounds() { - return __asan_internal_report_save( - DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS); +void __asan_report_out_of_shadow_bounds() { + __asan_internal_report_save(DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS); } /// From 195a800580ba1651c571a58342c512f1b80aa602 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 31 Jul 2024 23:49:06 -0700 Subject: [PATCH 4/8] fix test --- sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp index 73e37e0981ee0..ca1d5b45c7898 100644 --- a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp +++ b/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp @@ -20,7 +20,7 @@ int main() { }); // CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory // CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) - // CHECK: #0 {{.*}} {{.*null-pointer.cpp}}:[[@LINE-5]] + // CHECK: {{.*nullpointer.cpp}}:[[@LINE-5]] return 0; } From 7ee6dbb5cd77ab82f7f0cd893ae5246e9590a8f5 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 31 Jul 2024 23:51:03 -0700 Subject: [PATCH 5/8] fix test --- sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp index ca1d5b45c7898..b5d69cb5fc748 100644 --- a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp +++ b/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp @@ -1,6 +1,6 @@ // REQUIRES: linux -// RUN-NOT: %{build} %device_asan_flags -O0 -g -o %t -// RUN-NOT: %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O0 -g -o %t +// RUN: %{run} not %t 2>&1 | FileCheck %s // RUN: %{build} %device_asan_flags -O1 -g -o %t // RUN: %{run} not %t 2>&1 | FileCheck %s // RUN: %{build} %device_asan_flags -O2 -g -o %t @@ -11,7 +11,7 @@ int main() { sycl::queue Q; constexpr std::size_t N = 4; - int *array = 0; + int *array = nullptr; Q.submit([&](sycl::handler &h) { h.parallel_for( From 1792b64ea5875772883791da328a3d8d50f43f80 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Thu, 22 Aug 2024 08:19:00 +0200 Subject: [PATCH 6/8] enable gpu --- sycl/test-e2e/AddressSanitizer/lit.local.cfg | 3 --- 1 file changed, 3 deletions(-) diff --git a/sycl/test-e2e/AddressSanitizer/lit.local.cfg b/sycl/test-e2e/AddressSanitizer/lit.local.cfg index 8eb37fb1a7b43..7ffa3ee2056d7 100644 --- a/sycl/test-e2e/AddressSanitizer/lit.local.cfg +++ b/sycl/test-e2e/AddressSanitizer/lit.local.cfg @@ -9,6 +9,3 @@ config.substitutions.append( ) config.unsupported_features += ['cuda', 'hip'] - -# FIXME: Skip gen devices, waiting for gfx driver uplifting -config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-dg2', 'gpu-intel-pvc'] From 3ad24d72a6e63567c8aa421face96c3a98041f94 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Fri, 23 Aug 2024 03:23:03 +0200 Subject: [PATCH 7/8] enable dg2 only --- sycl/test-e2e/AddressSanitizer/lit.local.cfg | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/AddressSanitizer/lit.local.cfg b/sycl/test-e2e/AddressSanitizer/lit.local.cfg index 7ffa3ee2056d7..0f6650c8cf3fa 100644 --- a/sycl/test-e2e/AddressSanitizer/lit.local.cfg +++ b/sycl/test-e2e/AddressSanitizer/lit.local.cfg @@ -9,3 +9,6 @@ config.substitutions.append( ) config.unsupported_features += ['cuda', 'hip'] + +# FIXME: Skip some of gpu devices, waiting for gfx driver uplifting +config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-pvc'] From dd06c90ad62425be9e50bd012fdc5671585b5da4 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Fri, 13 Sep 2024 12:15:16 +0200 Subject: [PATCH 8/8] add tests --- .../{nullpointer.cpp => global_nullptr.cpp} | 10 ++++-- .../nullpointer/private_nullptr.cpp | 36 +++++++++++++++++++ 2 files changed, 44 insertions(+), 2 deletions(-) rename sycl/test-e2e/AddressSanitizer/nullpointer/{nullpointer.cpp => global_nullptr.cpp} (64%) create mode 100644 sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp similarity index 64% rename from sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp rename to sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp index b5d69cb5fc748..1904564b9e3be 100644 --- a/sycl/test-e2e/AddressSanitizer/nullpointer/nullpointer.cpp +++ b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp @@ -15,12 +15,18 @@ int main() { Q.submit([&](sycl::handler &h) { h.parallel_for( - sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { array[0] = 0; }); + sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { + auto private_array = + sycl::ext::oneapi::experimental::static_address_cast< + sycl::access::address_space::private_space, + sycl::access::decorated::no>(array); + private_array[0] = 0; + }); Q.wait(); }); // CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory // CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) - // CHECK: {{.*nullpointer.cpp}}:[[@LINE-5]] + // CHECK: {{.*global_nullptr.cpp}}:[[@LINE-5]] return 0; } diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp new file mode 100644 index 0000000000000..7d3455c43d4d8 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp @@ -0,0 +1,36 @@ +// REQUIRES: linux +// RUN: %{build} %device_asan_flags -O0 -g -o %t +// RUN: %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O1 -g -o %t +// RUN: %{run} not %t 2>&1 | FileCheck %s +// RUN: %{build} %device_asan_flags -O2 -g -o %t +// RUN: %{run} not %t 2>&1 | FileCheck %s + +// FIXME: There's an issue in gfx driver, so this test pending here. +// XFAIL: * + +#include +#include + +int main() { + sycl::queue Q; + constexpr std::size_t N = 4; + int *array = nullptr; + + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { + auto private_array = + sycl::ext::oneapi::experimental::static_address_cast< + sycl::access::address_space::private_space, + sycl::access::decorated::no>(array); + private_array[0] = 0; + }); + Q.wait(); + }); + // CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory + // CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0) + // CHECK: {{.*private_nullptr.cpp}}:[[@LINE-5]] + + return 0; +}