Skip to content

Commit 62acbce

Browse files
haonanya1haonanya1
committed
Update lit
1 parent fc2f34e commit 62acbce

File tree

5 files changed

+102
-82
lines changed

5 files changed

+102
-82
lines changed

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -29,10 +29,6 @@ constexpr unsigned kSpirOffloadGenericAS = 4;
2929

3030
TargetExtType *getTargetExtType(Type *Ty);
3131
bool isJointMatrixAccess(Value *V);
32-
bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I,
33-
bool InstrumentLocalPtr,
34-
bool InstrumentPrivatePtr);
35-
3632
} // namespace SanitizerCommonUtils
3733
} // namespace llvm
3834

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

Lines changed: 48 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1668,6 +1668,53 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) {
16681668
return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false");
16691669
}
16701670

1671+
static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) {
1672+
// Skip SPIR-V built-in varibles
1673+
auto *OrigValue = Addr->stripInBoundsOffsets();
1674+
if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
1675+
return true;
1676+
1677+
GlobalVariable *GV = dyn_cast<GlobalVariable>(OrigValue);
1678+
if (GV && isUnsupportedDeviceGlobal(GV))
1679+
return true;
1680+
1681+
// Ignore load/store for target ext type since we can't know exactly what size
1682+
// it is.
1683+
if (auto *SI = dyn_cast<StoreInst>(Inst))
1684+
if (SanitizerCommonUtils::getTargetExtType(
1685+
SI->getValueOperand()->getType()) ||
1686+
SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand()))
1687+
return true;
1688+
1689+
if (auto *LI = dyn_cast<LoadInst>(Inst))
1690+
if (SanitizerCommonUtils::getTargetExtType(Inst->getType()) ||
1691+
SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand()))
1692+
return true;
1693+
1694+
Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
1695+
switch (PtrTy->getPointerAddressSpace()) {
1696+
case SanitizerCommonUtils::kSpirOffloadPrivateAS: {
1697+
if (!ClSpirOffloadPrivates)
1698+
return true;
1699+
// Skip kernel arguments
1700+
return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL &&
1701+
isa<Argument>(Addr);
1702+
}
1703+
case SanitizerCommonUtils::kSpirOffloadGlobalAS: {
1704+
return !ClSpirOffloadGlobals;
1705+
}
1706+
case SanitizerCommonUtils::kSpirOffloadLocalAS: {
1707+
if (!ClSpirOffloadLocals)
1708+
return true;
1709+
return Addr->getName().starts_with("__Asan");
1710+
}
1711+
case SanitizerCommonUtils::kSpirOffloadGenericAS: {
1712+
return !ClSpirOffloadGenerics;
1713+
}
1714+
}
1715+
return true;
1716+
}
1717+
16711718
void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore,
16721719
Value *Addr,
16731720
SmallVectorImpl<Value *> &Args) {
@@ -1836,8 +1883,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) {
18361883
bool AddressSanitizer::ignoreAccess(Instruction *Inst, Value *Ptr) {
18371884
// SPIR has its own rules to filter the instrument accesses
18381885
if (TargetTriple.isSPIROrSPIRV()) {
1839-
if (SanitizerCommonUtils::isUnsupportedSPIRAccess(
1840-
Ptr, Inst, ClSpirOffloadLocals, ClSpirOffloadPrivates))
1886+
if (isUnsupportedSPIRAccess(Ptr, Inst))
18411887
return true;
18421888
} else {
18431889
// Instrument accesses from different address spaces only for AMDGPU.

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

Lines changed: 39 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1735,6 +1735,44 @@ static unsigned TypeSizeToSizeIndex(TypeSize TS) {
17351735
return Log2_32_Ceil((TypeSizeFixed + 7) / 8);
17361736
}
17371737

1738+
static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) {
1739+
if (isa<Instruction>(Addr) &&
1740+
cast<Instruction>(Addr)->getMetadata(LLVMContext::MD_nosanitize)) {
1741+
return true;
1742+
}
1743+
1744+
// Skip SPIR-V built-in varibles
1745+
auto *OrigValue = Addr->stripInBoundsOffsets();
1746+
assert(OrigValue != nullptr);
1747+
if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
1748+
return true;
1749+
1750+
// Ignore load/store for target ext type since we can't know exactly what size
1751+
// it is.
1752+
if (auto *SI = dyn_cast<StoreInst>(I))
1753+
if (SanitizerCommonUtils::getTargetExtType(
1754+
SI->getValueOperand()->getType()) ||
1755+
SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand()))
1756+
return true;
1757+
1758+
if (auto *LI = dyn_cast<LoadInst>(I))
1759+
if (SanitizerCommonUtils::getTargetExtType(I->getType()) ||
1760+
SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand()))
1761+
return true;
1762+
1763+
Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
1764+
switch (PtrTy->getPointerAddressSpace()) {
1765+
case kSpirOffloadPrivateAS:
1766+
return !ClSpirOffloadPrivates;
1767+
case kSpirOffloadLocalAS:
1768+
return !ClSpirOffloadLocals;
1769+
case kSpirOffloadGenericAS:
1770+
return false;
1771+
}
1772+
1773+
return false;
1774+
}
1775+
17381776
static void setNoSanitizedMetadataSPIR(Instruction &I) {
17391777
const Value *Addr = nullptr;
17401778
if (const auto *LI = dyn_cast<LoadInst>(&I))
@@ -1786,8 +1824,7 @@ static void setNoSanitizedMetadataSPIR(Instruction &I) {
17861824
}
17871825
}
17881826

1789-
if (Addr && SanitizerCommonUtils::isUnsupportedSPIRAccess(
1790-
Addr, &I, ClSpirOffloadLocals, ClSpirOffloadPrivates))
1827+
if (Addr && isUnsupportedSPIRAccess(Addr, &I))
17911828
I.setNoSanitizeMetadata();
17921829
}
17931830

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

Lines changed: 0 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -59,45 +59,5 @@ bool isJointMatrixAccess(Value *V) {
5959
}
6060
return false;
6161
}
62-
63-
bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I,
64-
bool InstrumentLocalPtr,
65-
bool InstrumentPrivatePtr) {
66-
if (isa<Instruction>(Addr) &&
67-
cast<Instruction>(Addr)->getMetadata(LLVMContext::MD_nosanitize)) {
68-
return true;
69-
}
70-
71-
// Skip SPIR-V built-in varibles
72-
auto *OrigValue = Addr->stripInBoundsOffsets();
73-
assert(OrigValue != nullptr);
74-
if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
75-
return true;
76-
77-
// Ignore load/store for target ext type since we can't know exactly what size
78-
// it is.
79-
if (auto *SI = dyn_cast<StoreInst>(I))
80-
if (getTargetExtType(SI->getValueOperand()->getType()) ||
81-
isJointMatrixAccess(SI->getPointerOperand()))
82-
return true;
83-
84-
if (auto *LI = dyn_cast<LoadInst>(I))
85-
if (getTargetExtType(I->getType()) ||
86-
isJointMatrixAccess(LI->getPointerOperand()))
87-
return true;
88-
89-
Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
90-
switch (PtrTy->getPointerAddressSpace()) {
91-
case kSpirOffloadPrivateAS:
92-
return !InstrumentPrivatePtr;
93-
case kSpirOffloadLocalAS:
94-
return !InstrumentLocalPtr;
95-
case kSpirOffloadGenericAS:
96-
return false;
97-
}
98-
99-
return false;
100-
}
101-
10262
} // namespace SanitizerCommonUtils
10363
} // namespace llvm

‎llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll‎

Lines changed: 15 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -3,42 +3,23 @@
33
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"
44
target triple = "spir64-unknown-unknown"
55

6-
%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1) }
6+
%"class.sycl::_V1::ext::oneapi::bfloat16" = type { i16 }
7+
%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.CooperativeMatrixKHR", i16, 3, 16, 32, 0) }
78

8-
; Function Attrs: sanitize_address
9-
define spir_kernel void @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE() #0 {
10-
entry:
11-
; CHECK-LABEL-DAG: @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE
12-
; CHECK-NOT: MyAlloc
13-
%a = alloca [2 x %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix"], i32 0, align 8
14-
br label %for.cond10.i
15-
16-
for.cond10.i: ; preds = %for.cond10.i, %entry
17-
%0 = load target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1), ptr null, align 8
18-
store target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) zeroinitializer, ptr null, align 8
19-
; CHECK-NOT: call void @asan_load
20-
; CHECK-NOT: call void @asan_store
21-
br label %for.cond10.i
22-
}
9+
; CHECK-LABEL: @test
10+
; CHECK-NOT: call i64 @__msan_get_shadow
11+
declare dso_local spir_func noundef ptr addrspace(4) @_Z19__spirv_AccessChainIN4sycl3_V13ext6oneapi8bfloat16ES4_Lm16ELm32ELN5__spv9MatrixUseE0ELNS5_5Scope4FlagE3EEPT_PPNS5_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr addrspace(4) noundef, i64 noundef)
2312

24-
; Function Attrs: sanitize_address
25-
define spir_kernel void @AccessChain() #0 {
13+
define weak_odr dso_local spir_kernel void @test() {
2614
entry:
27-
; CHECK-LABEL-DAG: @AccessChain
28-
%a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8
29-
%0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0
30-
%call.i35 = call spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0)
31-
%1 = getelementptr inbounds { i16 }, ptr %call.i35, i64 0, i32 0
32-
; CHECK-NOT: call void @__asan_load
33-
; CHECK-NOT: call void @__asan_store
34-
%2 = load i16, ptr %1, align 4
35-
%call.i42 = call spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0)
36-
%3 = getelementptr inbounds { i16 }, ptr %call.i42, i64 0, i32 0
37-
store i16 %2, ptr %3, align 4
15+
%sub_a.i = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8
16+
%element.i = alloca %"class.sycl::_V1::ext::oneapi::bfloat16", align 2
17+
%0 = getelementptr inbounds { i16 }, ptr %element.i, i64 0, i32 0
18+
%spvm.i = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %sub_a.i, i64 0, i32 0
19+
%addrcast = addrspacecast ptr %spvm.i to ptr addrspace(4)
20+
%call.i67 = call spir_func noundef ptr addrspace(4) @_Z19__spirv_AccessChainIN4sycl3_V13ext6oneapi8bfloat16ES4_Lm16ELm32ELN5__spv9MatrixUseE0ELNS5_5Scope4FlagE3EEPT_PPNS5_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr addrspace(4) noundef %addrcast, i64 1)
21+
%gep = getelementptr inbounds nuw { i16 }, ptr addrspace(4) %call.i67, i64 0, i32 0
22+
%val = load i16, ptr %0, align 2
23+
store i16 %val, ptr addrspace(4) %gep, align 2
3824
ret void
3925
}
40-
41-
declare spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr, i64)
42-
43-
attributes #0 = { sanitize_address }
44-

0 commit comments

Comments
 (0)