Skip to content

Commit 01b465a

Browse files
authored
[DeviceSanitizer] Ignore target extension type (#15484)
We don't know exactly what size it is, so just ignore that type.
1 parent e8c804b commit 01b465a

File tree

2 files changed

+57
-4
lines changed

2 files changed

+57
-4
lines changed

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

Lines changed: 37 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1506,19 +1506,49 @@ static bool isUnsupportedAMDGPUAddrspace(Value *Addr) {
15061506
return false;
15071507
}
15081508

1509-
static bool isUnsupportedSPIRAccess(Value *Addr, Function *Func) {
1509+
static bool containsTargetExtType(const Type *Ty) {
1510+
if (isa<TargetExtType>(Ty))
1511+
return true;
1512+
1513+
if (Ty->isVectorTy())
1514+
return containsTargetExtType(Ty->getScalarType());
1515+
1516+
if (Ty->isArrayTy())
1517+
return containsTargetExtType(Ty->getArrayElementType());
1518+
1519+
if (auto *STy = dyn_cast<StructType>(Ty)) {
1520+
for (unsigned int i = 0; i < STy->getNumElements(); i++)
1521+
if (containsTargetExtType(STy->getElementType(i)))
1522+
return true;
1523+
return false;
1524+
}
1525+
1526+
return false;
1527+
}
1528+
1529+
static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) {
15101530
// Skip SPIR-V built-in varibles
15111531
auto *OrigValue = Addr->stripInBoundsOffsets();
15121532
if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
15131533
return true;
15141534

1535+
// Ignore load/store for target ext type since we can't know exactly what size
1536+
// it is.
1537+
if (isa<StoreInst>(Inst) &&
1538+
containsTargetExtType(
1539+
cast<StoreInst>(Inst)->getValueOperand()->getType()))
1540+
return true;
1541+
1542+
if (isa<LoadInst>(Inst) && containsTargetExtType(Inst->getType()))
1543+
return true;
1544+
15151545
Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
15161546
switch (PtrTy->getPointerAddressSpace()) {
15171547
case kSpirOffloadPrivateAS: {
15181548
if (!ClSpirOffloadPrivates)
15191549
return true;
15201550
// Skip kernel arguments
1521-
return Func->getCallingConv() == CallingConv::SPIR_KERNEL &&
1551+
return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL &&
15221552
isa<Argument>(Addr);
15231553
}
15241554
case kSpirOffloadGlobalAS: {
@@ -1756,7 +1786,10 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) {
17561786
// swifterror allocas are register promoted by ISel
17571787
!AI.isSwiftError() &&
17581788
// safe allocas are not interesting
1759-
!(SSGI && SSGI->isSafe(AI)));
1789+
!(SSGI && SSGI->isSafe(AI)) &&
1790+
// ignore alloc contains target ext type since we can't know exactly what
1791+
// size it is.
1792+
!containsTargetExtType(AI.getAllocatedType()));
17601793

17611794
ProcessedAllocas[&AI] = IsInteresting;
17621795
return IsInteresting;
@@ -1765,7 +1798,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) {
17651798
bool AddressSanitizer::ignoreAccess(Instruction *Inst, Value *Ptr) {
17661799
// SPIR has its own rules to filter the instrument accesses
17671800
if (TargetTriple.isSPIROrSPIRV()) {
1768-
if (isUnsupportedSPIRAccess(Ptr, Inst->getFunction()))
1801+
if (isUnsupportedSPIRAccess(Ptr, Inst))
17691802
return true;
17701803
} else {
17711804
// Instrument accesses from different address spaces only for AMDGPU.
Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
; RUN: opt < %s -passes=asan -asan-instrumentation-with-call-threshold=0 -asan-stack=0 -asan-globals=0 -asan-use-after-return=never -asan-stack-dynamic-alloca=0 -asan-mapping-scale=4 -S | FileCheck %s
2+
3+
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"
4+
target triple = "spir64-unknown-unknown"
5+
6+
%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1) }
7+
8+
define spir_kernel void @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE() {
9+
entry:
10+
; CHECK-NOT: MyAlloc
11+
%sub_a.i = alloca [2 x %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix"], i32 0, align 8
12+
br label %for.cond10.i
13+
14+
for.cond10.i: ; preds = %for.cond10.i, %entry
15+
%0 = load target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1), ptr null, align 8
16+
store target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) zeroinitializer, ptr null, align 8
17+
; CHECK-NOT: asan_load
18+
; CHECK-NOT: asan_store
19+
br label %for.cond10.i
20+
}

0 commit comments

Comments
 (0)