Skip to content

Commit fd5126c

Browse files
AllanZyneAllanZyne
committed
init
1 parent 466de9c commit fd5126c

File tree

6 files changed

+85
-10
lines changed

6 files changed

+85
-10
lines changed

‎libdevice/sanitizer/asan_rtl.cpp‎

Lines changed: 36 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,9 @@ static const __SYCL_CONSTANT__ char __asan_print_shadow_value2[] =
4040
static __SYCL_CONSTANT__ const char __generic_to[] =
4141
"[kernel] %p(4) - %p(%d)\n";
4242

43+
static __SYCL_CONSTANT__ const char __asan_print_shadow_bound[] =
44+
"[kernel] addr: %p, shadow: %p, lower: %p, uppper: %p\n";
45+
4346
#define ASAN_REPORT_NONE 0
4447
#define ASAN_REPORT_START 1
4548
#define ASAN_REPORT_FINISH 2
@@ -67,6 +70,8 @@ struct DebugInfo {
6770

6871
void ReportUnknownDevice(const DebugInfo *debug);
6972
void PrintShadowMemory(uptr addr, uptr shadow_address, uint32_t as);
73+
void SaveReport(ErrorType error_type, MemoryType memory_type, bool is_recover,
74+
const DebugInfo *debug);
7075

7176
__SYCL_GLOBAL__ void *ToGlobal(void *ptr) {
7277
return __spirv_GenericCastToPtrExplicit_ToGlobal(ptr, 5);
@@ -115,11 +120,22 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
115120
launch_info->GlobalShadowOffset + (addr >> ASAN_SHADOW_SCALE);
116121
}
117122

123+
if (shadow_ptr < launch_info->GlobalShadowLowerBound ||
124+
shadow_ptr > launch_info->GlobalShadowUpperBound) {
125+
__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
126+
launch_info->GlobalShadowLowerBound,
127+
launch_info->GlobalShadowUpperBound);
128+
SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
129+
return 0;
130+
}
131+
118132
ASAN_DEBUG(
119133
const auto shadow_offset_end = launch_info->GlobalShadowOffsetEnd;
120134
if (shadow_ptr > shadow_offset_end) {
121135
__spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
122136
(uptr)launch_info->GlobalShadowOffset);
137+
SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::GLOBAL, false,
138+
debug);
123139
return 0;
124140
});
125141

@@ -141,6 +157,8 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
141157
if (shadow_ptr > shadow_offset_end) {
142158
__spirv_ocl_printf(__local_shadow_out_of_bound, addr,
143159
shadow_ptr, wid, (uptr)shadow_offset);
160+
SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::LOCAL,
161+
false, debug);
144162
return 0;
145163
});
146164
return shadow_ptr;
@@ -167,8 +185,10 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
167185
if (shadow_ptr > shadow_offset_end) {
168186
__spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, sid,
169187
private_base);
188+
SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::PRIVATE, false,
189+
debug);
170190
return 0;
171-
};
191+
}
172192

173193
return shadow_ptr;
174194
}
@@ -193,11 +213,22 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
193213
((addr & 0x7FFFFFFFFFFF) >> ASAN_SHADOW_SCALE);
194214
}
195215

216+
if (shadow_ptr < launch_info->GlobalShadowLowerBound ||
217+
shadow_ptr > launch_info->GlobalShadowUpperBound) {
218+
__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
219+
launch_info->GlobalShadowLowerBound,
220+
launch_info->GlobalShadowUpperBound);
221+
SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
222+
return 0;
223+
}
224+
196225
ASAN_DEBUG(
197226
const auto shadow_offset_end = launch_info->GlobalShadowOffsetEnd;
198227
if (shadow_ptr > shadow_offset_end) {
199228
__spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
200229
(uptr)launch_info->GlobalShadowOffset);
230+
SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::GLOBAL, false,
231+
debug);
201232
return 0;
202233
});
203234
return shadow_ptr;
@@ -218,6 +249,8 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
218249
if (shadow_ptr > shadow_offset_end) {
219250
__spirv_ocl_printf(__local_shadow_out_of_bound, addr,
220251
shadow_ptr, wid, (uptr)shadow_offset);
252+
SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::LOCAL,
253+
false, debug);
221254
return 0;
222255
});
223256
return shadow_ptr;
@@ -244,6 +277,8 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
244277
if (shadow_ptr > shadow_offset_end) {
245278
__spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, sid,
246279
private_base);
280+
SaveReport(ErrorType::OUT_OF_SHADOW_BOUNDS, MemoryType::PRIVATE, false,
281+
debug);
247282
return 0;
248283
};
249284

‎sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/arbitary_access.cpp‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
// REQUIRES: linux, cpu || (gpu && level_zero)
2+
// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t1.out
3+
// RUN: %{run} not %t1.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O1 -g -o %t2.out
5+
// RUN: %{run} not %t2.out 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O2 -g -o %t3.out
7+
// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
8+
9+
#include <sycl/detail/core.hpp>
10+
#include <sycl/usm.hpp>
11+
12+
void out_of_bounds_function() { *(int *)0xdeadbeef = 42; }
13+
// CHECK: out-of-bounds-access
14+
// CHECK-SAME: 0xdeadbeef
15+
// CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}}
16+
// CHECK: {{.*arbitary.cpp}}:[[@LINE-4]]
17+
18+
int main() {
19+
sycl::queue Q;
20+
21+
Q.submit([&](sycl::handler &h) {
22+
h.single_task<class MyKernel>([=]() { out_of_bounds_function(); });
23+
});
24+
Q.wait();
25+
26+
return 0;
27+
}

‎unified-runtime/source/loader/layers/sanitizer/asan/asan_libdevice.hpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -59,6 +59,9 @@ struct AsanRuntimeData {
5959
uintptr_t GlobalShadowOffset = 0;
6060
uintptr_t GlobalShadowOffsetEnd = 0;
6161

62+
uintptr_t GlobalShadowLowerBound = 0;
63+
uintptr_t GlobalShadowUpperBound = 0;
64+
6265
uintptr_t *PrivateBase = nullptr;
6366
uintptr_t PrivateShadowOffset = 0;
6467
uintptr_t PrivateShadowOffsetEnd = 0;

‎unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.cpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -226,6 +226,11 @@ ur_result_t ShadowMemoryGPU::EnqueuePoisonShadow(ur_queue_handle_t Queue,
226226
VirtualMemMaps[MappedPtr] = PhysicalMem;
227227
}
228228
}
229+
230+
ShadowLowerBound =
231+
std::min(ShadowLowerBound, RoundDownTo(ShadowBegin, PageSize));
232+
ShadowUpperBound =
233+
std::max(ShadowUpperBound, RoundUpTo(ShadowEnd, PageSize));
229234
}
230235

231236
auto URes = EnqueueUSMSet(Queue, (void *)ShadowBegin, Value,

‎unified-runtime/source/loader/layers/sanitizer/asan/asan_shadow.hpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -64,8 +64,10 @@ struct ShadowMemory {
6464
ur_device_handle_t Device{};
6565

6666
uptr ShadowBegin = 0;
67-
6867
uptr ShadowEnd = 0;
68+
69+
uptr ShadowLowerBound = 0xffff'ffff'ffff'ffff;
70+
uptr ShadowUpperBound = 0;
6971
};
7072

7173
struct ShadowMemoryCPU final : public ShadowMemory {

‎unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_libdevice.hpp‎

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -68,16 +68,19 @@ inline const char *ToString(ErrorType ErrorType) {
6868
}
6969
}
7070

71+
// clang-format off
7172
enum class MemoryType : int32_t {
72-
UNKNOWN,
73-
USM_DEVICE,
74-
USM_HOST,
75-
USM_SHARED,
76-
LOCAL,
77-
PRIVATE,
78-
MEM_BUFFER,
79-
DEVICE_GLOBAL,
73+
UNKNOWN = 0x000000'00,
74+
GLOBAL = 0x000001'00,
75+
USM_DEVICE = 0x000001'01,
76+
USM_HOST = 0x000001'02,
77+
USM_SHARED = 0x000001'03,
78+
MEM_BUFFER = 0x000001'04,
79+
DEVICE_GLOBAL = 0x000001'05,
80+
LOCAL = 0x000002'00,
81+
PRIVATE = 0x000004'00,
8082
};
83+
// clang-format on
8184

8285
inline const char *ToString(MemoryType MemoryType) {
8386
switch (MemoryType) {

0 commit comments

Comments
 (0)