Skip to content

Commit 52baa2b

Browse files
AllanZyneAllanZyne
authored
[DeviceASAN] Enhance out-of-bound by checking shadow bounds on global memory (intel#20079)
Record the boundary of mapped shadow memory (VA), so that we can check if accessing unknown address address is an out-of-bound access. Usage: ```bash clang++ -fsycl -Xarch_device -fsanitize=address -Xarch_device -mllvm=-asan-spir-shadow-bounds=1 demo.cpp ``` Misc. - Runtime option `detect_locals` and `detect_privates` has been deprecated. Just using compiler flag to disable these checkings is enough
1 parent 2ea23b4 commit 52baa2b

File tree

14 files changed

+218
-66
lines changed

14 files changed

+218
-66
lines changed

‎libdevice/sanitizer/asan_rtl.cpp‎

Lines changed: 30 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,8 @@
99
#include "include/asan_rtl.hpp"
1010
#include "asan/asan_libdevice.hpp"
1111

12+
extern "C" __attribute__((weak)) const int __asan_check_shadow_bounds;
13+
1214
// Save the pointer to LaunchInfo
1315
__SYCL_GLOBAL__ uptr *__SYCL_LOCAL__ __AsanLaunchInfo;
1416

@@ -40,6 +42,9 @@ static const __SYCL_CONSTANT__ char __asan_print_shadow_value2[] =
4042
static __SYCL_CONSTANT__ const char __generic_to[] =
4143
"[kernel] %p(4) - %p(%d)\n";
4244

45+
static __SYCL_CONSTANT__ const char __asan_print_shadow_bound[] =
46+
"[kernel] addr: %p, shadow: %p, lower: %p, upper: %p\n";
47+
4348
#define ASAN_REPORT_NONE 0
4449
#define ASAN_REPORT_START 1
4550
#define ASAN_REPORT_FINISH 2
@@ -65,8 +70,12 @@ struct DebugInfo {
6570
uint32_t line;
6671
};
6772

73+
inline bool IsCheckShadowBounds() { return __asan_check_shadow_bounds; }
74+
6875
void ReportUnknownDevice(const DebugInfo *debug);
6976
void PrintShadowMemory(uptr addr, uptr shadow_address, uint32_t as);
77+
void SaveReport(ErrorType error_type, MemoryType memory_type, bool is_recover,
78+
const DebugInfo *debug);
7079

7180
__SYCL_GLOBAL__ void *ToGlobal(void *ptr) {
7281
return __spirv_GenericCastToPtrExplicit_ToGlobal(ptr, 5);
@@ -115,6 +124,16 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
115124
launch_info->GlobalShadowOffset + (addr >> ASAN_SHADOW_SCALE);
116125
}
117126

127+
if (IsCheckShadowBounds() &&
128+
(shadow_ptr < launch_info->GlobalShadowLowerBound ||
129+
shadow_ptr > launch_info->GlobalShadowUpperBound)) {
130+
ASAN_DEBUG(__spirv_ocl_printf(__asan_print_shadow_bound, addr, shadow_ptr,
131+
launch_info->GlobalShadowLowerBound,
132+
launch_info->GlobalShadowUpperBound));
133+
SaveReport(ErrorType::OUT_OF_BOUNDS, MemoryType::GLOBAL, false, debug);
134+
return 0;
135+
}
136+
118137
ASAN_DEBUG(
119138
const auto shadow_offset_end = launch_info->GlobalShadowOffsetEnd;
120139
if (shadow_ptr > shadow_offset_end) {
@@ -168,7 +187,7 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as,
168187
__spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr, sid,
169188
private_base);
170189
return 0;
171-
};
190+
}
172191

173192
return shadow_ptr;
174193
}
@@ -193,6 +212,16 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as,
193212
((addr & 0x7FFFFFFFFFFF) >> ASAN_SHADOW_SCALE);
194213
}
195214

215+
if (IsCheckShadowBounds() &&
216+
(shadow_ptr < launch_info->GlobalShadowLowerBound ||
217+
shadow_ptr > launch_info->GlobalShadowUpperBound)) {
218+
ASAN_DEBUG(__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) {

‎llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,7 @@ enum SanitizedKernelFlags : uint32_t {
5050
CHECK_PRIVATES = 1U << 3,
5151
CHECK_GENERICS = 1U << 4,
5252
MSAN_TRACK_ORIGINS = 1U << 5,
53+
ASAN_CHECK_SHADOW_BOUNDS = 1U << 6,
5354
};
5455

5556
} // namespace llvm

‎llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp‎

Lines changed: 58 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -441,27 +441,35 @@ static cl::opt<AsanDtorKind> ClOverrideDestructorKind(
441441
cl::init(AsanDtorKind::Invalid), cl::Hidden);
442442

443443
// SYCL flags
444+
static cl::opt<bool> ClSpirOffloadPrivates(
445+
"asan-spir-privates",
446+
cl::desc("Instrument private pointer on SPIR-V target"), cl::Hidden,
447+
cl::init(true));
448+
444449
static cl::opt<bool>
445-
ClSpirOffloadPrivates("asan-spir-privates",
446-
cl::desc("instrument private pointer"), cl::Hidden,
447-
cl::init(true));
450+
ClSpirOffloadGlobals("asan-spir-globals",
451+
cl::desc("Instrument global pointer on SPIR-V target"),
452+
cl::Hidden, cl::init(true));
448453

449-
static cl::opt<bool> ClSpirOffloadGlobals("asan-spir-globals",
450-
cl::desc("instrument global pointer"),
451-
cl::Hidden, cl::init(true));
454+
static cl::opt<bool>
455+
ClSpirOffloadLocals("asan-spir-locals",
456+
cl::desc("Instrument local pointer on SPIR-V target"),
457+
cl::Hidden, cl::init(true));
452458

453-
static cl::opt<bool> ClSpirOffloadLocals("asan-spir-locals",
454-
cl::desc("instrument local pointer"),
455-
cl::Hidden, cl::init(true));
459+
static cl::opt<bool> ClSpirOffloadGenerics(
460+
"asan-spir-generics",
461+
cl::desc("Instrument generic pointer on SPIR-V target"), cl::Hidden,
462+
cl::init(true));
456463

457464
static cl::opt<bool>
458-
ClSpirOffloadGenerics("asan-spir-generics",
459-
cl::desc("instrument generic pointer"), cl::Hidden,
460-
cl::init(true));
465+
ClDeviceGlobals("asan-device-globals",
466+
cl::desc("Instrument device globals on SPIR-V target"),
467+
cl::Hidden, cl::init(true));
461468

462-
static cl::opt<bool> ClDeviceGlobals("asan-device-globals",
463-
cl::desc("instrument device globals"),
464-
cl::Hidden, cl::init(true));
469+
static cl::opt<bool> ClSpirCheckShadowBounds(
470+
"asan-spir-shadow-bounds",
471+
cl::desc("Enable checking shadow bounds on SPIR-V target"), cl::Hidden,
472+
cl::init(false));
465473

466474
// Debug flags.
467475

@@ -1394,7 +1402,8 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM,
13941402
// following structure:
13951403
// uptr unmangled_kernel_name
13961404
// uptr unmangled_kernel_name_size
1397-
StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
1405+
// uptr sanitized_flags
1406+
StructType *StructTy = StructType::get(IntptrTy, IntptrTy, IntptrTy);
13981407

13991408
if (!HasESIMD)
14001409
for (Function &F : M) {
@@ -1425,9 +1434,21 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM,
14251434
KernelNamesBytes.append(KernelName.begin(), KernelName.end());
14261435
auto *KernelNameGV = GetOrCreateGlobalString(
14271436
M, "__asan_kernel", KernelName, kSpirOffloadConstantAS);
1437+
1438+
uintptr_t SanitizerFlags = 0;
1439+
SanitizerFlags |= ClSpirOffloadLocals ? SanitizedKernelFlags::CHECK_LOCALS
1440+
: SanitizedKernelFlags::NO_CHECK;
1441+
SanitizerFlags |= ClSpirOffloadPrivates
1442+
? SanitizedKernelFlags::CHECK_PRIVATES
1443+
: SanitizedKernelFlags::NO_CHECK;
1444+
SanitizerFlags |= ClSpirCheckShadowBounds != 0
1445+
? SanitizedKernelFlags::ASAN_CHECK_SHADOW_BOUNDS
1446+
: SanitizedKernelFlags::NO_CHECK;
1447+
14281448
SpirKernelsMetadata.emplace_back(ConstantStruct::get(
14291449
StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy),
1430-
ConstantInt::get(IntptrTy, KernelName.size())));
1450+
ConstantInt::get(IntptrTy, KernelName.size()),
1451+
ConstantInt::get(IntptrTy, SanitizerFlags)));
14311452
}
14321453

14331454
// Create global variable to record spirv kernels' information
@@ -1615,6 +1636,17 @@ PreservedAnalyses AddressSanitizerPass::run(Module &M,
16151636
ExtendSpirKernelArgs(M, FAM, HasESIMD);
16161637
Modified = true;
16171638

1639+
{
1640+
IRBuilder<> IRB(M.getContext());
1641+
M.getOrInsertGlobal("__asan_check_shadow_bounds", IRB.getInt32Ty(), [&] {
1642+
return new GlobalVariable(
1643+
M, IRB.getInt32Ty(), true, GlobalValue::WeakODRLinkage,
1644+
ConstantInt::get(IRB.getInt32Ty(), ClSpirCheckShadowBounds),
1645+
"__asan_check_shadow_bounds", nullptr,
1646+
llvm::GlobalValue::NotThreadLocal, kSpirOffloadGlobalAS);
1647+
});
1648+
}
1649+
16181650
if (HasESIMD) {
16191651
GlobalStringMap.clear();
16201652
return PreservedAnalyses::none();
@@ -1693,19 +1725,21 @@ static bool isUnsupportedAMDGPUAddrspace(Value *Addr) {
16931725
}
16941726

16951727
static bool isUnsupportedDeviceGlobal(GlobalVariable *G) {
1696-
// Non image scope device globals are implemented by device USM, and the
1697-
// out-of-bounds check for them will be done by sanitizer USM part. So we
1698-
// exclude them here.
1699-
if (!G->hasAttribute("sycl-device-image-scope"))
1700-
return true;
1701-
17021728
// Skip instrumenting on "__AsanKernelMetadata" etc.
1703-
if (G->getName().starts_with("__Asan"))
1729+
if (G->getName().starts_with("__Asan") || G->getName().starts_with("__asan"))
17041730
return true;
17051731

17061732
if (G->getAddressSpace() == kSpirOffloadLocalAS)
17071733
return !ClSpirOffloadLocals;
17081734

1735+
// When shadow bounds check is enabled, we need to instrument all global
1736+
// variables that user code can access
1737+
if (ClSpirCheckShadowBounds)
1738+
return false;
1739+
1740+
// Non image scope device globals are implemented by device USM, and the
1741+
// out-of-bounds check for them will be done by sanitizer USM part. So we
1742+
// exclude them here.
17091743
Attribute Attr = G->getAttribute("sycl-device-image-scope");
17101744
return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false");
17111745
}

‎llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,4 +26,4 @@ entry:
2626
attributes #0 = { sanitize_address }
2727
;; sycl-device-global-size = 16 * 2
2828
;; sycl-host-access = 0 read-only
29-
; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" }
29+
; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="48" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" }

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

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

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

Lines changed: 59 additions & 22 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;
@@ -453,7 +465,7 @@ ur_result_t AsanInterceptor::unregisterProgram(ur_program_handle_t Program) {
453465
}
454466
ProgramInfo->AllocInfoForGlobals.clear();
455467

456-
ProgramInfo->InstrumentedKernels.clear();
468+
ProgramInfo->KernelMetadataMap.clear();
457469

458470
return UR_RESULT_SUCCESS;
459471
}
@@ -508,14 +520,18 @@ ur_result_t AsanInterceptor::registerSpirKernels(ur_program_handle_t Program) {
508520

509521
std::string KernelName =
510522
std::string(KernelNameV.begin(), KernelNameV.end());
523+
bool CheckShadowBounds =
524+
SKI.Flags & SanitizedKernelFlags::ASAN_CHECK_SHADOW_BOUNDS;
511525

512526
UR_LOG_L(getContext()->logger, INFO,
513-
"SpirKernel(name='{}', isInstrumented={})", KernelName, true);
527+
"SpirKernel(name='{}', isInstrumented={}, checkShadowBounds={})",
528+
KernelName, true, CheckShadowBounds);
514529

515-
PI->InstrumentedKernels.insert(std::move(KernelName));
530+
PI->KernelMetadataMap[KernelName] =
531+
ProgramInfo::KernelMetadata{CheckShadowBounds};
516532
}
517533
UR_LOG_L(getContext()->logger, INFO, "Number of sanitized kernel: {}",
518-
PI->InstrumentedKernels.size());
534+
PI->KernelMetadataMap.size());
519535
}
520536

521537
return UR_RESULT_SUCCESS;
@@ -666,11 +682,16 @@ KernelInfo &AsanInterceptor::getOrCreateKernelInfo(ur_kernel_handle_t Kernel) {
666682
auto Program = GetProgram(Kernel);
667683
auto PI = getProgramInfo(Program);
668684
assert(PI != nullptr && "unregistered program!");
669-
bool IsInstrumented = PI->isKernelInstrumented(Kernel);
685+
686+
auto KI = std::make_unique<KernelInfo>(Kernel);
687+
KI->IsInstrumented = PI->isKernelInstrumented(Kernel);
688+
if (KI->IsInstrumented) {
689+
auto &KM = PI->getKernelMetadata(Kernel);
690+
KI->IsCheckShadowBounds = KM.CheckShadowBounds;
691+
}
670692

671693
std::scoped_lock<ur_shared_mutex> Guard(m_KernelMapMutex);
672-
m_KernelMap.emplace(Kernel,
673-
std::make_unique<KernelInfo>(Kernel, IsInstrumented));
694+
m_KernelMap.emplace(Kernel, std::move(KI));
674695
return *m_KernelMap[Kernel].get();
675696
}
676697

@@ -815,6 +836,13 @@ ur_result_t AsanInterceptor::prepareLaunch(
815836
LaunchInfo.Data.Host.DeviceTy = DeviceInfo->Type;
816837
LaunchInfo.Data.Host.Debug = getContext()->Options.Debug ? 1 : 0;
817838

839+
if (KernelInfo.IsCheckShadowBounds) {
840+
LaunchInfo.Data.Host.GlobalShadowLowerBound =
841+
DeviceInfo->Shadow->ShadowLowerBound;
842+
LaunchInfo.Data.Host.GlobalShadowUpperBound =
843+
DeviceInfo->Shadow->ShadowUpperBound;
844+
}
845+
818846
// Write shadow memory offset for local memory
819847
if (getContext()->Options.DetectLocals) {
820848
if (DeviceInfo->Shadow->AllocLocalShadow(
@@ -872,18 +900,20 @@ ur_result_t AsanInterceptor::prepareLaunch(
872900
// sync asan runtime data to device side
873901
UR_CALL(LaunchInfo.Data.syncToDevice(Queue));
874902

875-
UR_LOG_L(getContext()->logger, INFO,
876-
"LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateBase={}, "
877-
"PrivateShadow={}, LocalArgs={}, NumLocalArgs={}, "
878-
"Device={}, Debug={})",
879-
(void *)LaunchInfo.Data.getDevicePtr(),
880-
(void *)LaunchInfo.Data.Host.GlobalShadowOffset,
881-
(void *)LaunchInfo.Data.Host.LocalShadowOffset,
882-
(void *)LaunchInfo.Data.Host.PrivateBase,
883-
(void *)LaunchInfo.Data.Host.PrivateShadowOffset,
884-
(void *)LaunchInfo.Data.Host.LocalArgs,
885-
LaunchInfo.Data.Host.NumLocalArgs,
886-
ToString(LaunchInfo.Data.Host.DeviceTy), LaunchInfo.Data.Host.Debug);
903+
UR_LOG_L(
904+
getContext()->logger, INFO,
905+
"LaunchInfo {} (GlobalShadow={}, LocalShadow={}, PrivateBase={}, "
906+
"PrivateShadow={}, GlobalShadowLowerBound={}, GlobalShadowUpperBound={}, "
907+
"LocalArgs={}, NumLocalArgs={}, Device={}, 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+
ToString(LaunchInfo.Data.Host.DeviceTy), LaunchInfo.Data.Host.Debug);
887917

888918
return UR_RESULT_SUCCESS;
889919
}
@@ -920,7 +950,14 @@ AsanInterceptor::findAllocInfoByContext(ur_context_handle_t Context) {
920950

921951
bool ProgramInfo::isKernelInstrumented(ur_kernel_handle_t Kernel) const {
922952
const auto Name = GetKernelName(Kernel);
923-
return InstrumentedKernels.find(Name) != InstrumentedKernels.end();
953+
return KernelMetadataMap.find(Name) != KernelMetadataMap.end();
954+
}
955+
956+
const ProgramInfo::KernelMetadata &
957+
ProgramInfo::getKernelMetadata(ur_kernel_handle_t Kernel) const {
958+
const auto Name = GetKernelName(Kernel);
959+
assert(KernelMetadataMap.find(Name) != KernelMetadataMap.end());
960+
return KernelMetadataMap.at(Name);
924961
}
925962

926963
ContextInfo::~ContextInfo() {

0 commit comments

Comments
 (0)