Skip to content

Commit 0985116

Browse files
AllanZyneomarahmed1111
AllanZyne
and
omarahmed1111
authored
[DeviceSanitizer] Support nullpointer detection & enable GPU tests (#14891)
UR: oneapi-src/unified-runtime#1914 --------- Co-authored-by: omarahmed1111 <[email protected]>
1 parent a9b870b commit 0985116

File tree

5 files changed

+105
-32
lines changed

5 files changed

+105
-32
lines changed

‎libdevice/sanitizer_utils.cpp‎

Lines changed: 28 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,8 @@ __spirv_GenericCastToPtrExplicit_ToPrivate(void *, int);
5555

5656
extern SYCL_EXTERNAL __attribute__((convergent)) void
5757
__spirv_ControlBarrier(uint32_t Execution, uint32_t Memory, uint32_t Semantics);
58+
59+
extern "C" SYCL_EXTERNAL void __devicelib_exit();
5860
#endif // __USE_SPIR_BUILTIN__
5961

6062
static const __SYCL_CONSTANT__ char __asan_shadow_value_start[] =
@@ -104,8 +106,8 @@ enum ADDRESS_SPACE : uint32_t {
104106

105107
namespace {
106108

107-
bool __asan_report_unknown_device();
108-
bool __asan_report_out_of_shadow_bounds();
109+
void __asan_report_unknown_device();
110+
void __asan_report_out_of_shadow_bounds();
109111
void __asan_print_shadow_memory(uptr addr, uptr shadow_address, uint32_t as);
110112

111113
__SYCL_GLOBAL__ void *ToGlobal(void *ptr) {
@@ -182,10 +184,11 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
182184
((addr & (slm_size - 1)) >> ASAN_SHADOW_SCALE);
183185

184186
if (shadow_ptr > shadow_offset_end) {
185-
if (__asan_report_out_of_shadow_bounds()) {
187+
if (__AsanDebug) {
186188
__spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
187189
wg_lid, (uptr)shadow_offset);
188190
}
191+
__asan_report_out_of_shadow_bounds();
189192
return 0;
190193
}
191194
return shadow_ptr;
@@ -215,10 +218,11 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
215218
((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE);
216219

217220
if (shadow_ptr > shadow_offset_end) {
218-
if (__asan_report_out_of_shadow_bounds()) {
221+
if (__AsanDebug) {
219222
__spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr,
220223
WG_LID, (uptr)shadow_offset);
221224
}
225+
__asan_report_out_of_shadow_bounds();
222226
return 0;
223227
}
224228
return shadow_ptr;
@@ -245,10 +249,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
245249
}
246250

247251
if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
248-
if (__asan_report_out_of_shadow_bounds()) {
252+
if (__AsanDebug) {
249253
__spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
250254
(uptr)__AsanShadowMemoryGlobalStart);
251255
}
256+
__asan_report_out_of_shadow_bounds();
252257
return 0;
253258
}
254259
return shadow_ptr;
@@ -281,10 +286,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
281286
((addr & (SLM_SIZE - 1)) >> ASAN_SHADOW_SCALE);
282287

283288
if (shadow_ptr > shadow_offset_end) {
284-
if (__asan_report_out_of_shadow_bounds()) {
289+
if (__AsanDebug) {
285290
__spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
286291
wg_lid, (uptr)shadow_offset);
287292
}
293+
__asan_report_out_of_shadow_bounds();
288294
return 0;
289295
}
290296
return shadow_ptr;
@@ -314,10 +320,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
314320
((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE);
315321

316322
if (shadow_ptr > shadow_offset_end) {
317-
if (__asan_report_out_of_shadow_bounds()) {
323+
if (__AsanDebug) {
318324
__spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr,
319325
WG_LID, (uptr)shadow_offset);
320326
}
327+
__asan_report_out_of_shadow_bounds();
321328
return 0;
322329
}
323330
return shadow_ptr;
@@ -336,14 +343,13 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
336343
} else if (__DeviceType == DeviceType::GPU_DG2) {
337344
shadow_ptr = MemToShadow_DG2(addr, as);
338345
} else {
339-
if (__asan_report_unknown_device() && __AsanDebug) {
346+
if (__AsanDebug) {
340347
__spirv_ocl_printf(__asan_print_unsupport_device_type, (int)__DeviceType);
341348
}
342-
return shadow_ptr;
349+
__asan_report_unknown_device();
350+
return 0;
343351
}
344352

345-
// FIXME: OCL "O2" optimizer doesn't work well with following code
346-
#if 0
347353
if (__AsanDebug) {
348354
if (shadow_ptr) {
349355
if (as == ADDRESS_SPACE_PRIVATE)
@@ -355,7 +361,6 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
355361
__spirv_ocl_printf(__asan_print_shadow_value2, addr, as, shadow_ptr);
356362
}
357363
}
358-
#endif
359364

360365
return shadow_ptr;
361366
}
@@ -398,7 +403,7 @@ bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) {
398403
static __SYCL_CONSTANT__ const char __mem_sanitizer_report[] =
399404
"[kernel] SanitizerReport (ErrorType=%d, IsRecover=%d)\n";
400405

401-
bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
406+
void __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
402407
const int Expected = ASAN_REPORT_NONE;
403408
int Desired = ASAN_REPORT_START;
404409

@@ -423,12 +428,11 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
423428
if (__AsanDebug)
424429
__spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType,
425430
SanitizerReport.IsRecover);
426-
return true;
427431
}
428-
return false;
432+
__devicelib_exit();
429433
}
430434

431-
bool __asan_internal_report_save(
435+
void __asan_internal_report_save(
432436
uptr ptr, uint32_t as, const char __SYCL_CONSTANT__ *file, uint32_t line,
433437
const char __SYCL_CONSTANT__ *func, bool is_write, uint32_t access_size,
434438
DeviceSanitizerMemoryType memory_type, DeviceSanitizerErrorType error_type,
@@ -505,9 +509,8 @@ bool __asan_internal_report_save(
505509
if (__AsanDebug)
506510
__spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType,
507511
SanitizerReport.IsRecover);
508-
return true;
509512
}
510-
return false;
513+
__devicelib_exit();
511514
}
512515

513516
///
@@ -575,6 +578,9 @@ void __asan_report_access_error(uptr addr, uint32_t as, size_t size,
575578
case kUsmSharedDeallocatedMagic:
576579
error_type = DeviceSanitizerErrorType::USE_AFTER_FREE;
577580
break;
581+
case kNullPointerRedzoneMagic:
582+
error_type = DeviceSanitizerErrorType::NULL_POINTER;
583+
break;
578584
default:
579585
error_type = DeviceSanitizerErrorType::UNKNOWN;
580586
}
@@ -604,13 +610,12 @@ void __asan_report_misalign_error(uptr addr, uint32_t as, size_t size,
604610
memory_type, error_type, is_recover);
605611
}
606612

607-
bool __asan_report_unknown_device() {
608-
return __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE);
613+
void __asan_report_unknown_device() {
614+
__asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE);
609615
}
610616

611-
bool __asan_report_out_of_shadow_bounds() {
612-
return __asan_internal_report_save(
613-
DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS);
617+
void __asan_report_out_of_shadow_bounds() {
618+
__asan_internal_report_save(DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS);
614619
}
615620

616621
///

‎sycl/cmake/modules/FetchUnifiedRuntime.cmake‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -117,13 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT)
117117
endfunction()
118118

119119
set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
120-
# commit 6298474e628889d3598b9416303a52e67a2b66aa
121-
# Merge: 3cd6eaeb 4bb6a103
122-
# Author: Piotr Balcer <piotr.balcer@intel.com>
123-
# Date: Wed Sep 18 09:20:05 2024 +0200
124-
# Merge pull request #2093 from lslusarczyk/memleak-fix
125-
# fixed issue #1990, L0 leaks checker counts successful create/destroy only
126-
set(UNIFIED_RUNTIME_TAG 6298474e628889d3598b9416303a52e67a2b66aa)
120+
# commit 4517290650a9938537666e6409fb8e0db73ff4d8
121+
# Merge: 6298474e 3dbb7a2a
122+
# Author: Omar Ahmed <omar.ahmed@codeplay.com>
123+
# Date: Wed Sep 18 08:48:03 2024 +0100
124+
# Merge pull request #1914 from AllanZyne/review/yang/dsan_nullpointer
125+
# [DeviceSanitizer] Support nullpointer detection
126+
set(UNIFIED_RUNTIME_TAG 4517290650a9938537666e6409fb8e0db73ff4d8)
127127

128128
set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
129129
# Due to the use of dependentloadflag and no installer for UMF and hwloc we need

‎sycl/test-e2e/AddressSanitizer/lit.local.cfg‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,5 +10,5 @@ config.substitutions.append(
1010

1111
config.unsupported_features += ['cuda', 'hip']
1212

13-
# FIXME: Skip gen devices, waiting for gfx driver uplifting
14-
config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-dg2', 'gpu-intel-pvc']
13+
# FIXME: Skip some of gpu devices, waiting for gfx driver uplifting
14+
config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-pvc']

‎sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
// REQUIRES: linux
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t
3+
// RUN: %{run} not %t 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t
5+
// RUN: %{run} not %t 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_asan_flags -O2 -g -o %t
7+
// RUN: %{run} not %t 2>&1 | FileCheck %s
8+
9+
#include <sycl/detail/core.hpp>
10+
11+
int main() {
12+
sycl::queue Q;
13+
constexpr std::size_t N = 4;
14+
int *array = nullptr;
15+
16+
Q.submit([&](sycl::handler &h) {
17+
h.parallel_for<class MyKernel>(
18+
sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) {
19+
auto private_array =
20+
sycl::ext::oneapi::experimental::static_address_cast<
21+
sycl::access::address_space::private_space,
22+
sycl::access::decorated::no>(array);
23+
private_array[0] = 0;
24+
});
25+
Q.wait();
26+
});
27+
// CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory
28+
// CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0)
29+
// CHECK: {{.*global_nullptr.cpp}}:[[@LINE-5]]
30+
31+
return 0;
32+
}

‎sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp‎

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// REQUIRES: linux
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t
3+
// RUN: %{run} not %t 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -O1 -g -o %t
5+
// RUN: %{run} not %t 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_asan_flags -O2 -g -o %t
7+
// RUN: %{run} not %t 2>&1 | FileCheck %s
8+
9+
// FIXME: There's an issue in gfx driver, so this test pending here.
10+
// XFAIL: *
11+
12+
#include <sycl/detail/core.hpp>
13+
#include <sycl/ext/oneapi/experimental/address_cast.hpp>
14+
15+
int main() {
16+
sycl::queue Q;
17+
constexpr std::size_t N = 4;
18+
int *array = nullptr;
19+
20+
Q.submit([&](sycl::handler &h) {
21+
h.parallel_for<class MyKernel>(
22+
sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) {
23+
auto private_array =
24+
sycl::ext::oneapi::experimental::static_address_cast<
25+
sycl::access::address_space::private_space,
26+
sycl::access::decorated::no>(array);
27+
private_array[0] = 0;
28+
});
29+
Q.wait();
30+
});
31+
// CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory
32+
// CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0)
33+
// CHECK: {{.*private_nullptr.cpp}}:[[@LINE-5]]
34+
35+
return 0;
36+
}

0 commit comments

Comments
 (0)