diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp index ea26fd9c4bb68..1074b6ac7e3cd 100644 --- a/libdevice/sanitizer_utils.cpp +++ b/libdevice/sanitizer_utils.cpp @@ -55,6 +55,8 @@ __spirv_GenericCastToPtrExplicit_ToPrivate(void *, int); extern SYCL_EXTERNAL __attribute__((convergent)) void __spirv_ControlBarrier(uint32_t Execution, uint32_t Memory, uint32_t Semantics); + +extern "C" SYCL_EXTERNAL void __devicelib_exit(); #endif // __USE_SPIR_BUILTIN__ static const __SYCL_CONSTANT__ char __asan_shadow_value_start[] = @@ -104,8 +106,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) { @@ -182,10 +184,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; @@ -215,10 +218,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; @@ -245,10 +249,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; @@ -281,10 +286,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; @@ -314,10 +320,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; @@ -336,14 +343,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) @@ -355,7 +361,6 @@ inline uptr MemToShadow(uptr addr, uint32_t as) { __spirv_ocl_printf(__asan_print_shadow_value2, addr, as, shadow_ptr); } } -#endif return shadow_ptr; } @@ -398,7 +403,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; @@ -423,12 +428,11 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) { if (__AsanDebug) __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType, SanitizerReport.IsRecover); - return true; } - return false; + __devicelib_exit(); } -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,9 +509,8 @@ bool __asan_internal_report_save( if (__AsanDebug) __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType, SanitizerReport.IsRecover); - return true; } - return false; + __devicelib_exit(); } /// @@ -575,6 +578,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; } @@ -604,13 +610,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); } /// diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 503db98e0ea09..a4c32b10be2de 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -117,13 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT) endfunction() set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 6298474e628889d3598b9416303a52e67a2b66aa - # Merge: 3cd6eaeb 4bb6a103 - # Author: Piotr Balcer - # Date: Wed Sep 18 09:20:05 2024 +0200 - # Merge pull request #2093 from lslusarczyk/memleak-fix - # fixed issue #1990, L0 leaks checker counts successful create/destroy only - set(UNIFIED_RUNTIME_TAG 6298474e628889d3598b9416303a52e67a2b66aa) + # commit 4517290650a9938537666e6409fb8e0db73ff4d8 + # Merge: 6298474e 3dbb7a2a + # Author: Omar Ahmed + # Date: Wed Sep 18 08:48:03 2024 +0100 + # Merge pull request #1914 from AllanZyne/review/yang/dsan_nullpointer + # [DeviceSanitizer] Support nullpointer detection + set(UNIFIED_RUNTIME_TAG 4517290650a9938537666e6409fb8e0db73ff4d8) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") # Due to the use of dependentloadflag and no installer for UMF and hwloc we need diff --git a/sycl/test-e2e/AddressSanitizer/lit.local.cfg b/sycl/test-e2e/AddressSanitizer/lit.local.cfg index 8eb37fb1a7b43..0f6650c8cf3fa 100644 --- a/sycl/test-e2e/AddressSanitizer/lit.local.cfg +++ b/sycl/test-e2e/AddressSanitizer/lit.local.cfg @@ -10,5 +10,5 @@ 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'] +# 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'] diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp new file mode 100644 index 0000000000000..1904564b9e3be --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp @@ -0,0 +1,32 @@ +// 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 + +#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: {{.*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; +}