From 177826a9ddc365a11f8f70e44c9359f8d8d66041 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 5 Mar 2025 07:01:02 +0100 Subject: [PATCH 01/14] wip --- libdevice/sanitizer/msan_rtl.cpp | 35 ++++++++++++++---- .../Instrumentation/MemorySanitizer.cpp | 4 +-- .../private/single_private.cpp | 32 +++++++++++++++++ .../sanitizer/msan/msan_interceptor.cpp | 25 +++++++++++-- .../layers/sanitizer/msan/msan_libdevice.hpp | 3 ++ .../layers/sanitizer/msan/msan_shadow.cpp | 36 +++++++++++++++++++ .../layers/sanitizer/msan/msan_shadow.hpp | 14 ++++++++ 7 files changed, 138 insertions(+), 11 deletions(-) create mode 100644 sycl/test-e2e/MemorySanitizer/private/single_private.cpp diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp index 0abb77aa60e98..0691978815da0 100644 --- a/libdevice/sanitizer/msan_rtl.cpp +++ b/libdevice/sanitizer/msan_rtl.cpp @@ -48,6 +48,10 @@ const __SYCL_CONSTANT__ char __msan_print_func_beg[] = const __SYCL_CONSTANT__ char __msan_print_func_end[] = "[kernel] ===== %s() end\n"; +const __SYCL_CONSTANT__ char __msan_print_private_shadow_out_of_bound[] = + "[kernel] Private shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: " + "%p)\n"; + } // namespace #if defined(__SPIR__) || defined(__SPIRV__) @@ -61,6 +65,13 @@ const __SYCL_CONSTANT__ char __msan_print_func_end[] = namespace { +inline size_t workgroup_linear_id() { + return __spirv_BuiltInWorkgroupId.x * __spirv_BuiltInNumWorkgroups.y * + __spirv_BuiltInNumWorkgroups.z + + __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z + + __spirv_BuiltInWorkgroupId.z; +} + inline void ConvertGenericPointer(uptr &addr, uint32_t &as) { auto old = addr; if ((addr = (uptr)ToPrivate((void *)old))) { @@ -181,18 +192,28 @@ inline uptr __msan_get_shadow_pvc(uptr addr, uint32_t as) { } else if (as == ADDRESS_SPACE_LOCAL) { // The size of SLM is 128KB on PVC constexpr unsigned SLM_SIZE = 128 * 1024; - // work-group linear id - const auto wg_lid = - __spirv_BuiltInWorkgroupId.x * __spirv_BuiltInNumWorkgroups.y * - __spirv_BuiltInNumWorkgroups.z + - __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z + - __spirv_BuiltInWorkgroupId.z; - + const auto wg_lid = workgroup_linear_id(); const auto shadow_offset = GetMsanLaunchInfo->LocalShadowOffset; if (shadow_offset != 0) { return shadow_offset + (wg_lid * SLM_SIZE) + (addr & (SLM_SIZE - 1)); } + } else if (as == ADDRESS_SPACE_PRIVATE) { + const auto wg_lid = workgroup_linear_id(); + const auto shadow_offset = GetMsanLaunchInfo->PrivateShadowOffset; + + if (shadow_offset != 0) { + uptr shadow_ptr = shadow_offset + (wg_lid * MSAN_PRIVATE_SIZE) + + (addr & (MSAN_PRIVATE_SIZE - 1)); + MSAN_DEBUG( + const auto shadow_offset_end = GetMsanLaunchInfo->PrivateShadowOffsetEnd; + if (shadow_ptr > shadow_offset_end) { + __spirv_ocl_printf(__msan_print_private_shadow_out_of_bound, addr, + shadow_ptr, wg_lid, (uptr)shadow_offset); + return 0; + }); + return shadow_ptr; + } } return GetMsanLaunchInfo->CleanShadow; diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 0d054e6cc8a8a..f991a893b7946 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -390,7 +390,7 @@ static cl::opt ClSpirOffloadLocals("msan-spir-locals", static cl::opt ClSpirOffloadPrivates("msan-spir-privates", cl::desc("instrument private pointer"), cl::Hidden, - cl::init(false)); + cl::init(true)); const char kMsanModuleCtorName[] = "msan.module_ctor"; const char kMsanInitName[] = "__msan_init"; @@ -1607,7 +1607,7 @@ static void setNoSanitizedMetadataSPIR(Instruction &I) { Addr = XCHG->getPointerOperand(); else if (const auto *ASC = dyn_cast(&I)) Addr = ASC->getPointerOperand(); - else if (isa(&I)) + else if (isa(&I) && !ClSpirOffloadPrivates) I.setNoSanitizeMetadata(); else if (const auto *CI = dyn_cast(&I)) { auto *Func = CI->getCalledFunction(); diff --git a/sycl/test-e2e/MemorySanitizer/private/single_private.cpp b/sycl/test-e2e/MemorySanitizer/private/single_private.cpp new file mode 100644 index 0000000000000..f4f8784a095b6 --- /dev/null +++ b/sycl/test-e2e/MemorySanitizer/private/single_private.cpp @@ -0,0 +1,32 @@ +// REQUIRES: linux, cpu || (gpu && level_zero) +// RUN: %{build} %device_msan_flags -O0 -g -o %t1.out +// RUN: %{run} not %t1.out 2>&1 | FileCheck %s +// RUN: %{build} %device_msan_flags -O1 -g -o %t2.out +// RUN: %{run} not %t2.out 2>&1 | FileCheck %s +// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out +// RUN: %{run} not %t3.out 2>&1 | FileCheck %s + +#include +#include + +__attribute__((noinline)) int check(int p) { return p; } +__attribute__((noinline)) int foo(int *p) { return check(*p); } +// CHECK-NOT: [kernel] +// CHECK: DeviceSanitizer: use-of-uninitialized-value +// CHECK: #0 {{foo.*}} {{.*single_private.cpp}}:[[@LINE-3]] + +int main() { + sycl::queue Q; + auto *array = sycl::malloc_device(1, Q); + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { + int p[4]; + *array += foo(p); + }); + }); + Q.wait(); + + sycl::free(array, Q); + return 0; +} diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp index 646fd813c1594..c25a3345467a5 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp @@ -524,11 +524,32 @@ ur_result_t MsanInterceptor::prepareLaunch( } } + // Write shadow memory offset for private memory + if (KernelInfo.IsCheckPrivates) { + if (DeviceInfo->Shadow->AllocPrivateShadow( + Queue, NumWG, LaunchInfo.Data->PrivateShadowOffset, + LaunchInfo.Data->PrivateShadowOffsetEnd) != UR_RESULT_SUCCESS) { + getContext()->logger.warning( + "Failed to allocate shadow memory for private " + "memory, maybe the number of workgroup ({}) is too " + "large", + NumWG); + getContext()->logger.warning( + "Skip checking private memory of kernel <{}>", GetKernelName(Kernel)); + } else { + getContext()->logger.info( + "ShadowMemory(Private, WorkGroup={}, {} - {})", NumWG, + (void *)LaunchInfo.Data->PrivateShadowOffset, + (void *)LaunchInfo.Data->PrivateShadowOffsetEnd); + } + } + getContext()->logger.info( - "LaunchInfo {} (GlobalShadow={}, LocalShadow={}, CleanShadow={}, " + "LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateShadow={}, " + "CleanShadow={}, " "Device={}, Debug={})", (void *)LaunchInfo.Data, (void *)LaunchInfo.Data->GlobalShadowOffset, - (void *)LaunchInfo.Data->LocalShadowOffset, + (void *)LaunchInfo.Data->LocalShadowOffset, (void *)LaunchInfo.Data->PrivateShadowOffset, (void *)LaunchInfo.Data->CleanShadow, ToString(LaunchInfo.Data->DeviceTy), LaunchInfo.Data->Debug); diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp index 2ad9c917b56c4..b0fb94a27067b 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp @@ -51,6 +51,9 @@ struct MsanLaunchInfo { uintptr_t LocalShadowOffset = 0; uintptr_t LocalShadowOffsetEnd = 0; + uintptr_t PrivateShadowOffset = 0; + uintptr_t PrivateShadowOffsetEnd = 0; + uintptr_t CleanShadow = 0; DeviceType DeviceTy = DeviceType::UNKNOWN; diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.cpp index 19d7f4fa08536..70bba4b565796 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.cpp @@ -333,6 +333,42 @@ ur_result_t MsanShadowMemoryGPU::AllocLocalShadow(ur_queue_handle_t Queue, return UR_RESULT_SUCCESS; } +ur_result_t MsanShadowMemoryGPU::AllocPrivateShadow(ur_queue_handle_t Queue, + uint32_t NumWG, uptr &Begin, + uptr &End) { + const size_t RequiredShadowSize = NumWG * MSAN_PRIVATE_SIZE; + static size_t LastAllocedSize = 0; + if (RequiredShadowSize > LastAllocedSize) { + auto ContextInfo = getMsanInterceptor()->getContextInfo(Context); + if (PrivateShadowOffset) { + UR_CALL(getContext()->urDdiTable.USM.pfnFree( + Context, (void *)PrivateShadowOffset)); + PrivateShadowOffset = 0; + LastAllocedSize = 0; + } + + UR_CALL(getContext()->urDdiTable.USM.pfnDeviceAlloc( + Context, Device, nullptr, nullptr, RequiredShadowSize, + (void **)&PrivateShadowOffset)); + + // Initialize shadow memory + ur_result_t URes = EnqueueUSMBlockingSet(Queue, (void *)PrivateShadowOffset, + 0, RequiredShadowSize); + if (URes != UR_RESULT_SUCCESS) { + UR_CALL(getContext()->urDdiTable.USM.pfnFree( + Context, (void *)PrivateShadowOffset)); + PrivateShadowOffset = 0; + LastAllocedSize = 0; + } + + LastAllocedSize = RequiredShadowSize; + } + + Begin = PrivateShadowOffset; + End = PrivateShadowOffset + RequiredShadowSize - 1; + return UR_RESULT_SUCCESS; +} + uptr MsanShadowMemoryPVC::MemToShadow(uptr Ptr) { assert(MsanShadowMemoryPVC::IsDeviceUSM(Ptr) && "Ptr must be device USM"); if (Ptr < ShadowBegin) { diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp index c737ee249182b..4b25a8709d45f 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp @@ -46,6 +46,10 @@ struct MsanShadowMemory { virtual ur_result_t AllocLocalShadow(ur_queue_handle_t Queue, uint32_t NumWG, uptr &Begin, uptr &End) = 0; + virtual ur_result_t AllocPrivateShadow(ur_queue_handle_t Queue, + uint32_t NumWG, uptr &Begin, + uptr &End) = 0; + ur_context_handle_t Context{}; ur_device_handle_t Device{}; @@ -93,6 +97,13 @@ struct MsanShadowMemoryCPU final : public MsanShadowMemory { End = ShadowEnd; return UR_RESULT_SUCCESS; } + + ur_result_t AllocPrivateShadow(ur_queue_handle_t, uint32_t, uptr &Begin, + uptr &End) override { + Begin = ShadowBegin; + End = ShadowEnd; + return UR_RESULT_SUCCESS; + } }; struct MsanShadowMemoryGPU : public MsanShadowMemory { @@ -114,6 +125,9 @@ struct MsanShadowMemoryGPU : public MsanShadowMemory { ur_result_t AllocLocalShadow(ur_queue_handle_t Queue, uint32_t NumWG, uptr &Begin, uptr &End) override final; + ur_result_t AllocPrivateShadow(ur_queue_handle_t Queue, uint32_t NumWG, + uptr &Begin, uptr &End) override final; + virtual size_t GetShadowSize() = 0; virtual uptr GetStartAddress() { return 0; } From fc73783a989b2f205748b6bf9fde8eef4795bec0 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 5 Mar 2025 07:19:01 +0100 Subject: [PATCH 02/14] add lit test --- .../MemorySanitizer/SPIRV/check_unsupported_access.ll | 2 +- .../SPIRV/instrument_static_local_mem.ll | 2 +- .../loader/layers/sanitizer/msan/msan_interceptor.cpp | 10 +++++----- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_unsupported_access.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_unsupported_access.ll index f679e7a43cf4d..3ca2ab9542ea8 100644 --- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_unsupported_access.ll +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_unsupported_access.ll @@ -1,4 +1,4 @@ -; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -S | FileCheck %s +; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -msan-spir-privates=0 -S | FileCheck %s target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" target triple = "spir64-unknown-unknown" diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_static_local_mem.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_static_local_mem.ll index 6130edfde9de7..15cf8a1584399 100644 --- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_static_local_mem.ll +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_static_local_mem.ll @@ -1,4 +1,4 @@ -; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -msan-spir-locals=1 -S | FileCheck %s +; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -msan-spir-locals=1 -msan-spir-privates=0 -S | FileCheck %s target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" target triple = "spir64-unknown-unknown" diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp index c25a3345467a5..0a3c098b51341 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp @@ -517,10 +517,10 @@ ur_result_t MsanInterceptor::prepareLaunch( getContext()->logger.warning("Skip checking local memory of kernel <{}> ", GetKernelName(Kernel)); } else { - getContext()->logger.info("ShadowMemory(Local, WorkGroup={}, {} - {})", - NumWG, - (void *)LaunchInfo.Data->LocalShadowOffset, - (void *)LaunchInfo.Data->LocalShadowOffsetEnd); + getContext()->logger.debug("ShadowMemory(Local, WorkGroup={}, {} - {})", + NumWG, + (void *)LaunchInfo.Data->LocalShadowOffset, + (void *)LaunchInfo.Data->LocalShadowOffsetEnd); } } @@ -537,7 +537,7 @@ ur_result_t MsanInterceptor::prepareLaunch( getContext()->logger.warning( "Skip checking private memory of kernel <{}>", GetKernelName(Kernel)); } else { - getContext()->logger.info( + getContext()->logger.debug( "ShadowMemory(Private, WorkGroup={}, {} - {})", NumWG, (void *)LaunchInfo.Data->PrivateShadowOffset, (void *)LaunchInfo.Data->PrivateShadowOffsetEnd); From 74f6ae3d374726b393204eaab9723f1787416429 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 5 Mar 2025 07:19:36 +0100 Subject: [PATCH 03/14] add lit test --- .../SPIRV/instrument_private_mem.ll | 14 ++++++++++++++ 1 file changed, 14 insertions(+) create mode 100644 llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll new file mode 100644 index 0000000000000..133a916026126 --- /dev/null +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll @@ -0,0 +1,14 @@ +; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -msan-spir-privates=1 -S | FileCheck %s +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" +target triple = "spir64-unknown-unknown" + +define spir_kernel void @MyKernel() sanitize_memory { +; CHECK-LABEL: @MyKernel +entry: + %array = alloca [4 x i32], align 4 + ; CHECK: %0 = ptrtoint ptr %array to i64 + ; CHECK: %1 = call i64 @__msan_get_shadow(i64 %0, i32 0) + ; CHECK: %2 = inttoptr i64 %1 to ptr addrspace(1) + ; CHECK: call void @llvm.memset.p1.i64(ptr addrspace(1) align 4 %2, i8 -1, i64 16, i1 false) + ret void +} From fa33675cd8d0f3fbb19265c18bd814fdc7afa627 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 5 Mar 2025 07:20:30 +0100 Subject: [PATCH 04/14] fix format --- .../source/loader/layers/sanitizer/msan/msan_interceptor.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp index 0a3c098b51341..a74f29683a8df 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp @@ -549,7 +549,8 @@ ur_result_t MsanInterceptor::prepareLaunch( "CleanShadow={}, " "Device={}, Debug={})", (void *)LaunchInfo.Data, (void *)LaunchInfo.Data->GlobalShadowOffset, - (void *)LaunchInfo.Data->LocalShadowOffset, (void *)LaunchInfo.Data->PrivateShadowOffset, + (void *)LaunchInfo.Data->LocalShadowOffset, + (void *)LaunchInfo.Data->PrivateShadowOffset, (void *)LaunchInfo.Data->CleanShadow, ToString(LaunchInfo.Data->DeviceTy), LaunchInfo.Data->Debug); From 2c6e46ccf112f775dcf0c72a333517a9f902b3e7 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 10 Mar 2025 04:47:57 +0100 Subject: [PATCH 05/14] wip --- clang/lib/Driver/SanitizerArgs.cpp | 3 ++ libdevice/sanitizer/asan_rtl.cpp | 2 + libdevice/sanitizer/msan_rtl.cpp | 66 ++++++++++++++++++++---------- 3 files changed, 50 insertions(+), 21 deletions(-) diff --git a/clang/lib/Driver/SanitizerArgs.cpp b/clang/lib/Driver/SanitizerArgs.cpp index 93d1781d134b7..f5a97ae4e93c0 100644 --- a/clang/lib/Driver/SanitizerArgs.cpp +++ b/clang/lib/Driver/SanitizerArgs.cpp @@ -1278,6 +1278,9 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args, CmdArgs.push_back("-mllvm"); CmdArgs.push_back("-msan-eager-checks=1"); + + CmdArgs.push_back("-mllvm"); + CmdArgs.push_back("-msan-poison-stack-with-call=1"); } else if (Sanitizers.has(SanitizerKind::Thread)) { CmdArgs.push_back("-fsanitize=thread"); // The tsan function entry/exit builtins are used to record stack diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index 6507aa3a93d8e..f6fc29a4f407c 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/libdevice/sanitizer/asan_rtl.cpp @@ -881,6 +881,8 @@ static __SYCL_CONSTANT__ const char __mem_set_shadow_private_end[] = static __SYCL_CONSTANT__ const char __mem_set_shadow_private[] = "[kernel] set_shadow_private(beg=%p, end=%p, val:%02X)\n"; +// We outline the function of setting shadow memory of private memory, because +// it may allocate failed on UR DEVICE_EXTERN_C_NOINLINE void __asan_set_shadow_private(uptr begin, uptr size, char val) { if (!__AsanLaunchInfo) diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp index 0691978815da0..7ad49c75036ac 100644 --- a/libdevice/sanitizer/msan_rtl.cpp +++ b/libdevice/sanitizer/msan_rtl.cpp @@ -9,6 +9,7 @@ #include "include/msan_rtl.hpp" #include "atomic.hpp" #include "device.h" +#include "include/sanitizer_defs.hpp" #include "msan/msan_libdevice.hpp" #include "spirv_vars.h" @@ -34,7 +35,8 @@ const __SYCL_CONSTANT__ char __msan_print_shadow[] = "[kernel] __msan_get_shadow(addr=%p, as=%d) = %p: %02X\n"; const __SYCL_CONSTANT__ char __msan_print_launchinfo[] = - "[kernel] !!! launchinfo %p (GlobalShadow=%p)\n"; + "[kernel] !!! launchinfo %p (GlobalShadow=%p, LocalShadow=%p, " + "PrivateShadow=%p)\n"; const __SYCL_CONSTANT__ char __msan_print_unsupport_device_type[] = "[kernel] Unsupport device type: %d\n"; @@ -65,7 +67,7 @@ const __SYCL_CONSTANT__ char __msan_print_private_shadow_out_of_bound[] = namespace { -inline size_t workgroup_linear_id() { +inline size_t WorkGroupLinearId() { return __spirv_BuiltInWorkgroupId.x * __spirv_BuiltInNumWorkgroups.y * __spirv_BuiltInNumWorkgroups.z + __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z + @@ -190,28 +192,27 @@ inline uptr __msan_get_shadow_pvc(uptr addr, uint32_t as) { return addr - (PVC_DEVICE_USM_END - shadow_end); } } else if (as == ADDRESS_SPACE_LOCAL) { - // The size of SLM is 128KB on PVC - constexpr unsigned SLM_SIZE = 128 * 1024; - const auto wg_lid = workgroup_linear_id(); const auto shadow_offset = GetMsanLaunchInfo->LocalShadowOffset; - if (shadow_offset != 0) { + // The size of SLM is 128KB on PVC + constexpr unsigned SLM_SIZE = 128 * 1024; + const auto wg_lid = WorkGroupLinearId(); return shadow_offset + (wg_lid * SLM_SIZE) + (addr & (SLM_SIZE - 1)); } } else if (as == ADDRESS_SPACE_PRIVATE) { - const auto wg_lid = workgroup_linear_id(); const auto shadow_offset = GetMsanLaunchInfo->PrivateShadowOffset; - if (shadow_offset != 0) { + const auto wg_lid = WorkGroupLinearId(); uptr shadow_ptr = shadow_offset + (wg_lid * MSAN_PRIVATE_SIZE) + (addr & (MSAN_PRIVATE_SIZE - 1)); - MSAN_DEBUG( - const auto shadow_offset_end = GetMsanLaunchInfo->PrivateShadowOffsetEnd; - if (shadow_ptr > shadow_offset_end) { - __spirv_ocl_printf(__msan_print_private_shadow_out_of_bound, addr, - shadow_ptr, wg_lid, (uptr)shadow_offset); - return 0; - }); + MSAN_DEBUG(const auto shadow_offset_end = + GetMsanLaunchInfo->PrivateShadowOffsetEnd; + if (shadow_ptr > shadow_offset_end) { + __spirv_ocl_printf(__msan_print_private_shadow_out_of_bound, + addr, shadow_ptr, wg_lid, + (uptr)shadow_offset); + return 0; + }); return shadow_ptr; } } @@ -260,7 +261,9 @@ DEVICE_EXTERN_C_NOINLINE uptr __msan_get_shadow(uptr addr, uint32_t as) { return shadow_ptr; MSAN_DEBUG(__spirv_ocl_printf(__msan_print_launchinfo, GetMsanLaunchInfo, - GetMsanLaunchInfo->GlobalShadowOffset)); + GetMsanLaunchInfo->GlobalShadowOffset, + GetMsanLaunchInfo->LocalShadowOffset, + GetMsanLaunchInfo->PrivateShadowOffset)); #if defined(__LIBDEVICE_PVC__) shadow_ptr = __msan_get_shadow_pvc(addr, as); @@ -372,7 +375,7 @@ MSAN_MEMCPY(3) MSAN_MEMCPY(4) /// -/// Initialize shdadow memory of local memory +/// Initialize shadow memory of local memory /// static __SYCL_CONSTANT__ const char __mem_set_shadow_local[] = @@ -411,16 +414,13 @@ DEVICE_EXTERN_C_NOINLINE void __msan_unpoison_shadow_static_local(uptr ptr, if (__spirv_LocalInvocationId_x() + __spirv_LocalInvocationId_y() + __spirv_LocalInvocationId_z() == 0) { - if (!GetMsanLaunchInfo) + if (!GetMsanLaunchInfo || GetMsanLaunchInfo->LocalShadowOffset == 0) return; MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, "__msan_unpoison_shadow_static_local")); auto shadow_address = __msan_get_shadow(ptr, ADDRESS_SPACE_LOCAL); - if (shadow_address == GetMsanLaunchInfo->CleanShadow) - return; - for (size_t i = 0; i < size; ++i) { ((__SYCL_GLOBAL__ u8 *)shadow_address)[i] = 0; } @@ -438,4 +438,28 @@ DEVICE_EXTERN_C_INLINE void __msan_barrier() { __spv::MemorySemanticsMask::WorkgroupMemory); } +static __SYCL_CONSTANT__ const char __msan_print_set_shadow_private[] = + "[kernel] set_shadow_private(beg=%p, end=%p, val:%02X)\n"; + +// We outline the function of setting shadow memory of private memory, because +// it may allocate failed on UR +DEVICE_EXTERN_C_NOINLINE void __msan_poison_stack(uptr ptr, uptr size) { + if (!GetMsanLaunchInfo || GetMsanLaunchInfo->PrivateShadowOffset == 0) + return; + + MSAN_DEBUG( + __spirv_ocl_printf(__msan_print_func_beg, "__msan_set_shadow_private")); + + auto shadow_address = __msan_get_shadow(ptr, ADDRESS_SPACE_PRIVATE); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_private, + (void *)shadow_address, + (void *)(shadow_address + size), 0xff)); + + for (size_t i = 0; i < size; i++) + ((__SYCL_GLOBAL__ u8 *)shadow_address)[i] = 0xff; + + MSAN_DEBUG( + __spirv_ocl_printf(__msan_print_func_beg, "__msan_set_shadow_end")); +} + #endif // __SPIR__ || __SPIRV__ From 1e376acdd4c8644c261ff3f346548d4d62b3531b Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 17 Mar 2025 11:04:35 +0100 Subject: [PATCH 06/14] wip --- libdevice/sanitizer/msan_rtl.cpp | 82 ++++-- .../Instrumentation/MemorySanitizer.cpp | 238 ++++++++++++++---- .../test-e2e/MemorySanitizer/check_buffer.cpp | 7 +- .../layers/sanitizer/msan/msan_libdevice.hpp | 1 + .../layers/sanitizer/msan/msan_report.cpp | 12 +- 5 files changed, 258 insertions(+), 82 deletions(-) diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp index 9cb9a30144f86..14afc32fc1bc8 100644 --- a/libdevice/sanitizer/msan_rtl.cpp +++ b/libdevice/sanitizer/msan_rtl.cpp @@ -32,7 +32,7 @@ constexpr uptr DG2_DEVICE_USM_BEGIN = 0xffff'8000'0000'0000ULL; constexpr uptr DG2_DEVICE_USM_END = 0xffff'ffff'ffff'ffffULL; const __SYCL_CONSTANT__ char __msan_print_shadow[] = - "[kernel] __msan_get_shadow(addr=%p, as=%d) = %p: %02X\n"; + "[kernel] __msan_get_shadow(addr=%p, as=%d) = %p: %02X <%s>\n"; const __SYCL_CONSTANT__ char __msan_print_launchinfo[] = "[kernel] !!! launchinfo %p (GlobalShadow=%p, LocalShadow=%p, " @@ -45,15 +45,16 @@ const __SYCL_CONSTANT__ char __msan_print_generic_to[] = "[kernel] %p(4) - %p(%d)\n"; const __SYCL_CONSTANT__ char __msan_print_func_beg[] = - "[kernel] ===== %s() begin\n"; + "[kernel] ===== BEGIN %s()\n"; const __SYCL_CONSTANT__ char __msan_print_func_end[] = - "[kernel] ===== %s() end\n"; + "[kernel] ===== END %s()\n"; const __SYCL_CONSTANT__ char __msan_print_private_shadow_out_of_bound[] = "[kernel] Private shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: " "%p)\n"; +const __SYCL_CONSTANT__ char __msan_print_unknown[] = "unknown"; } // namespace #if defined(__SPIR__) || defined(__SPIRV__) @@ -92,7 +93,8 @@ inline void ConvertGenericPointer(uptr &addr, uint32_t &as) { void __msan_internal_report_save(const uint32_t size, const char __SYCL_CONSTANT__ *file, const uint32_t line, - const char __SYCL_CONSTANT__ *func) { + const char __SYCL_CONSTANT__ *func, + const uptr origin) { const int Expected = MSAN_REPORT_NONE; int Desired = MSAN_REPORT_START; @@ -128,6 +130,7 @@ void __msan_internal_report_save(const uint32_t size, SanitizerReport.Func[MaxFuncIdx] = '\0'; SanitizerReport.AccessSize = size; + SanitizerReport.Origin = origin; SanitizerReport.Line = line; SanitizerReport.GID0 = __spirv_GlobalInvocationId_x(); SanitizerReport.GID1 = __spirv_GlobalInvocationId_y(); @@ -144,8 +147,8 @@ void __msan_internal_report_save(const uint32_t size, void __msan_report_error(const uint32_t size, const char __SYCL_CONSTANT__ *file, const uint32_t line, - const char __SYCL_CONSTANT__ *func) { - __msan_internal_report_save(size, file, line, func); + const char __SYCL_CONSTANT__ *func, uptr origin = 0) { + __msan_internal_report_save(size, file, line, func, origin); if (!GetMsanLaunchInfo->IsRecover) { __devicelib_exit(); @@ -224,12 +227,12 @@ inline uptr __msan_get_shadow_pvc(uptr addr, uint32_t as) { #define MSAN_MAYBE_WARNING(type, size) \ DEVICE_EXTERN_C_NOINLINE void __msan_maybe_warning_##size( \ - type s, u32 o, const char __SYCL_CONSTANT__ *file, uint32_t line, \ + type s, uptr o, const char __SYCL_CONSTANT__ *file, uint32_t line, \ const char __SYCL_CONSTANT__ *func) { \ if (!GetMsanLaunchInfo) \ return; \ if (UNLIKELY(s)) { \ - __msan_report_error(size, file, line, func); \ + __msan_report_error(size, file, line, func, o); \ } \ } @@ -247,13 +250,14 @@ __msan_warning(const char __SYCL_CONSTANT__ *file, uint32_t line, DEVICE_EXTERN_C_NOINLINE void __msan_warning_noreturn(const char __SYCL_CONSTANT__ *file, uint32_t line, const char __SYCL_CONSTANT__ *func) { - __msan_internal_report_save(1, file, line, func); + __msan_internal_report_save(1, file, line, func, 0); __devicelib_exit(); } // For mapping detail, ref to // "unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp" -DEVICE_EXTERN_C_NOINLINE uptr __msan_get_shadow(uptr addr, uint32_t as) { +DEVICE_EXTERN_C_NOINLINE uptr __msan_get_shadow( + uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *func = nullptr) { // Return clean shadow (0s) by default uptr shadow_ptr = GetMsanLaunchInfo->CleanShadow; @@ -283,12 +287,13 @@ DEVICE_EXTERN_C_NOINLINE uptr __msan_get_shadow(uptr addr, uint32_t as) { #endif MSAN_DEBUG(__spirv_ocl_printf(__msan_print_shadow, (void *)addr, as, - (void *)shadow_ptr, *(u8 *)shadow_ptr)); + (void *)shadow_ptr, *(u8 *)shadow_ptr, + func ? func : __msan_print_unknown)); return shadow_ptr; } -static __SYCL_CONSTANT__ const char __mem_memset[] = +static __SYCL_CONSTANT__ const char __msan_print_memset[] = "[kernel] memset(beg=%p, shadow_beg=%p, shadow_end=%p)\n"; #define MSAN_MEMSET(as) \ @@ -301,8 +306,8 @@ static __SYCL_CONSTANT__ const char __mem_memset[] = dest[i] = val; \ ((__SYCL_GLOBAL__ char *)shadow)[i] = 0; \ } \ - MSAN_DEBUG( \ - __spirv_ocl_printf(__mem_memset, dest, shadow, shadow + size - 1)); \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_memset, dest, shadow, \ + shadow + size - 1)); \ MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, "__msan_memset")); \ return dest; \ } @@ -312,11 +317,15 @@ MSAN_MEMSET(1) MSAN_MEMSET(3) MSAN_MEMSET(4) +static __SYCL_CONSTANT__ const char __msan_print_memmove[] = + "[kernel] memmove(dst=%p, src=%p, shadow_dst=%p, shadow_src=%p, size=%p)\n"; + #define MSAN_MEMMOVE_BASE(dst_as, src_as) \ DEVICE_EXTERN_C_NOINLINE __attribute__((address_space(dst_as))) void \ *__msan_memmove_p##dst_as##_p##src_as( \ __attribute__((address_space(dst_as))) char *dest, \ __attribute__((address_space(src_as))) char *src, size_t size) { \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, "__msan_memmove")); \ uptr dest_shadow = __msan_get_shadow((uptr)dest, dst_as); \ uptr src_shadow = __msan_get_shadow((uptr)src, src_as); \ if ((uptr)dest > (uptr)src) { \ @@ -332,6 +341,9 @@ MSAN_MEMSET(4) ((__SYCL_GLOBAL__ char *)src_shadow)[i]; \ } \ } \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_memmove, dest, src, \ + dest_shadow, src_shadow, size)); \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, "__msan_memmove")); \ return dest; \ } @@ -347,11 +359,15 @@ MSAN_MEMMOVE(1) MSAN_MEMMOVE(3) MSAN_MEMMOVE(4) +static __SYCL_CONSTANT__ const char __msan_print_memcpy[] = + "[kernel] memcpy(dst=%p, src=%p, shadow_dst=%p, shadow_src=%p, size=%p)\n"; + #define MSAN_MEMCPY_BASE(dst_as, src_as) \ DEVICE_EXTERN_C_NOINLINE __attribute__((address_space(dst_as))) void \ *__msan_memcpy_p##dst_as##_p##src_as( \ __attribute__((address_space(dst_as))) char *dest, \ __attribute__((address_space(src_as))) char *src, size_t size) { \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, "__msan_memcpy")); \ uptr dest_shadow = __msan_get_shadow((uptr)dest, dst_as); \ uptr src_shadow = __msan_get_shadow((uptr)src, src_as); \ for (size_t i = 0; i < size; i++) { \ @@ -359,6 +375,9 @@ MSAN_MEMMOVE(4) ((__SYCL_GLOBAL__ char *)dest_shadow)[i] = \ ((__SYCL_GLOBAL__ char *)src_shadow)[i]; \ } \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_memmove, dest, src, \ + dest_shadow, src_shadow, size)); \ + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, "__msan_memcpy")); \ return dest; \ } @@ -456,7 +475,8 @@ __msan_poison_shadow_dynamic_local(uptr ptr, uint32_t num_args) { if (!GetMsanLaunchInfo) return; - MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_dynamic_local_begin)); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, + "__msan_poison_shadow_dynamic_local")); if (num_args != GetMsanLaunchInfo->NumLocalArgs) { __spirv_ocl_printf(__msan_print_report_arg_count_incorrect, num_args, @@ -473,7 +493,8 @@ __msan_poison_shadow_dynamic_local(uptr ptr, uint32_t num_args) { __msan_poison_shadow_static_local(args[i], local_arg->Size); } - MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_dynamic_local_end)); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, + "__msan_poison_shadow_dynamic_local")); } static __SYCL_CONSTANT__ const char @@ -487,7 +508,8 @@ __msan_unpoison_shadow_dynamic_local(uptr ptr, uint32_t num_args) { if (!GetMsanLaunchInfo) return; - MSAN_DEBUG(__spirv_ocl_printf(__mem_unpoison_shadow_dynamic_local_begin)); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, + "__msan_unpoison_shadow_dynamic_local")); if (num_args != GetMsanLaunchInfo->NumLocalArgs) { __spirv_ocl_printf(__msan_print_report_arg_count_incorrect, num_args, @@ -504,7 +526,8 @@ __msan_unpoison_shadow_dynamic_local(uptr ptr, uint32_t num_args) { __msan_unpoison_shadow_static_local(args[i], local_arg->Size); } - MSAN_DEBUG(__spirv_ocl_printf(__mem_unpoison_shadow_dynamic_local_end)); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, + "__msan_unpoison_shadow_dynamic_local")); } static __SYCL_CONSTANT__ const char __msan_print_set_shadow_private[] = @@ -516,8 +539,7 @@ DEVICE_EXTERN_C_NOINLINE void __msan_poison_stack(uptr ptr, uptr size) { if (!GetMsanLaunchInfo || GetMsanLaunchInfo->PrivateShadowOffset == 0) return; - MSAN_DEBUG( - __spirv_ocl_printf(__msan_print_func_beg, "__msan_set_shadow_private")); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, "__msan_poison_stack")); auto shadow_address = __msan_get_shadow(ptr, ADDRESS_SPACE_PRIVATE); MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_private, @@ -527,8 +549,26 @@ DEVICE_EXTERN_C_NOINLINE void __msan_poison_stack(uptr ptr, uptr size) { for (size_t i = 0; i < size; i++) ((__SYCL_GLOBAL__ u8 *)shadow_address)[i] = 0xff; + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, "__msan_poison_stack")); +} + +DEVICE_EXTERN_C_NOINLINE void __msan_unpoison_stack(uptr ptr, uptr size) { + if (!GetMsanLaunchInfo || GetMsanLaunchInfo->PrivateShadowOffset == 0) + return; + + MSAN_DEBUG( + __spirv_ocl_printf(__msan_print_func_beg, "__msan_unpoison_stack")); + + auto shadow_address = __msan_get_shadow(ptr, ADDRESS_SPACE_PRIVATE); + MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_private, + (void *)shadow_address, + (void *)(shadow_address + size), 0x0)); + + for (size_t i = 0; i < size; i++) + ((__SYCL_GLOBAL__ u8 *)shadow_address)[i] = 0; + MSAN_DEBUG( - __spirv_ocl_printf(__msan_print_func_beg, "__msan_set_shadow_end")); + __spirv_ocl_printf(__msan_print_func_end, "__msan_unpoison_stack")); } #endif // __SPIR__ || __SPIRV__ diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 0aaf8af5d99be..143b13090de25 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -392,6 +392,10 @@ static cl::opt cl::desc("instrument private pointer"), cl::Hidden, cl::init(true)); +static cl::opt ClSpirOffloadDebug("msan-spir-debug", + cl::desc("enhance debug for spirv"), + cl::Hidden, cl::init(false)); + const char kMsanModuleCtorName[] = "msan.module_ctor"; const char kMsanInitName[] = "__msan_init"; @@ -781,7 +785,8 @@ class MemorySanitizerOnSpirv { } bool instrumentModule(); - void instrumentFunction(Function &F); + void beforeInstrumentFunction(Function &F, Instruction *FnPrologueEnd); + void afterInstrumentFunction(Function &F); Constant *getOrCreateGlobalString(StringRef Name, StringRef Value, unsigned AddressSpace); @@ -792,6 +797,7 @@ class MemorySanitizerOnSpirv { void instrumentStaticLocalMemory(); void instrumentDynamicLocalMemory(Function &F); void instrumentKernelsMetadata(); + void instrumentPrivateArguments(Function &F, Instruction *FnPrologueEnd); void initializeRetVecMap(Function *F); void initializeKernelCallerMap(Function *F); @@ -822,6 +828,7 @@ class MemorySanitizerOnSpirv { FunctionCallee MsanPoisonShadowDynamicLocalFunc; FunctionCallee MsanUnpoisonShadowDynamicLocalFunc; FunctionCallee MsanBarrierFunc; + FunctionCallee MsanUnpoisonStackFunc; }; } // end anonymous namespace @@ -904,6 +911,13 @@ void MemorySanitizerOnSpirv::initializeCallbacks() { // __msan_barrier() MsanBarrierFunc = M.getOrInsertFunction("__msan_barrier", IRB.getVoidTy()); + + // __msan_unpoison_stack( + // uptr ptr, + // size_t size + // ) + MsanUnpoisonStackFunc = M.getOrInsertFunction( + "__msan_unpoison_stack", IRB.getVoidTy(), IntptrTy, IntptrTy); } // Handle global variables: @@ -1076,6 +1090,46 @@ void MemorySanitizerOnSpirv::instrumentDynamicLocalMemory(Function &F) { InsertBarrier[&F] = true; } +void MemorySanitizerOnSpirv::instrumentPrivateArguments( + Function &F, Instruction *FnPrologueEnd) { + if (!ClSpirOffloadPrivates) + return; + + // We need to copy and replace all byval arguments to alloca in kernel because + // we need to make sure that all byval arguments have shadow memory. + // This change needs to be inserted after the prologue because instructions + // in prologue don't have shadow memory. + IRBuilder<> IRB(FnPrologueEnd->getNextNode()); + + for (auto &Arg : F.args()) { + PointerType *PtrTy = dyn_cast(Arg.getType()->getScalarType()); + if (PtrTy && PtrTy->getPointerAddressSpace() == kSpirOffloadPrivateAS && + Arg.hasByValAttr()) { + Type *Ty = Arg.getParamByValType(); + const Align Alignment = + DL.getValueOrABITypeAlignment(Arg.getParamAlign(), Ty); + + AllocaInst *AI = IRB.CreateAlloca( + Ty, kSpirOffloadPrivateAS, nullptr, + (Arg.hasName() ? Arg.getName() : "Arg" + Twine(Arg.getArgNo())) + + ".byval"); + AI->setAlignment(Alignment); + Arg.replaceAllUsesWith(AI); + + auto *UnpoisonStack = + IRB.CreateCall(MsanUnpoisonStackFunc, + {IRB.CreatePointerCast(AI, IntptrTy), + ConstantInt::get(IntptrTy, DL.getTypeAllocSize(Ty))}); + UnpoisonStack->setNoSanitizeMetadata(); + + uint64_t AllocSize = DL.getTypeAllocSize(Ty); + auto *Memcpy = + IRB.CreateMemCpy(AI, Alignment, &Arg, Alignment, AllocSize); + Memcpy->setNoSanitizeMetadata(); + } + } +} + // Instrument __MsanKernelMetadata, which records information of sanitized // kernel void MemorySanitizerOnSpirv::instrumentKernelsMetadata() { @@ -1162,7 +1216,16 @@ bool MemorySanitizerOnSpirv::instrumentModule() { return true; } -void MemorySanitizerOnSpirv::instrumentFunction(Function &F) { +void MemorySanitizerOnSpirv::beforeInstrumentFunction( + Function &F, Instruction *FnPrologueEnd) { + if (!IsSPIRV) + return; + + if (F.getCallingConv() == CallingConv::SPIR_KERNEL) + instrumentPrivateArguments(F, FnPrologueEnd); +} + +void MemorySanitizerOnSpirv::afterInstrumentFunction(Function &F) { if (!IsSPIRV) return; @@ -1193,7 +1256,6 @@ PreservedAnalyses MemorySanitizerPass::run(Module &M, MemorySanitizer Msan(*F.getParent(), MsanSpirv, Options); Modified |= Msan.sanitizeFunction(F, FAM.getResult(F)); - MsanSpirv.instrumentFunction(F); } if (!Modified) @@ -1352,31 +1414,35 @@ void MemorySanitizer::createUserspaceApi(Module &M, } } - // Create the global TLS variables. - RetvalTLS = - getOrInsertGlobal(M, "__msan_retval_tls", - ArrayType::get(IRB.getInt64Ty(), kRetvalTLSSize / 8)); + // SPIR-V doesn't support TLS variables + if (!TargetTriple.isSPIROrSPIRV()) { + // Create the global TLS variables. + RetvalTLS = + getOrInsertGlobal(M, "__msan_retval_tls", + ArrayType::get(IRB.getInt64Ty(), kRetvalTLSSize / 8)); - RetvalOriginTLS = getOrInsertGlobal(M, "__msan_retval_origin_tls", OriginTy); + RetvalOriginTLS = + getOrInsertGlobal(M, "__msan_retval_origin_tls", OriginTy); - ParamTLS = - getOrInsertGlobal(M, "__msan_param_tls", - ArrayType::get(IRB.getInt64Ty(), kParamTLSSize / 8)); + ParamTLS = + getOrInsertGlobal(M, "__msan_param_tls", + ArrayType::get(IRB.getInt64Ty(), kParamTLSSize / 8)); - ParamOriginTLS = - getOrInsertGlobal(M, "__msan_param_origin_tls", - ArrayType::get(OriginTy, kParamTLSSize / 4)); + ParamOriginTLS = + getOrInsertGlobal(M, "__msan_param_origin_tls", + ArrayType::get(OriginTy, kParamTLSSize / 4)); - VAArgTLS = - getOrInsertGlobal(M, "__msan_va_arg_tls", - ArrayType::get(IRB.getInt64Ty(), kParamTLSSize / 8)); + VAArgTLS = + getOrInsertGlobal(M, "__msan_va_arg_tls", + ArrayType::get(IRB.getInt64Ty(), kParamTLSSize / 8)); - VAArgOriginTLS = - getOrInsertGlobal(M, "__msan_va_arg_origin_tls", - ArrayType::get(OriginTy, kParamTLSSize / 4)); + VAArgOriginTLS = + getOrInsertGlobal(M, "__msan_va_arg_origin_tls", + ArrayType::get(OriginTy, kParamTLSSize / 4)); - VAArgOverflowSizeTLS = - getOrInsertGlobal(M, "__msan_va_arg_overflow_size_tls", IRB.getInt64Ty()); + VAArgOverflowSizeTLS = getOrInsertGlobal( + M, "__msan_va_arg_overflow_size_tls", IRB.getInt64Ty()); + } for (size_t AccessSizeIndex = 0; AccessSizeIndex < kNumberOfAccessSizes; AccessSizeIndex++) { @@ -1389,14 +1455,15 @@ void MemorySanitizer::createUserspaceApi(Module &M, IRB.getVoidTy(), IRB.getIntNTy(AccessSize * 8), IRB.getInt32Ty()); } else { // SPIR or SPIR-V // __msan_maybe_warning_N( - // ... + // intN_t status, + // uptr origin, // possible shadow address of status // char* file, // unsigned int line, // char* func // ) MaybeWarningFn[AccessSizeIndex] = M.getOrInsertFunction( FunctionName, TLI.getAttrList(C, {0, 1}, /*Signed=*/false), - IRB.getVoidTy(), IRB.getIntNTy(AccessSize * 8), IRB.getInt32Ty(), + IRB.getVoidTy(), IRB.getIntNTy(AccessSize * 8), IntptrTy, IRB.getInt8PtrTy(kSpirOffloadConstantAS), IRB.getInt32Ty(), IRB.getInt8PtrTy(kSpirOffloadConstantAS)); } @@ -1469,8 +1536,9 @@ void MemorySanitizer::initializeCallbacks(Module &M, MsanInstrumentAsmStoreFn = M.getOrInsertFunction( "__msan_instrument_asm_store", IRB.getVoidTy(), PtrTy, IntptrTy); - MsanGetShadowFn = M.getOrInsertFunction("__msan_get_shadow", IntptrTy, - IntptrTy, IRB.getInt32Ty()); + MsanGetShadowFn = M.getOrInsertFunction( + "__msan_get_shadow", IntptrTy, IntptrTy, IRB.getInt32Ty(), + IRB.getInt8PtrTy(kSpirOffloadConstantAS)); if (CompileKernel) { createKernelApi(M, TLI); @@ -2070,9 +2138,38 @@ struct MemorySanitizerVisitor : public InstVisitor { CB->addParamAttr(0, Attribute::ZExt); CB->addParamAttr(1, Attribute::ZExt); } else { // SPIR or SPIR-V - SmallVector Args = { - ConvertedShadow2, - MS.TrackOrigins && Origin ? Origin : (Value *)IRB.getInt32(0)}; + // Pass the pointer of shadow memory to the report function + SmallVector Args = {ConvertedShadow2}; + + if (ClSpirOffloadDebug) { + // Attempt to get the shadow memory + if (auto *LoadShadow = dyn_cast(ConvertedShadow)) { + Args.emplace_back(IRB.CreatePointerCast( + LoadShadow->getPointerOperand(), MS.IntptrTy)); + } else if (auto *BinaryOp = + dyn_cast(ConvertedShadow)) { + Value *LastOperand = nullptr; + do { + LastOperand = BinaryOp->getOperand(0); + // TODO: assert second operand is 0 + BinaryOp = dyn_cast(LastOperand); + } while (BinaryOp && BinaryOp->getOpcode() == Instruction::Or); + + if (auto *LoadShadow = dyn_cast(LastOperand)) { + Args.emplace_back(IRB.CreatePointerCast( + LoadShadow->getPointerOperand(), MS.IntptrTy)); + } + } else if (auto *Trunc = dyn_cast(ConvertedShadow)) { + if (auto *LoadShadow = dyn_cast(Trunc->getOperand(0))) { + Args.emplace_back(IRB.CreatePointerCast( + LoadShadow->getPointerOperand(), MS.IntptrTy)); + } + } + } + + if (Args.size() == 1) { + Args.emplace_back(ConstantInt::get(MS.IntptrTy, 0)); + } appendDebugInfoToArgs(IRB, Args); @@ -2097,7 +2194,7 @@ struct MemorySanitizerVisitor : public InstVisitor { const DataLayout &DL = F.getDataLayout(); // Disable combining in some cases. TrackOrigins checks each shadow to pick // correct origin. - bool Combine = !MS.TrackOrigins; + bool Combine = !(MS.TrackOrigins || ClSpirOffloadDebug); Instruction *Instruction = InstructionChecks.front().OrigIns; Value *Shadow = nullptr; for (const auto &ShadowData : InstructionChecks) { @@ -2398,10 +2495,17 @@ struct MemorySanitizerVisitor : public InstVisitor { OffsetLong = IRB.CreateXor(OffsetLong, constToIntPtr(IntptrTy, XorMask)); } else { // SPIR or SPIR-V + auto *ConstASPtrTy = PointerType::get(Type::getInt8Ty(Addr->getContext()), + kSpirOffloadConstantAS); + auto *FuncNameGV = MS.Spirv.getOrCreateGlobalString( + "__msan_func", F.getName(), kSpirOffloadConstantAS); + OffsetLong = IRB.CreateCall( MS.MsanGetShadowFn, - {OffsetLong, - IRB.getInt32(Addr->getType()->getPointerAddressSpace())}); + {OffsetLong, IRB.getInt32(Addr->getType()->getPointerAddressSpace()), + ClSpirOffloadDebug + ? ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy) + : ConstantPointerNull::get(ConstASPtrTy)}); } return OffsetLong; @@ -2659,10 +2763,7 @@ struct MemorySanitizerVisitor : public InstVisitor { unsigned ArgOffset = 0; const DataLayout &DL = F->getDataLayout(); for (auto &FArg : F->args()) { - // FIXME: Need to find a reasonable way to handle byval arguments for - // spirv target. - if (!FArg.getType()->isSized() || FArg.getType()->isScalableTy() || - (SpirOrSpirv && FArg.hasByValAttr())) { + if (!FArg.getType()->isSized() || FArg.getType()->isScalableTy()) { LLVM_DEBUG(dbgs() << (FArg.getType()->isScalableTy() ? "vscale not fully supported\n" : "Arg is not sized\n")); @@ -2679,6 +2780,24 @@ struct MemorySanitizerVisitor : public InstVisitor { : DL.getTypeAllocSize(FArg.getType()); if (A == &FArg) { + // SPIR-V doesn't propagate shadow for arguments. + if (SpirOrSpirv) { + if (FArg.hasByValAttr()) { + const Align ArgAlign = DL.getValueOrABITypeAlignment( + FArg.getParamAlign(), FArg.getParamByValType()); + Value *CpShadowPtr, *CpOriginPtr; + std::tie(CpShadowPtr, CpOriginPtr) = getShadowOriginPtr( + V, EntryIRB, EntryIRB.getInt8Ty(), ArgAlign, + /*isStore*/ true); + EntryIRB.CreateMemSet( + CpShadowPtr, Constant::getNullValue(EntryIRB.getInt8Ty()), + Size, ArgAlign); + } + ShadowPtr = getCleanShadow(V); + setOrigin(A, getCleanOrigin()); + break; + } + bool Overflow = ArgOffset + Size > kParamTLSSize; if (FArg.hasByValAttr()) { // ByVal pointer itself has clean shadow. We copy the actual @@ -2690,6 +2809,7 @@ struct MemorySanitizerVisitor : public InstVisitor { std::tie(CpShadowPtr, CpOriginPtr) = getShadowOriginPtr(V, EntryIRB, EntryIRB.getInt8Ty(), ArgAlign, /*isStore*/ true); + // Spirv needn't to consider overflow since it doesn't use TLS if (!PropagateShadow || Overflow) { // ParamTLS overflow. EntryIRB.CreateMemSet( @@ -2788,6 +2908,10 @@ struct MemorySanitizerVisitor : public InstVisitor { } #ifndef NDEBUG Type *ShadowTy = Shadow->getType(); + if (!(isa(ShadowTy) || isa(ShadowTy) || + isa(ShadowTy) || isa(ShadowTy))) { + ShadowTy->dump(); + } assert((isa(ShadowTy) || isa(ShadowTy) || isa(ShadowTy) || isa(ShadowTy)) && "Can only insert checks for integer, vector, and aggregate shadow " @@ -3593,22 +3717,24 @@ struct MemorySanitizerVisitor : public InstVisitor { /// __msan_memcpy(). Should this be wrong, such as when implementing memcpy() /// itself, instrumentation should be disabled with the no_sanitize attribute. void visitMemCpyInst(MemCpyInst &I) { - if (SpirOrSpirv && ((isa(I.getArgOperand(0)) && - cast(I.getArgOperand(0)) - ->getMetadata(LLVMContext::MD_nosanitize)) || - (isa(I.getArgOperand(1)) && - cast(I.getArgOperand(1)) - ->getMetadata(LLVMContext::MD_nosanitize)))) - return; + if (SpirOrSpirv) { + if ((isa(I.getArgOperand(0)) && + cast(I.getArgOperand(0)) + ->getMetadata(LLVMContext::MD_nosanitize)) || + (isa(I.getArgOperand(1)) && + cast(I.getArgOperand(1)) + ->getMetadata(LLVMContext::MD_nosanitize))) + return; - if (SpirOrSpirv && !ClSpirOffloadLocals && - (I.getSourceAddressSpace() == kSpirOffloadLocalAS || - I.getDestAddressSpace() == kSpirOffloadLocalAS)) - return; - if (SpirOrSpirv && !ClSpirOffloadPrivates && - (I.getSourceAddressSpace() == kSpirOffloadPrivateAS || - I.getDestAddressSpace() == kSpirOffloadPrivateAS)) - return; + // If we disable checking local/private memory, we needn't update its + // shadow memory, and treat its shadow value as zeros at rtl + if (!ClSpirOffloadLocals && + I.getDestAddressSpace() == kSpirOffloadLocalAS) + return; + if (!ClSpirOffloadPrivates && + I.getDestAddressSpace() == kSpirOffloadPrivateAS) + return; + } getShadow(I.getArgOperand(1)); // Ensure shadow initialized IRBuilder<> IRB(&I); @@ -5775,7 +5901,6 @@ struct MemorySanitizerVisitor : public InstVisitor { // Don't emit the epilogue for musttail call returns. if (isAMustTailRetVal(RetVal)) return; - Value *ShadowPtr = getShadowPtrForRetval(IRB); bool HasNoUndef = F.hasRetAttribute(Attribute::NoUndef); bool StoreShadow = !(MS.EagerChecks && HasNoUndef); // FIXME: Consider using SpecialCaseList to specify a list of functions that @@ -5793,6 +5918,7 @@ struct MemorySanitizerVisitor : public InstVisitor { // The caller may still expect information passed over TLS if we pass our // check if (StoreShadow) { + Value *ShadowPtr = getShadowPtrForRetval(IRB); IRB.CreateAlignedStore(Shadow, ShadowPtr, kShadowTLSAlignment); if (MS.TrackOrigins && StoreOrigin) IRB.CreateStore(getOrigin(RetVal), getOriginPtrForRetval()); @@ -5829,7 +5955,8 @@ struct MemorySanitizerVisitor : public InstVisitor { void poisonAllocaUserspace(AllocaInst &I, IRBuilder<> &IRB, Value *Len) { if (PoisonStack && ClPoisonStackWithCall) { - IRB.CreateCall(MS.MsanPoisonStackFn, {&I, Len}); + if (SpirOrSpirv && !I.getName().ends_with(".byval")) + IRB.CreateCall(MS.MsanPoisonStackFn, {&I, Len}); } else { Value *ShadowBase, *OriginBase; std::tie(ShadowBase, OriginBase) = getShadowOriginPtr( @@ -7396,5 +7523,8 @@ bool MemorySanitizer::sanitizeFunction(Function &F, TargetLibraryInfo &TLI) { B.addAttribute(Attribute::Memory).addAttribute(Attribute::Speculatable); F.removeFnAttrs(B); - return Visitor.runOnFunction(); + Spirv.beforeInstrumentFunction(F, Visitor.FnPrologueEnd); + bool Modified = Visitor.runOnFunction(); + Spirv.afterInstrumentFunction(F); + return Modified; } diff --git a/sycl/test-e2e/MemorySanitizer/check_buffer.cpp b/sycl/test-e2e/MemorySanitizer/check_buffer.cpp index df4c6ef7bfc5a..54117d002b9b8 100644 --- a/sycl/test-e2e/MemorySanitizer/check_buffer.cpp +++ b/sycl/test-e2e/MemorySanitizer/check_buffer.cpp @@ -8,7 +8,7 @@ #include -__attribute__((noinline)) long long foo(int data1, long long data2) { +__attribute__((noinline)) int foo(int data1, int data2) { return data1 + data2; } @@ -16,12 +16,9 @@ int main() { sycl::queue q; sycl::buffer buf1(sycl::range<1>(1)); - sycl::buffer buf2(sycl::range<1>(1)); q.submit([&](sycl::handler &h) { auto array1 = buf1.get_access(h); - auto array2 = buf2.get_access(h); - h.single_task( - [=]() { array1[0] = foo(array1[0], array2[0]); }); + h.single_task([=]() { foo(array1[0], array1[0]); }); }).wait(); // CHECK: use-of-uninitialized-value // CHECK: kernel <{{.*MyKernel}}> diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp index f9e601e6bc6d9..8d65f54daa242 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_libdevice.hpp @@ -37,6 +37,7 @@ struct MsanErrorReport { uint32_t AccessSize = 0; ErrorType ErrorTy = ErrorType::UNKNOWN; + uintptr_t Origin; }; struct MsanLocalArgsInfo { diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp index bdf2adf2b377e..7001e2324b791 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp @@ -30,8 +30,16 @@ void ReportUsesUninitializedValue(const MsanErrorReport &Report, // Try to demangle the kernel name KernelName = DemangleName(KernelName); - getContext()->logger.always( - "====WARNING: DeviceSanitizer: use-of-uninitialized-value"); + if (Report.Origin) { + getContext()->logger.always( + "====WARNING: DeviceSanitizer: use-of-uninitialized-value (shadow: {})", + (void *)Report.Origin); + } else { + getContext()->logger.always( + "====WARNING: DeviceSanitizer: use-of-uninitialized-value)", + (void *)Report.Origin); + } + getContext()->logger.always( "use of size {} at kernel <{}> LID({}, {}, {}) GID({}, " "{}, {})", From b5973518d28c26a31949d41d8067b4d41aa4d34d Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 18 Mar 2025 03:42:38 +0100 Subject: [PATCH 07/14] wip --- .../Instrumentation/MemorySanitizer.cpp | 71 +++++++++---------- 1 file changed, 35 insertions(+), 36 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 143b13090de25..eb838c248787f 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -2809,7 +2809,6 @@ struct MemorySanitizerVisitor : public InstVisitor { std::tie(CpShadowPtr, CpOriginPtr) = getShadowOriginPtr(V, EntryIRB, EntryIRB.getInt8Ty(), ArgAlign, /*isStore*/ true); - // Spirv needn't to consider overflow since it doesn't use TLS if (!PropagateShadow || Overflow) { // ParamTLS overflow. EntryIRB.CreateMemSet( @@ -2908,10 +2907,6 @@ struct MemorySanitizerVisitor : public InstVisitor { } #ifndef NDEBUG Type *ShadowTy = Shadow->getType(); - if (!(isa(ShadowTy) || isa(ShadowTy) || - isa(ShadowTy) || isa(ShadowTy))) { - ShadowTy->dump(); - } assert((isa(ShadowTy) || isa(ShadowTy) || isa(ShadowTy) || isa(ShadowTy)) && "Can only insert checks for integer, vector, and aggregate shadow " @@ -3675,22 +3670,23 @@ struct MemorySanitizerVisitor : public InstVisitor { /// /// Similar situation exists for memcpy and memset. void visitMemMoveInst(MemMoveInst &I) { - if (SpirOrSpirv && ((isa(I.getArgOperand(0)) && - cast(I.getArgOperand(0)) - ->getMetadata(LLVMContext::MD_nosanitize)) || - (isa(I.getArgOperand(1)) && - cast(I.getArgOperand(1)) - ->getMetadata(LLVMContext::MD_nosanitize)))) - return; + if (SpirOrSpirv) { + // If the destination is a nosanitize value, we don't need to update its + // shadow memory + if (isa(I.getArgOperand(0)) && + cast(I.getArgOperand(0)) + ->getMetadata(LLVMContext::MD_nosanitize)) + return; - if (SpirOrSpirv && !ClSpirOffloadLocals && - (I.getSourceAddressSpace() == kSpirOffloadLocalAS || - I.getDestAddressSpace() == kSpirOffloadLocalAS)) - return; - if (SpirOrSpirv && !ClSpirOffloadPrivates && - (I.getSourceAddressSpace() == kSpirOffloadPrivateAS || - I.getDestAddressSpace() == kSpirOffloadPrivateAS)) - return; + // If we disable checking local/private memory, we needn't update its + // shadow memory + if (!ClSpirOffloadLocals && + I.getDestAddressSpace() == kSpirOffloadLocalAS) + return; + if (!ClSpirOffloadPrivates && + I.getDestAddressSpace() == kSpirOffloadPrivateAS) + return; + } getShadow(I.getArgOperand(1)); // Ensure shadow initialized IRBuilder<> IRB(&I); @@ -3718,12 +3714,11 @@ struct MemorySanitizerVisitor : public InstVisitor { /// itself, instrumentation should be disabled with the no_sanitize attribute. void visitMemCpyInst(MemCpyInst &I) { if (SpirOrSpirv) { - if ((isa(I.getArgOperand(0)) && - cast(I.getArgOperand(0)) - ->getMetadata(LLVMContext::MD_nosanitize)) || - (isa(I.getArgOperand(1)) && - cast(I.getArgOperand(1)) - ->getMetadata(LLVMContext::MD_nosanitize))) + // If the destination is a nosanitize value, we don't need to update its + // shadow memory + if (isa(I.getArgOperand(0)) && + cast(I.getArgOperand(0)) + ->getMetadata(LLVMContext::MD_nosanitize)) return; // If we disable checking local/private memory, we needn't update its @@ -3748,17 +3743,21 @@ struct MemorySanitizerVisitor : public InstVisitor { // Same as memcpy. void visitMemSetInst(MemSetInst &I) { - if (SpirOrSpirv && isa(I.getArgOperand(0)) && - cast(I.getArgOperand(0)) - ->getMetadata(LLVMContext::MD_nosanitize)) - return; + if (SpirOrSpirv) { + if (isa(I.getArgOperand(0)) && + cast(I.getArgOperand(0)) + ->getMetadata(LLVMContext::MD_nosanitize)) + return; - if (SpirOrSpirv && !ClSpirOffloadLocals && - I.getDestAddressSpace() == kSpirOffloadLocalAS) - return; - if (SpirOrSpirv && !ClSpirOffloadPrivates && - I.getDestAddressSpace() == kSpirOffloadPrivateAS) - return; + // If we disable checking local/private memory, we needn't update its + // shadow memory, and treat its shadow value as zeros at rtl + if (!ClSpirOffloadLocals && + I.getDestAddressSpace() == kSpirOffloadLocalAS) + return; + if (!ClSpirOffloadPrivates && + I.getDestAddressSpace() == kSpirOffloadPrivateAS) + return; + } IRBuilder<> IRB(&I); IRB.CreateCall( From e13305cb9880c1d5465213a9bcf27ea769361dfd Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 18 Mar 2025 04:49:12 +0100 Subject: [PATCH 08/14] add test --- libdevice/sanitizer/msan_rtl.cpp | 10 ++++--- .../Instrumentation/MemorySanitizer.cpp | 10 +++---- .../SPIRV/check_large_access_size.ll | 19 ------------ .../SPIRV/instrument_private_mem.ll | 29 ++++++++++++++++--- 4 files changed, 36 insertions(+), 32 deletions(-) delete mode 100644 llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_large_access_size.ll diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp index 14afc32fc1bc8..90b309316e0fb 100644 --- a/libdevice/sanitizer/msan_rtl.cpp +++ b/libdevice/sanitizer/msan_rtl.cpp @@ -535,13 +535,14 @@ static __SYCL_CONSTANT__ const char __msan_print_set_shadow_private[] = // We outline the function of setting shadow memory of private memory, because // it may allocate failed on UR -DEVICE_EXTERN_C_NOINLINE void __msan_poison_stack(uptr ptr, uptr size) { +DEVICE_EXTERN_C_NOINLINE void __msan_poison_stack(__SYCL_PRIVATE__ void *ptr, + uptr size) { if (!GetMsanLaunchInfo || GetMsanLaunchInfo->PrivateShadowOffset == 0) return; MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_beg, "__msan_poison_stack")); - auto shadow_address = __msan_get_shadow(ptr, ADDRESS_SPACE_PRIVATE); + auto shadow_address = __msan_get_shadow((uptr)ptr, ADDRESS_SPACE_PRIVATE); MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_private, (void *)shadow_address, (void *)(shadow_address + size), 0xff)); @@ -552,14 +553,15 @@ DEVICE_EXTERN_C_NOINLINE void __msan_poison_stack(uptr ptr, uptr size) { MSAN_DEBUG(__spirv_ocl_printf(__msan_print_func_end, "__msan_poison_stack")); } -DEVICE_EXTERN_C_NOINLINE void __msan_unpoison_stack(uptr ptr, uptr size) { +DEVICE_EXTERN_C_NOINLINE void __msan_unpoison_stack(__SYCL_PRIVATE__ void *ptr, + uptr size) { if (!GetMsanLaunchInfo || GetMsanLaunchInfo->PrivateShadowOffset == 0) return; MSAN_DEBUG( __spirv_ocl_printf(__msan_print_func_beg, "__msan_unpoison_stack")); - auto shadow_address = __msan_get_shadow(ptr, ADDRESS_SPACE_PRIVATE); + auto shadow_address = __msan_get_shadow((uptr)ptr, ADDRESS_SPACE_PRIVATE); MSAN_DEBUG(__spirv_ocl_printf(__msan_print_set_shadow_private, (void *)shadow_address, (void *)(shadow_address + size), 0x0)); diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index eb838c248787f..874ce7312b36f 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -878,6 +878,7 @@ MemorySanitizerOnSpirv::getOrCreateGlobalString(StringRef Name, StringRef Value, // Initialize MSan runtime functions and globals void MemorySanitizerOnSpirv::initializeCallbacks() { IRBuilder<> IRB(C); + auto *PtrTy = IRB.getPtrTy(); // __msan_set_shadow_static_local( // uptr ptr, @@ -917,7 +918,7 @@ void MemorySanitizerOnSpirv::initializeCallbacks() { // size_t size // ) MsanUnpoisonStackFunc = M.getOrInsertFunction( - "__msan_unpoison_stack", IRB.getVoidTy(), IntptrTy, IntptrTy); + "__msan_unpoison_stack", IRB.getVoidTy(), PtrTy, IntptrTy); } // Handle global variables: @@ -1116,10 +1117,9 @@ void MemorySanitizerOnSpirv::instrumentPrivateArguments( AI->setAlignment(Alignment); Arg.replaceAllUsesWith(AI); - auto *UnpoisonStack = - IRB.CreateCall(MsanUnpoisonStackFunc, - {IRB.CreatePointerCast(AI, IntptrTy), - ConstantInt::get(IntptrTy, DL.getTypeAllocSize(Ty))}); + auto *UnpoisonStack = IRB.CreateCall( + MsanUnpoisonStackFunc, + {AI, ConstantInt::get(IntptrTy, DL.getTypeAllocSize(Ty))}); UnpoisonStack->setNoSanitizeMetadata(); uint64_t AllocSize = DL.getTypeAllocSize(Ty); diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_large_access_size.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_large_access_size.ll deleted file mode 100644 index 8a255418ec8bb..0000000000000 --- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/check_large_access_size.ll +++ /dev/null @@ -1,19 +0,0 @@ -; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -S | FileCheck %s -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" -target triple = "spir64-unknown-unknown" - -; Function Attrs: sanitize_memory -define spir_kernel void @MyKernel(<3 x i32> %extractVec.i8.i.i.i) #0 { -entry: - br label %for.body.i -for.body.i: ; preds = %for.body.i, %entry - %div.i.i.i.i.i.i = sdiv <3 x i32> zeroinitializer, %extractVec.i8.i.i.i - br label %for.body.i -} -; CHECK-LABEL: @MyKernel -; CHECK: %0 = load <3 x i32>, ptr inttoptr (i64 ptrtoint (ptr addrspace(1) @__msan_param_tls to i64) to ptr), align 8 -; CHECK: %1 = bitcast <3 x i32> %0 to i96 -; CHECK: %_mscmp = icmp ne i96 %1, 0 -; CHECK: call void @__msan_warning_noreturn(ptr addrspace(2) null, i32 0, ptr addrspace(2) @__msan_kernel) - -attributes #0 = { sanitize_memory } diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll index 133a916026126..7a68f58eadd9a 100644 --- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_private_mem.ll @@ -1,4 +1,4 @@ -; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -msan-spir-privates=1 -S | FileCheck %s +; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -msan-spir-privates=1 -msan-poison-stack-with-call=1 -S | FileCheck %s target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" target triple = "spir64-unknown-unknown" @@ -6,9 +6,30 @@ define spir_kernel void @MyKernel() sanitize_memory { ; CHECK-LABEL: @MyKernel entry: %array = alloca [4 x i32], align 4 - ; CHECK: %0 = ptrtoint ptr %array to i64 - ; CHECK: %1 = call i64 @__msan_get_shadow(i64 %0, i32 0) + ; CHECK: call void @__msan_poison_stack(ptr %array, i64 16) + ret void +} + +%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" } +%"class.sycl::_V1::detail::array" = type { [1 x i64] } + +define spir_func void @ByValFunc(ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_array12) sanitize_memory { +; CHECK-LABEL: @ByValFunc +entry: + ; CHECK: %0 = ptrtoint ptr %_arg_array12 to i64 + ; CHECK: %1 = call i64 @__msan_get_shadow(i64 %0, i32 0, ptr addrspace(2) null) ; CHECK: %2 = inttoptr i64 %1 to ptr addrspace(1) - ; CHECK: call void @llvm.memset.p1.i64(ptr addrspace(1) align 4 %2, i8 -1, i64 16, i1 false) + ; CHECK: call void @llvm.memset.p1.i64(ptr addrspace(1) align 8 %2, i8 0, i64 8, i1 false) + %_arg_array12.ascast = addrspacecast ptr %_arg_array12 to ptr addrspace(4) + ret void +} + +define spir_kernel void @ByValKernel(ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_array12) sanitize_memory { +; CHECK-LABEL: @ByValKernel +entry: + ; CHECK: %_arg_array12.byval = alloca %"class.sycl::_V1::range", align 8 + ; CHECK: call void @__msan_unpoison_stack(ptr %_arg_array12.byval, i64 8), !nosanitize + ; CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 %_arg_array12.byval, ptr align 8 %_arg_array12, i64 8, i1 false), !nosanitize + call void @ByValFunc(ptr %_arg_array12) ret void } From b7f27751e8b0b0d160c424d5e7f665087c00c91f Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 18 Mar 2025 06:38:11 +0100 Subject: [PATCH 09/14] add comments --- libdevice/sanitizer/msan_rtl.cpp | 1 - .../Instrumentation/MemorySanitizer.cpp | 17 ++++++++++------- .../layers/sanitizer/msan/msan_report.cpp | 3 +-- 3 files changed, 11 insertions(+), 10 deletions(-) diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp index 90b309316e0fb..9f367a846f969 100644 --- a/libdevice/sanitizer/msan_rtl.cpp +++ b/libdevice/sanitizer/msan_rtl.cpp @@ -9,7 +9,6 @@ #include "include/msan_rtl.hpp" #include "atomic.hpp" #include "device.h" -#include "include/sanitizer_defs.hpp" #include "msan/msan_libdevice.hpp" #include "spirv_vars.h" diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 874ce7312b36f..4324a33648e00 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -392,6 +392,12 @@ static cl::opt cl::desc("instrument private pointer"), cl::Hidden, cl::init(true)); +// This flag is used to enhance debug for spirv (internal use only) +// - Add function name (not demangled for easily debug) on __msan_get_shadow +// (but not work on GPU) +// - Diable combination of origin and shadow propagation +// - The "origin" parameter of "__msan_maybe_warning_N" is the shadow address +// of UUM static cl::opt ClSpirOffloadDebug("msan-spir-debug", cl::desc("enhance debug for spirv"), cl::Hidden, cl::init(false)); @@ -3714,15 +3720,12 @@ struct MemorySanitizerVisitor : public InstVisitor { /// itself, instrumentation should be disabled with the no_sanitize attribute. void visitMemCpyInst(MemCpyInst &I) { if (SpirOrSpirv) { - // If the destination is a nosanitize value, we don't need to update its - // shadow memory + // Same as memmove if (isa(I.getArgOperand(0)) && cast(I.getArgOperand(0)) ->getMetadata(LLVMContext::MD_nosanitize)) return; - // If we disable checking local/private memory, we needn't update its - // shadow memory, and treat its shadow value as zeros at rtl if (!ClSpirOffloadLocals && I.getDestAddressSpace() == kSpirOffloadLocalAS) return; @@ -3744,13 +3747,12 @@ struct MemorySanitizerVisitor : public InstVisitor { // Same as memcpy. void visitMemSetInst(MemSetInst &I) { if (SpirOrSpirv) { + // Same as memmove if (isa(I.getArgOperand(0)) && cast(I.getArgOperand(0)) ->getMetadata(LLVMContext::MD_nosanitize)) return; - // If we disable checking local/private memory, we needn't update its - // shadow memory, and treat its shadow value as zeros at rtl if (!ClSpirOffloadLocals && I.getDestAddressSpace() == kSpirOffloadLocalAS) return; @@ -5900,6 +5902,7 @@ struct MemorySanitizerVisitor : public InstVisitor { // Don't emit the epilogue for musttail call returns. if (isAMustTailRetVal(RetVal)) return; + Value *ShadowPtr = !SpirOrSpirv ? getShadowPtrForRetval(IRB) : nullptr; bool HasNoUndef = F.hasRetAttribute(Attribute::NoUndef); bool StoreShadow = !(MS.EagerChecks && HasNoUndef); // FIXME: Consider using SpecialCaseList to specify a list of functions that @@ -5917,7 +5920,6 @@ struct MemorySanitizerVisitor : public InstVisitor { // The caller may still expect information passed over TLS if we pass our // check if (StoreShadow) { - Value *ShadowPtr = getShadowPtrForRetval(IRB); IRB.CreateAlignedStore(Shadow, ShadowPtr, kShadowTLSAlignment); if (MS.TrackOrigins && StoreOrigin) IRB.CreateStore(getOrigin(RetVal), getOriginPtrForRetval()); @@ -5954,6 +5956,7 @@ struct MemorySanitizerVisitor : public InstVisitor { void poisonAllocaUserspace(AllocaInst &I, IRBuilder<> &IRB, Value *Len) { if (PoisonStack && ClPoisonStackWithCall) { + // ".byval" alloca is updated by MsanUnpoisonStackFn if (SpirOrSpirv && !I.getName().ends_with(".byval")) IRB.CreateCall(MS.MsanPoisonStackFn, {&I, Len}); } else { diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp index 7001e2324b791..c9c984fbc6058 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_report.cpp @@ -36,8 +36,7 @@ void ReportUsesUninitializedValue(const MsanErrorReport &Report, (void *)Report.Origin); } else { getContext()->logger.always( - "====WARNING: DeviceSanitizer: use-of-uninitialized-value)", - (void *)Report.Origin); + "====WARNING: DeviceSanitizer: use-of-uninitialized-value)"); } getContext()->logger.always( From be43892902361ac3db63e51d5fd91cad482289d3 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 18 Mar 2025 06:53:53 +0100 Subject: [PATCH 10/14] fix build --- llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp | 5 +++-- .../source/loader/layers/sanitizer/msan/msan_shadow.hpp | 2 ++ 2 files changed, 5 insertions(+), 2 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 4324a33648e00..26eccf175e8af 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -5956,8 +5956,9 @@ struct MemorySanitizerVisitor : public InstVisitor { void poisonAllocaUserspace(AllocaInst &I, IRBuilder<> &IRB, Value *Len) { if (PoisonStack && ClPoisonStackWithCall) { - // ".byval" alloca is updated by MsanUnpoisonStackFn - if (SpirOrSpirv && !I.getName().ends_with(".byval")) + // SPIR-V: ".byval" alloca is updated by "__msan_unpoison_stack", so + // skiped here + if (!SpirOrSpirv || !I.getName().ends_with(".byval")) IRB.CreateCall(MS.MsanPoisonStackFn, {&I, Len}); } else { Value *ShadowBase, *OriginBase; diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp index 4ec8ebfd38938..9a3ec3d4a4c04 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_shadow.hpp @@ -146,6 +146,8 @@ struct MsanShadowMemoryGPU : public MsanShadowMemory { ur_mutex VirtualMemMapsMutex; uptr LocalShadowOffset = 0; + + uptr PrivateShadowOffset = 0; }; // clang-format off From 5e265739ce274c6db9a52089692e7b6c4890af12 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 18 Mar 2025 06:57:53 +0100 Subject: [PATCH 11/14] fix merge --- .../loader/layers/sanitizer/msan/msan_interceptor.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp index 49b0c94a1e8ff..a83ec5aea6cee 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp @@ -556,12 +556,12 @@ ur_result_t MsanInterceptor::prepareLaunch( getContext()->logger.info( "LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateShadow={}, " - "CleanShadow={}, " - "Device={}, Debug={})", + "CleanShadow={}, LocalArgs={}, NumLocalArgs={}, Device={}, Debug={})", (void *)LaunchInfo.Data, (void *)LaunchInfo.Data->GlobalShadowOffset, (void *)LaunchInfo.Data->LocalShadowOffset, (void *)LaunchInfo.Data->PrivateShadowOffset, - (void *)LaunchInfo.Data->CleanShadow, ToString(LaunchInfo.Data->DeviceTy), + (void *)LaunchInfo.Data->CleanShadow, (void *)LaunchInfo.Data->LocalArgs, + LaunchInfo.Data->NumLocalArgs, ToString(LaunchInfo.Data->DeviceTy), LaunchInfo.Data->Debug); ur_result_t URes = From fd846cc9f129b925c8f354f615116bf770c855c0 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 19 Mar 2025 03:05:06 +0100 Subject: [PATCH 12/14] fix format --- .../source/loader/layers/validation/ur_valddi.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 169c2a437ab15..972912b4be75d 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -9437,12 +9437,12 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendNativeCommandExp( } if (getContext()->enableParameterValidation) { - if (NULL == hCommandBuffer) - return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - if (NULL == pfnNativeCommand) return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (NULL == hCommandBuffer) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; @@ -9682,11 +9682,11 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferGetNativeHandleExp( } if (getContext()->enableParameterValidation) { - if (NULL == hCommandBuffer) - return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - if (NULL == phNativeCommandBuffer) return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hCommandBuffer) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; } ur_result_t result = From 45f0111798ce54e733d65309811c970bbdc0a69a Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 19 Mar 2025 03:13:35 +0100 Subject: [PATCH 13/14] Revert "fix format" This reverts commit fd846cc9f129b925c8f354f615116bf770c855c0. --- .../source/loader/layers/validation/ur_valddi.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 972912b4be75d..169c2a437ab15 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -9437,12 +9437,12 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendNativeCommandExp( } if (getContext()->enableParameterValidation) { - if (NULL == pfnNativeCommand) - return UR_RESULT_ERROR_INVALID_NULL_POINTER; - if (NULL == hCommandBuffer) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (NULL == pfnNativeCommand) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; @@ -9682,11 +9682,11 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferGetNativeHandleExp( } if (getContext()->enableParameterValidation) { - if (NULL == phNativeCommandBuffer) - return UR_RESULT_ERROR_INVALID_NULL_POINTER; - if (NULL == hCommandBuffer) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == phNativeCommandBuffer) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; } ur_result_t result = From f667cc65a84b1fdb54464e1909308e6e066ddce6 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Wed, 19 Mar 2025 03:16:02 +0100 Subject: [PATCH 14/14] Revert "Revert "fix format"" This reverts commit 45f0111798ce54e733d65309811c970bbdc0a69a. --- .../source/loader/layers/validation/ur_valddi.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 169c2a437ab15..972912b4be75d 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -9437,12 +9437,12 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendNativeCommandExp( } if (getContext()->enableParameterValidation) { - if (NULL == hCommandBuffer) - return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - if (NULL == pfnNativeCommand) return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (NULL == hCommandBuffer) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; @@ -9682,11 +9682,11 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferGetNativeHandleExp( } if (getContext()->enableParameterValidation) { - if (NULL == hCommandBuffer) - return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - if (NULL == phNativeCommandBuffer) return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hCommandBuffer) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; } ur_result_t result =