Skip to content

Commit 199451d

Browse files
AllanZyneAllanZyne
committed
wip
1 parent fd5126c commit 199451d

File tree

4 files changed

+45
-25
lines changed

4 files changed

+45
-25
lines changed

‎libdevice/sanitizer/asan_rtl.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -122,9 +122,9 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
122122

123123
if (shadow_ptr < launch_info->GlobalShadowLowerBound ||
124124
shadow_ptr > launch_info->GlobalShadowUpperBound) {
125-
__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
126-
launch_info->GlobalShadowLowerBound,
127-
launch_info->GlobalShadowUpperBound);
125+
ASAN_DEBUG(__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
126+
launch_info->GlobalShadowLowerBound,
127+
launch_info->GlobalShadowUpperBound));
128128
SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
129129
return 0;
130130
}
@@ -215,9 +215,9 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
215215

216216
if (shadow_ptr < launch_info->GlobalShadowLowerBound ||
217217
shadow_ptr > launch_info->GlobalShadowUpperBound) {
218-
__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
219-
launch_info->GlobalShadowLowerBound,
220-
launch_info->GlobalShadowUpperBound);
218+
ASAN_DEBUG(__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
219+
launch_info->GlobalShadowLowerBound,
220+
launch_info->GlobalShadowUpperBound));
221221
SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
222222
return 0;
223223
}

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

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,7 @@
1-
// REQUIRES: linux, cpu || (gpu && level_zero)
2-
// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O0 -g -o %t1.out
1+
// REQUIRES: linux, gpu && level_zero
2+
// RUN: %{build} %device_asan_flags -O0 -g -o %t1.out
33
// 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
4+
// RUN: %{build} %device_asan_flags -O2 -g -o %t3.out
75
// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
86

97
#include <sycl/detail/core.hpp>
@@ -13,7 +11,7 @@ void out_of_bounds_function() { *(int *)0xdeadbeef = 42; }
1311
// CHECK: out-of-bounds-access
1412
// CHECK-SAME: 0xdeadbeef
1513
// CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}}
16-
// CHECK: {{.*arbitary.cpp}}:[[@LINE-4]]
14+
// CHECK: {{.*arbitary_access.cpp}}:[[@LINE-4]]
1715

1816
int main() {
1917
sycl::queue Q;

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

Lines changed: 32 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -267,11 +267,23 @@ ur_result_t AsanInterceptor::preLaunchKernel(ur_kernel_handle_t Kernel,
267267

268268
ur_queue_handle_t InternalQueue = ContextInfo->getInternalQueue(Device);
269269

270+
// To get right shadow boundary, shadow memory should be updated before
271+
// prepareLaunch
272+
{
273+
// Force to allocate membuffer before prepareLaunch
274+
auto &KernelInfo = getOrCreateKernelInfo(Kernel);
275+
std::shared_lock<ur_shared_mutex> Guard(KernelInfo.Mutex);
276+
for (const auto &[ArgIndex, MemBuffer] : KernelInfo.BufferArgs) {
277+
char *ArgPointer = nullptr;
278+
UR_CALL(MemBuffer->getHandle(DeviceInfo->Handle, ArgPointer));
279+
(void)ArgPointer;
280+
}
281+
}
282+
UR_CALL(updateShadowMemory(ContextInfo, DeviceInfo, InternalQueue));
283+
270284
UR_CALL(prepareLaunch(ContextInfo, DeviceInfo, InternalQueue, Kernel,
271285
LaunchInfo));
272286

273-
UR_CALL(updateShadowMemory(ContextInfo, DeviceInfo, InternalQueue));
274-
275287
UR_CALL(getContext()->urDdiTable.Queue.pfnFinish(InternalQueue));
276288

277289
return UR_RESULT_SUCCESS;
@@ -825,6 +837,10 @@ ur_result_t AsanInterceptor::prepareLaunch(
825837
// Prepare asan runtime data
826838
LaunchInfo.Data.Host.GlobalShadowOffset = DeviceInfo->Shadow->ShadowBegin;
827839
LaunchInfo.Data.Host.GlobalShadowOffsetEnd = DeviceInfo->Shadow->ShadowEnd;
840+
LaunchInfo.Data.Host.GlobalShadowLowerBound =
841+
DeviceInfo->Shadow->ShadowLowerBound;
842+
LaunchInfo.Data.Host.GlobalShadowUpperBound =
843+
DeviceInfo->Shadow->ShadowUpperBound;
828844
LaunchInfo.Data.Host.Debug = getContext()->Options.Debug ? 1 : 0;
829845

830846
// Write shadow memory offset for local memory
@@ -884,16 +900,20 @@ ur_result_t AsanInterceptor::prepareLaunch(
884900
// sync asan runtime data to device side
885901
UR_CALL(LaunchInfo.Data.syncToDevice(Queue));
886902

887-
UR_LOG_L(getContext()->logger, INFO,
888-
"LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateBase={}, "
889-
"PrivateShadow={}, LocalArgs={}, NumLocalArgs={}, Debug={})",
890-
(void *)LaunchInfo.Data.getDevicePtr(),
891-
(void *)LaunchInfo.Data.Host.GlobalShadowOffset,
892-
(void *)LaunchInfo.Data.Host.LocalShadowOffset,
893-
(void *)LaunchInfo.Data.Host.PrivateBase,
894-
(void *)LaunchInfo.Data.Host.PrivateShadowOffset,
895-
(void *)LaunchInfo.Data.Host.LocalArgs,
896-
LaunchInfo.Data.Host.NumLocalArgs, LaunchInfo.Data.Host.Debug);
903+
UR_LOG_L(
904+
getContext()->logger, INFO,
905+
"LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateBase={}, "
906+
"PrivateShadow={}, GlobalShadowLowerBound={}, GlobalShadowUpperBound={}, "
907+
"LocalArgs={}, NumLocalArgs={}, Debug={})",
908+
(void *)LaunchInfo.Data.getDevicePtr(),
909+
(void *)LaunchInfo.Data.Host.GlobalShadowOffset,
910+
(void *)LaunchInfo.Data.Host.LocalShadowOffset,
911+
(void *)LaunchInfo.Data.Host.PrivateBase,
912+
(void *)LaunchInfo.Data.Host.PrivateShadowOffset,
913+
(void *)LaunchInfo.Data.Host.GlobalShadowLowerBound,
914+
(void *)LaunchInfo.Data.Host.GlobalShadowUpperBound,
915+
(void *)LaunchInfo.Data.Host.LocalArgs, LaunchInfo.Data.Host.NumLocalArgs,
916+
LaunchInfo.Data.Host.Debug);
897917

898918
return UR_RESULT_SUCCESS;
899919
}

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

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ inline const char *ToString(ErrorType ErrorType) {
6969
}
7070

7171
// clang-format off
72-
enum class MemoryType : int32_t {
72+
enum MemoryType : uint32_t {
7373
UNKNOWN = 0x000000'00,
7474
GLOBAL = 0x000001'00,
7575
USM_DEVICE = 0x000001'01,
@@ -79,6 +79,8 @@ enum class MemoryType : int32_t {
7979
DEVICE_GLOBAL = 0x000001'05,
8080
LOCAL = 0x000002'00,
8181
PRIVATE = 0x000004'00,
82+
CONSTANT = 0x000008'00,
83+
GENERIC = 0x000010'00,
8284
};
8385
// clang-format on
8486

0 commit comments

Comments
 (0)