Skip to content

Commit

Permalink
Merge remote-tracking branch 'my_remote/sycl' into sycl
Browse files Browse the repository at this point in the history
  • Loading branch information
@smanna12
smanna12 committed Jun 28, 2024
2 parents ceb0eab + c9842c1 commit 4d9dfa1
Show file tree
Hide file tree
Showing 16 changed files with 311 additions and 467 deletions.
6 changes: 3 additions & 3 deletions libdevice/include/asan_libdevice.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -66,8 +66,8 @@ struct LocalArgsInfo {
constexpr std::size_t ASAN_MAX_NUM_REPORTS = 10;

struct LaunchInfo {
// Don't move this field, we use it in AddressSanitizerPass
uintptr_t PrivateShadowOffset = 0;
uintptr_t PrivateShadowOffsetEnd = 0;

uintptr_t LocalShadowOffset = 0;
uintptr_t LocalShadowOffsetEnd = 0;
Expand All @@ -82,8 +82,8 @@ constexpr unsigned ASAN_SHADOW_SCALE = 4;
constexpr unsigned ASAN_SHADOW_GRANULARITY = 1ULL << ASAN_SHADOW_SCALE;

// Based on the observation, only the last 24 bits of the address of the private
// variable have changed, we use 31 bits(2G) to be safe.
constexpr std::size_t ASAN_PRIVATE_SIZE = 0x7fffffffULL + 1;
// variable have changed
constexpr std::size_t ASAN_PRIVATE_SIZE = 0xffffffULL + 1;

// These magic values are written to shadow for better error
// reporting.
Expand Down
53 changes: 50 additions & 3 deletions libdevice/sanitizer_utils.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -65,6 +65,9 @@ static const __SYCL_CONSTANT__ char __global_shadow_out_of_bound[] =
static const __SYCL_CONSTANT__ char __local_shadow_out_of_bound[] =
"[kernel] Local shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: "
"%p)\n";
static const __SYCL_CONSTANT__ char __private_shadow_out_of_bound[] =
"[kernel] Private shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: "
"%p)\n";

static const __SYCL_CONSTANT__ char __asan_print_unsupport_device_type[] =
"[kernel] Unsupport device type: %d\n";
Expand Down Expand Up @@ -123,7 +126,7 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
}

if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
if (__asan_report_out_of_shadow_bounds() && __AsanDebug) {
if (__asan_report_out_of_shadow_bounds()) {
__spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr);
}
}
Expand Down Expand Up @@ -171,7 +174,7 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
}

if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
if (__asan_report_out_of_shadow_bounds() && __AsanDebug) {
if (__asan_report_out_of_shadow_bounds()) {
__spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
(uptr)__AsanShadowMemoryGlobalStart);
}
Expand Down Expand Up @@ -207,13 +210,46 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
((addr & (SLM_SIZE - 1)) >> ASAN_SHADOW_SCALE);

if (shadow_ptr > shadow_offset_end) {
if (__asan_report_out_of_shadow_bounds() && __AsanDebug) {
if (__asan_report_out_of_shadow_bounds()) {
__spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
wg_lid, (uptr)shadow_offset);
}
return 0;
}
return shadow_ptr;
} else if (as == ADDRESS_SPACE_PRIVATE) { // private
// 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;

auto launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
const auto shadow_offset = launch_info->PrivateShadowOffset;
const auto shadow_offset_end = launch_info->PrivateShadowOffsetEnd;

if (shadow_offset == 0) {
return 0;
}

if (__AsanDebug)
__spirv_ocl_printf(__mem_launch_info, launch_info,
launch_info->PrivateShadowOffset, 0,
launch_info->NumLocalArgs, launch_info->LocalArgs);

uptr shadow_ptr = shadow_offset +
((WG_LID * ASAN_PRIVATE_SIZE) >> ASAN_SHADOW_SCALE) +
((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE);

if (shadow_ptr > shadow_offset_end) {
if (__asan_report_out_of_shadow_bounds()) {
__spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr,
WG_LID, (uptr)shadow_offset);
}
return 0;
}
return shadow_ptr;
}

return 0;
Expand All @@ -233,6 +269,8 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
return shadow_ptr;
}

// FIXME: OCL "O2" optimizer doesn't work well with following code
#if 0
if (__AsanDebug) {
if (shadow_ptr) {
if (as == ADDRESS_SPACE_PRIVATE)
Expand All @@ -244,6 +282,7 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
__spirv_ocl_printf(__asan_print_shadow_value2, addr, as, shadow_ptr);
}
}
#endif

return shadow_ptr;
}
Expand Down Expand Up @@ -606,6 +645,14 @@ ASAN_REPORT_ERROR(store, true, 16)
ASAN_REPORT_ERROR_N(load, false)
ASAN_REPORT_ERROR_N(store, true)

///
/// ASAN convert memory address to shadow memory address
///

DEVICE_EXTERN_C_NOINLINE uptr __asan_mem_to_shadow(uptr ptr, uint32_t as) {
return MemToShadow(ptr, as);
}

///
/// ASAN initialize shdadow memory of local memory
///
Expand Down
102 changes: 55 additions & 47 deletions llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -179,6 +179,8 @@ const char kAMDGPUAddressPrivateName[] = "llvm.amdgcn.is.private";
const char kAMDGPUBallotName[] = "llvm.amdgcn.ballot.i64";
const char kAMDGPUUnreachableName[] = "llvm.amdgcn.unreachable";

const char kAsanMemToShadow[] = "__asan_mem_to_shadow";

// Accesses sizes are powers of two: 1, 2, 4, 8, 16.
static const size_t kNumberOfAccessSizes = 5;

Expand Down Expand Up @@ -447,7 +449,7 @@ static cl::opt<AsanDtorKind> ClOverrideDestructorKind(
static cl::opt<bool>
ClSpirOffloadPrivates("asan-spir-privates",
cl::desc("instrument private pointer"), cl::Hidden,
cl::init(false));
cl::init(true));

static cl::opt<bool> ClSpirOffloadGlobals("asan-spir-globals",
cl::desc("instrument global pointer"),
Expand Down Expand Up @@ -820,14 +822,15 @@ struct AddressSanitizer {
Value *SizeArgument, uint32_t Exp,
RuntimeCallInserter &RTCI);
void instrumentMemIntrinsic(MemIntrinsic *MI, RuntimeCallInserter &RTCI);
Value *memToShadow(Value *Shadow, IRBuilder<> &IRB);
Value *memToShadow(Value *Shadow, IRBuilder<> &IRB,
uint32_t AddressSpace = kSpirOffloadPrivateAS);
bool suppressInstrumentationSiteForDebug(int &Instrumented);
bool instrumentFunction(Function &F, const TargetLibraryInfo *TLI);
bool maybeInsertAsanInitAtFunctionEntry(Function &F);
bool maybeInsertDynamicShadowAtFunctionEntry(Function &F);
void markEscapedLocalAllocas(Function &F);
void instrumentSyclStaticLocalMemory(CallInst *CI);
void instrumentSyclDynamicLocalMemory(Function &F);
bool instrumentSyclDynamicLocalMemory(Function &F);

GlobalVariable *GetOrCreateGlobalString(Module &M, StringRef Name,
StringRef Value,
Expand Down Expand Up @@ -899,6 +902,8 @@ struct AddressSanitizer {
FunctionCallee AMDGPUAddressPrivate;
int InstrumentationWithCallsThreshold;
uint32_t MaxInlinePoisoningSize;

FunctionCallee AsanMemToShadow;
};

class ModuleAddressSanitizer {
Expand Down Expand Up @@ -1067,7 +1072,7 @@ struct FunctionStackPoisoner : public InstVisitor<FunctionStackPoisoner> {
DIB(*F.getParent(), /*AllowUnresolved*/ false), C(ASan.C),
IntptrTy(ASan.IntptrTy), IntptrPtrTy(PointerType::get(IntptrTy, 0)),
Mapping(ASan.Mapping),
PoisonStack(ClStack &&
PoisonStack((ClStack || ClSpirOffloadPrivates) &&
!Triple(F.getParent()->getTargetTriple()).isAMDGPU()) {}

bool runOnFunction() {
Expand Down Expand Up @@ -1350,7 +1355,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM) {
}

// Fixup all users
for (auto [F, NewF] : SpirFuncs) {
for (auto &[F, NewF] : SpirFuncs) {
SmallVector<User *, 16> Users(F->users());
for (User *U : Users) {
if (auto *CI = dyn_cast<CallInst>(U)) {
Expand Down Expand Up @@ -1544,13 +1549,13 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore,
Args.push_back(ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy));
}

Value *AddressSanitizer::memToShadow(Value *Shadow, IRBuilder<> &IRB) {
Value *AddressSanitizer::memToShadow(Value *Shadow, IRBuilder<> &IRB,
uint32_t AddressSpace) {
if (TargetTriple.isSPIR()) {
// ((Shadow & 0xffffffff) >> 3) + __AsanShadowMemoryPrivateStart;
Shadow = IRB.CreateAnd(Shadow, ConstantInt::get(IntptrTy, 0xffffffff));
Shadow = IRB.CreateLShr(Shadow, Mapping.Scale);
Value *ShadowBase = IRB.CreateLoad(IntptrTy, AsanShadowDevicePrivate);
return IRB.CreateAdd(Shadow, ShadowBase);
return IRB.CreateCall(
AsanMemToShadow,
{Shadow, ConstantInt::get(IRB.getInt32Ty(), AddressSpace)},
"shadow_ptr");
}
// Shadow >> scale
Shadow = IRB.CreateLShr(Shadow, Mapping.Scale);
Expand Down Expand Up @@ -1619,7 +1624,7 @@ void AddressSanitizer::instrumentSyclStaticLocalMemory(CallInst *CI) {
}

// Instument dynamic local memory
void AddressSanitizer::instrumentSyclDynamicLocalMemory(Function &F) {
bool AddressSanitizer::instrumentSyclDynamicLocalMemory(Function &F) {
InstrumentationIRBuilder IRB(F.getEntryBlock().getFirstNonPHI());

// Save "__asan_launch" into local memory "__AsanLaunchInfo"
Expand All @@ -1631,13 +1636,12 @@ void AddressSanitizer::instrumentSyclDynamicLocalMemory(Function &F) {
SmallVector<Argument *> LocalArgs;
for (auto &Arg : F.args()) {
Type *PtrTy = dyn_cast<PointerType>(Arg.getType()->getScalarType());
// Local address space
if (PtrTy && PtrTy->getPointerAddressSpace() == 3)
if (PtrTy && PtrTy->getPointerAddressSpace() == kSpirOffloadLocalAS)
LocalArgs.push_back(&Arg);
}

if (LocalArgs.empty())
return;
return false;

AllocaInst *ArgsArray = IRB.CreateAlloca(
IntptrTy, ConstantInt::get(Int32Ty, LocalArgs.size()), "local_args");
Expand All @@ -1649,6 +1653,7 @@ void AddressSanitizer::instrumentSyclDynamicLocalMemory(Function &F) {
IRB.CreateCall(AsanSetShadowDynamicLocalFunc,
{IRB.CreatePointerCast(ArgsArray, IntptrTy),
ConstantInt::get(Int32Ty, LocalArgs.size())});
return true;
}

// Instrument memset/memmove/memcpy
Expand Down Expand Up @@ -3232,14 +3237,6 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T
ArrayType::get(IRB.getInt8Ty(), 0));

if (TargetTriple.isSPIR()) {
AsanShadowDevicePrivate =
M.getOrInsertGlobal("__AsanShadowMemoryPrivateStart", IntptrTy, [&] {
return new GlobalVariable(M, IntptrTy, true,
GlobalVariable::ExternalLinkage, nullptr,
"__AsanShadowMemoryPrivateStart", nullptr,
GlobalVariable::NotThreadLocal, 1);
});

// __asan_set_shadow_static_local(
// uptr ptr,
// size_t size,
Expand All @@ -3263,6 +3260,9 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T
GlobalVariable::ExternalLinkage, nullptr, "__AsanLaunchInfo",
nullptr, GlobalVariable::NotThreadLocal, kSpirOffloadLocalAS);
});

AsanMemToShadow = M.getOrInsertFunction(kAsanMemToShadow, IntptrTy,
IntptrTy, Type::getInt32Ty(*C));
}

AMDGPUAddressShared =
Expand Down Expand Up @@ -3391,10 +3391,6 @@ bool AddressSanitizer::instrumentFunction(Function &F,
// can be passed to that intrinsic.
markEscapedLocalAllocas(F);

if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
instrumentSyclDynamicLocalMemory(F);
}

// We want to instrument every address only once per basic block (unless there
// are calls between uses).
SmallPtrSet<Value *, 16> TempsToInstrument;
Expand Down Expand Up @@ -3514,6 +3510,11 @@ bool AddressSanitizer::instrumentFunction(Function &F,
if (ChangedStack || !NoReturnCalls.empty())
FunctionModified = true;

// We need to instrument dynamic local arguments after stack poisoner
if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
FunctionModified |= instrumentSyclDynamicLocalMemory(F);
}

LLVM_DEBUG(dbgs() << "ASAN done instrumenting: " << FunctionModified << " "
<< F << "\n");

Expand Down Expand Up @@ -3999,32 +4000,39 @@ void FunctionStackPoisoner::processStaticAllocas() {
AI->replaceAllUsesWith(NewAllocaPtr);
}

auto TargetTriple = Triple(F.getParent()->getTargetTriple());

// The left-most redzone has enough space for at least 4 pointers.
// Write the Magic value to redzone[0].
Value *BasePlus0 = IRB.CreateIntToPtr(LocalStackBase, IntptrPtrTy);
IRB.CreateStore(ConstantInt::get(IntptrTy, kCurrentStackFrameMagic),
BasePlus0);
// Write the frame description constant to redzone[1].
Value *BasePlus1 = IRB.CreateIntToPtr(
IRB.CreateAdd(LocalStackBase,
ConstantInt::get(IntptrTy, ASan.LongSize / 8)),
IntptrPtrTy);
GlobalVariable *StackDescriptionGlobal =
createPrivateGlobalForString(*F.getParent(), DescriptionString,
/*AllowMerging*/ true, kAsanGenPrefix);
Value *Description = IRB.CreatePointerCast(StackDescriptionGlobal, IntptrTy);
IRB.CreateStore(Description, BasePlus1);
// Write the PC to redzone[2].
Value *BasePlus2 = IRB.CreateIntToPtr(
IRB.CreateAdd(LocalStackBase,
ConstantInt::get(IntptrTy, 2 * ASan.LongSize / 8)),
IntptrPtrTy);
IRB.CreateStore(IRB.CreatePointerCast(&F, IntptrTy), BasePlus2);
// SPIRV doesn't use the following metadata
if (!TargetTriple.isSPIR()) {
// Write the Magic value to redzone[0].
IRB.CreateStore(ConstantInt::get(IntptrTy, kCurrentStackFrameMagic),
BasePlus0);
// Write the frame description constant to redzone[1].
Value *BasePlus1 = IRB.CreateIntToPtr(
IRB.CreateAdd(LocalStackBase,
ConstantInt::get(IntptrTy, ASan.LongSize / 8)),
IntptrPtrTy);
GlobalVariable *StackDescriptionGlobal =
createPrivateGlobalForString(*F.getParent(), DescriptionString,
/*AllowMerging*/ true, kAsanGenPrefix);
Value *Description =
IRB.CreatePointerCast(StackDescriptionGlobal, IntptrTy);
IRB.CreateStore(Description, BasePlus1);
// Write the PC to redzone[2].
Value *BasePlus2 = IRB.CreateIntToPtr(
IRB.CreateAdd(LocalStackBase,
ConstantInt::get(IntptrTy, 2 * ASan.LongSize / 8)),
IntptrPtrTy);
IRB.CreateStore(IRB.CreatePointerCast(&F, IntptrTy), BasePlus2);
}

const auto &ShadowAfterScope = GetShadowBytesAfterScope(SVD, L);

// Poison the stack red zones at the entry.
Value *ShadowBase = ASan.memToShadow(LocalStackBase, IRB);
Value *ShadowBase =
ASan.memToShadow(LocalStackBase, IRB, kSpirOffloadPrivateAS);
// As mask we must use most poisoned case: red zones and after scope.
// As bytes we can use either the same or just red zones only.
copyToShadow(ShadowAfterScope, ShadowAfterScope, IRB, ShadowBase);
Expand Down
32 changes: 32 additions & 0 deletions llvm/test/Instrumentation/AddressSanitizer/SPIRV/instrument_private_address_space.ll
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,32 @@
; RUN: opt < %s -passes=asan -asan-instrumentation-with-call-threshold=0 -asan-stack=0 -asan-globals=0 -asan-constructor-kind=none -asan-spir-privates=1 -asan-use-after-return=never -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"
target triple = "spir64-unknown-unknown"

%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
%"class.sycl::_V1::detail::array" = type { [1 x i64] }
%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }

@__const._ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv.p = private unnamed_addr addrspace(1) constant [4 x i32] [i32 1, i32 2, i32 3, i32 4], align 4

define spir_func i32 @_Z3fooPii(ptr addrspace(4) %p) {
entry:
%arrayidx = getelementptr inbounds i32, ptr addrspace(4) %p, i64 0
%0 = load i32, ptr addrspace(4) %arrayidx, align 4
ret i32 %0
}

define spir_kernel void @kernel() #0 {
; CHECK-LABEL: define spir_kernel void @kernel
entry:
%p.i = alloca [4 x i32], align 4
; CHECK: %shadow_ptr = call i64 @__asan_mem_to_shadow(i64 %0, i32 0)
call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %p.i)
call void @llvm.memcpy.p0.p1.i64(ptr align 4 %p.i, ptr addrspace(1) align 4 @__const._ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv.p, i64 16, i1 false)
%arraydecay.i = getelementptr inbounds [4 x i32], ptr %p.i, i64 0, i64 0
%0 = addrspacecast ptr %arraydecay.i to ptr addrspace(4)
%call.i = call spir_func i32 @_Z3fooPii(ptr addrspace(4) %0)
ret void
}

attributes #0 = { mustprogress norecurse nounwind sanitize_address uwtable }
Loading

0 comments on commit 4d9dfa1

Please sign in to comment.