Skip to content

Commit

Permalink
[DeviceSanitizer] Support nullpointer detection & enable GPU tests (#…
Browse files Browse the repository at this point in the history
…14891)

UR: oneapi-src/unified-runtime#1914

---------

Co-authored-by: omarahmed1111 <[email protected]>
  • Loading branch information
AllanZyne and omarahmed1111 authored Sep 19, 2024
1 parent a9b870b commit 0985116
Show file tree
Hide file tree
Showing 5 changed files with 105 additions and 32 deletions.
51 changes: 28 additions & 23 deletions libdevice/sanitizer_utils.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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[] =
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand All @@ -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;
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand All @@ -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)
Expand All @@ -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;
}
Expand Down Expand Up @@ -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;

Expand All @@ -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,
Expand Down Expand Up @@ -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();
}

///
Expand Down Expand Up @@ -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;
}
Expand Down Expand Up @@ -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);
}

///
Expand Down
14 changes: 7 additions & 7 deletions sycl/cmake/modules/FetchUnifiedRuntime.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -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 <piotr.balcer@intel.com>
# 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 <omar.ahmed@codeplay.com>
# 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
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/AddressSanitizer/lit.local.cfg
Original file line number Diff line number Diff line change
Expand Up @@ -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']
32 changes: 32 additions & 0 deletions sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
Original file line number Diff line number Diff line change
@@ -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 <sycl/detail/core.hpp>

int main() {
sycl::queue Q;
constexpr std::size_t N = 4;
int *array = nullptr;

Q.submit([&](sycl::handler &h) {
h.parallel_for<class MyKernel>(
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;
}
36 changes: 36 additions & 0 deletions sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp
Original file line number Diff line number Diff line change
@@ -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 <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/experimental/address_cast.hpp>

int main() {
sycl::queue Q;
constexpr std::size_t N = 4;
int *array = nullptr;

Q.submit([&](sycl::handler &h) {
h.parallel_for<class MyKernel>(
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;
}

0 comments on commit 0985116

Please sign in to comment.