From ba592fac0201d1656222991c5d92728d0144743a Mon Sep 17 00:00:00 2001 From: haonanya1 Date: Sun, 23 Mar 2025 22:20:06 -0700 Subject: [PATCH 1/6] [DeviceSanitizer] Move isUnsupportedSPIRAccess to SanitizerCommonUtils namespace. Both AddressSanitizer and MemorySanitizer ignore target extension type. --- .../Instrumentation/SanitizerCommonUtils.h | 39 +++++++ .../Instrumentation/AddressSanitizer.cpp | 95 +--------------- .../Transforms/Instrumentation/CMakeLists.txt | 1 + .../Instrumentation/MemorySanitizer.cpp | 29 +---- .../Instrumentation/SanitizerCommonUtils.cpp | 103 ++++++++++++++++++ .../SPIRV/ignore_target_ext_type.ll | 44 ++++++++ 6 files changed, 194 insertions(+), 117 deletions(-) create mode 100644 llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h create mode 100644 llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp create mode 100644 llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll diff --git a/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h b/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h new file mode 100644 index 0000000000000..011bd449b0629 --- /dev/null +++ b/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h @@ -0,0 +1,39 @@ +//===- SanitizerCommonUtils.h - Sanitizer commnon utils ---------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file declares common infrastructure for Sanitizer. +// +//===----------------------------------------------------------------------===// +#ifndef LLVM_TRANSFORMS_INSTRUMENTATION_SANITIZERCOMMONUTILS_H +#define LLVM_TRANSFORMS_INSTRUMENTATION_SANITIZERCOMMONUTILS_H + +#include "llvm/IR/DerivedTypes.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/Type.h" +#include "llvm/IR/Value.h" + +namespace llvm { +namespace SanitizerCommonUtils { + +// Spir memory address space +constexpr unsigned kSpirOffloadPrivateAS = 0; +constexpr unsigned kSpirOffloadGlobalAS = 1; +constexpr unsigned kSpirOffloadConstantAS = 2; +constexpr unsigned kSpirOffloadLocalAS = 3; +constexpr unsigned kSpirOffloadGenericAS = 4; + +TargetExtType *getTargetExtType(Type *Ty); +bool isJointMatrixAccess(Value *V); +bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I, + bool InstrumentLocalPtr, + bool InstrumentPrivatePtr); + +} // namespace SanitizerCommonUtils +} // namespace llvm + +#endif // LLVM_TRANSFORMS_INSTRUMENTATION_SANITIZERCOMMONUTILS_H diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 710c1ac05b4b3..da0405f7a7a9e 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -74,6 +74,7 @@ #include "llvm/TargetParser/Triple.h" #include "llvm/Transforms/Instrumentation/AddressSanitizerCommon.h" #include "llvm/Transforms/Instrumentation/AddressSanitizerOptions.h" +#include "llvm/Transforms/Instrumentation/SanitizerCommonUtils.h" #include "llvm/Transforms/Utils/ASanStackFrameLayout.h" #include "llvm/Transforms/Utils/BasicBlockUtils.h" #include "llvm/Transforms/Utils/Instrumentation.h" @@ -1653,49 +1654,6 @@ static bool isUnsupportedAMDGPUAddrspace(Value *Addr) { return false; } -static TargetExtType *getTargetExtType(Type *Ty) { - if (auto *TargetTy = dyn_cast(Ty)) - return TargetTy; - - if (Ty->isVectorTy()) - return getTargetExtType(Ty->getScalarType()); - - if (Ty->isArrayTy()) - return getTargetExtType(Ty->getArrayElementType()); - - if (auto *STy = dyn_cast(Ty)) { - for (unsigned int i = 0; i < STy->getNumElements(); i++) - if (auto *TargetTy = getTargetExtType(STy->getElementType(i))) - return TargetTy; - return nullptr; - } - - return nullptr; -} - -// Skip pointer operand that is sycl joint matrix access since it isn't from -// user code, e.g. %call: -// clang-format off -// %a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8 -// %0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0 -// %call = call spir_func ptr -// @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0) -// %1 = load float, ptr %call, align 4 -// store float %1, ptr %call, align 4 -// clang-format on -static bool isJointMatrixAccess(Value *V) { - auto *ActualV = V->stripInBoundsOffsets(); - if (auto *CI = dyn_cast(ActualV)) { - for (Value *Op : CI->args()) { - if (auto *AI = dyn_cast(Op->stripInBoundsOffsets())) - if (auto *TargetTy = getTargetExtType(AI->getAllocatedType())) - return TargetTy->getName().starts_with("spirv.") && - TargetTy->getName().contains("Matrix"); - } - } - return false; -} - static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { // Non image scope device globals are implemented by device USM, and the // out-of-bounds check for them will be done by sanitizer USM part. So we @@ -1714,52 +1672,6 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false"); } -static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) { - // Skip SPIR-V built-in varibles - auto *OrigValue = Addr->stripInBoundsOffsets(); - if (OrigValue->getName().starts_with("__spirv_BuiltIn")) - return true; - - GlobalVariable *GV = dyn_cast(OrigValue); - if (GV && isUnsupportedDeviceGlobal(GV)) - return true; - - // Ignore load/store for target ext type since we can't know exactly what size - // it is. - if (auto *SI = dyn_cast(Inst)) - if (getTargetExtType(SI->getValueOperand()->getType()) || - isJointMatrixAccess(SI->getPointerOperand())) - return true; - - if (auto *LI = dyn_cast(Inst)) - if (getTargetExtType(Inst->getType()) || - isJointMatrixAccess(LI->getPointerOperand())) - return true; - - Type *PtrTy = cast(Addr->getType()->getScalarType()); - switch (PtrTy->getPointerAddressSpace()) { - case kSpirOffloadPrivateAS: { - if (!ClSpirOffloadPrivates) - return true; - // Skip kernel arguments - return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL && - isa(Addr); - } - case kSpirOffloadGlobalAS: { - return !ClSpirOffloadGlobals; - } - case kSpirOffloadLocalAS: { - if (!ClSpirOffloadLocals) - return true; - return Addr->getName().starts_with("__Asan"); - } - case kSpirOffloadGenericAS: { - return !ClSpirOffloadGenerics; - } - } - return true; -} - void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, Value *Addr, SmallVectorImpl &Args) { @@ -1916,7 +1828,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) { !(SSGI && SSGI->isSafe(AI)) && // ignore alloc contains target ext type since we can't know exactly what // size it is. - !getTargetExtType(AI.getAllocatedType())); + !SanitizerCommonUtils::getTargetExtType(AI.getAllocatedType())); It->second = IsInteresting; return IsInteresting; @@ -1925,7 +1837,8 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) { bool AddressSanitizer::ignoreAccess(Instruction *Inst, Value *Ptr) { // SPIR has its own rules to filter the instrument accesses if (TargetTriple.isSPIROrSPIRV()) { - if (isUnsupportedSPIRAccess(Ptr, Inst)) + if (SanitizerCommonUtils::isUnsupportedSPIRAccess( + Ptr, Inst, ClSpirOffloadLocals, ClSpirOffloadPrivates)) return true; } else { // Instrument accesses from different address spaces only for AMDGPU. diff --git a/llvm/lib/Transforms/Instrumentation/CMakeLists.txt b/llvm/lib/Transforms/Instrumentation/CMakeLists.txt index ffd940fa85942..b42b349e4be4f 100644 --- a/llvm/lib/Transforms/Instrumentation/CMakeLists.txt +++ b/llvm/lib/Transforms/Instrumentation/CMakeLists.txt @@ -26,6 +26,7 @@ add_llvm_component_library(LLVMInstrumentation TypeSanitizer.cpp HWAddressSanitizer.cpp RealtimeSanitizer.cpp + SanitizerCommonUtils.cpp ADDITIONAL_HEADER_DIRS ${LLVM_MAIN_INCLUDE_DIR}/llvm/Transforms diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 50c66d03b855c..8e4086a7148dc 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -200,6 +200,7 @@ #include "llvm/Support/Path.h" #include "llvm/Support/raw_ostream.h" #include "llvm/TargetParser/Triple.h" +#include "llvm/Transforms/Instrumentation/SanitizerCommonUtils.h" #include "llvm/Transforms/Utils/BasicBlockUtils.h" #include "llvm/Transforms/Utils/Instrumentation.h" #include "llvm/Transforms/Utils/Local.h" @@ -1734,31 +1735,6 @@ static unsigned TypeSizeToSizeIndex(TypeSize TS) { return Log2_32_Ceil((TypeSizeFixed + 7) / 8); } -static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) { - if (isa(Addr) && - cast(Addr)->getMetadata(LLVMContext::MD_nosanitize)) { - return true; - } - - // Skip SPIR-V built-in varibles - auto *OrigValue = Addr->stripInBoundsOffsets(); - assert(OrigValue != nullptr); - if (OrigValue->getName().starts_with("__spirv_BuiltIn")) - return true; - - Type *PtrTy = cast(Addr->getType()->getScalarType()); - switch (PtrTy->getPointerAddressSpace()) { - case kSpirOffloadPrivateAS: - return !ClSpirOffloadPrivates; - case kSpirOffloadLocalAS: - return !ClSpirOffloadLocals; - case kSpirOffloadGenericAS: - return false; - } - - return false; -} - static void setNoSanitizedMetadataSPIR(Instruction &I) { const Value *Addr = nullptr; if (const auto *LI = dyn_cast(&I)) @@ -1810,7 +1786,8 @@ static void setNoSanitizedMetadataSPIR(Instruction &I) { } } - if (Addr && isUnsupportedSPIRAccess(Addr, &I)) + if (Addr && SanitizerCommonUtils::isUnsupportedSPIRAccess( + Addr, &I, ClSpirOffloadLocals, ClSpirOffloadPrivates)) I.setNoSanitizeMetadata(); } diff --git a/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp b/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp new file mode 100644 index 0000000000000..358ad2a2404b9 --- /dev/null +++ b/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp @@ -0,0 +1,103 @@ +//===- SanitizerCommonUtils.cpp- Sanitizer commnon utils------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file declares common infrastructure for Sanitizer. +// +//===----------------------------------------------------------------------===// + +#include "llvm/Transforms/Instrumentation/SanitizerCommonUtils.h" + +using namespace llvm; + +namespace llvm { +namespace SanitizerCommonUtils { + +TargetExtType *getTargetExtType(Type *Ty) { + if (auto *TargetTy = dyn_cast(Ty)) + return TargetTy; + + if (Ty->isVectorTy()) + return getTargetExtType(Ty->getScalarType()); + + if (Ty->isArrayTy()) + return getTargetExtType(Ty->getArrayElementType()); + + if (auto *STy = dyn_cast(Ty)) { + for (unsigned int i = 0; i < STy->getNumElements(); i++) + if (auto *TargetTy = getTargetExtType(STy->getElementType(i))) + return TargetTy; + return nullptr; + } + + return nullptr; +} + +// Skip pointer operand that is sycl joint matrix access since it isn't from +// user code, e.g. %call: +// clang-format off +// %a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8 +// %0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0 +// %call = call spir_func ptr +// @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0) +// %1 = load float, ptr %call, align 4 +// store float %1, ptr %call, align 4 +// clang-format on +bool isJointMatrixAccess(Value *V) { + auto *ActualV = V->stripInBoundsOffsets(); + if (auto *CI = dyn_cast(ActualV)) { + for (Value *Op : CI->args()) { + if (auto *AI = dyn_cast(Op->stripInBoundsOffsets())) + if (auto *TargetTy = getTargetExtType(AI->getAllocatedType())) + return TargetTy->getName().starts_with("spirv.") && + TargetTy->getName().contains("Matrix"); + } + } + return false; +} + +bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I, + bool InstrumentLocalPtr, + bool InstrumentPrivatePtr) { + if (isa(Addr) && + cast(Addr)->getMetadata(LLVMContext::MD_nosanitize)) { + return true; + } + + // Skip SPIR-V built-in varibles + auto *OrigValue = Addr->stripInBoundsOffsets(); + assert(OrigValue != nullptr); + if (OrigValue->getName().starts_with("__spirv_BuiltIn")) + return true; + + // Ignore load/store for target ext type since we can't know exactly what size + // it is. + if (auto *SI = dyn_cast(I)) + if (getTargetExtType(SI->getValueOperand()->getType()) || + isJointMatrixAccess(SI->getPointerOperand())) + return true; + + if (auto *LI = dyn_cast(I)) + if (getTargetExtType(I->getType()) || + isJointMatrixAccess(LI->getPointerOperand())) + return true; + + Type *PtrTy = cast(Addr->getType()->getScalarType()); + switch (PtrTy->getPointerAddressSpace()) { + case kSpirOffloadPrivateAS: + return !InstrumentPrivatePtr; + case kSpirOffloadLocalAS: + return !InstrumentLocalPtr; + case kSpirOffloadGenericAS: + return false; + } + + return false; +} + +} // namespace SanitizerCommonUtils +} // namespace llvm diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll new file mode 100644 index 0000000000000..845589c161314 --- /dev/null +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll @@ -0,0 +1,44 @@ +; 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" + +%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1) } + +; Function Attrs: sanitize_address +define spir_kernel void @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE() #0 { +entry: +; CHECK-LABEL-DAG: @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE +; CHECK-NOT: MyAlloc + %a = alloca [2 x %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix"], i32 0, align 8 + br label %for.cond10.i + +for.cond10.i: ; preds = %for.cond10.i, %entry + %0 = load target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1), ptr null, align 8 + store target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) zeroinitializer, ptr null, align 8 +; CHECK-NOT: call void @asan_load +; CHECK-NOT: call void @asan_store + br label %for.cond10.i +} + +; Function Attrs: sanitize_address +define spir_kernel void @AccessChain() #0 { +entry: +; CHECK-LABEL-DAG: @AccessChain + %a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8 + %0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0 + %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) + %1 = getelementptr inbounds { i16 }, ptr %call.i35, i64 0, i32 0 +; CHECK-NOT: call void @__asan_load +; CHECK-NOT: call void @__asan_store + %2 = load i16, ptr %1, align 4 + %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) + %3 = getelementptr inbounds { i16 }, ptr %call.i42, i64 0, i32 0 + store i16 %2, ptr %3, align 4 + ret void +} + +declare spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr, i64) + +attributes #0 = { sanitize_address } + From fc2f34e352443ac4e77dfc50a0adbd067516e34b Mon Sep 17 00:00:00 2001 From: haonanya1 Date: Sun, 23 Mar 2025 23:01:40 -0700 Subject: [PATCH 2/6] Remove unused var --- .../Instrumentation/AddressSanitizer.cpp | 71 ++++++++++--------- 1 file changed, 38 insertions(+), 33 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index da0405f7a7a9e..d3b805b2a0fd2 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -196,13 +196,6 @@ constexpr size_t kAccessSizeIndexMask = 0xf; constexpr size_t kIsWriteShift = 5; constexpr size_t kIsWriteMask = 0x1; -// Spir memory address space -static constexpr unsigned kSpirOffloadPrivateAS = 0; -static constexpr unsigned kSpirOffloadGlobalAS = 1; -static constexpr unsigned kSpirOffloadConstantAS = 2; -static constexpr unsigned kSpirOffloadLocalAS = 3; -static constexpr unsigned kSpirOffloadGenericAS = 4; - // Command-line flags. static cl::opt ClEnableKasan( @@ -872,8 +865,9 @@ struct AddressSanitizer { Value *SizeArgument, uint32_t Exp, RuntimeCallInserter &RTCI); void instrumentMemIntrinsic(MemIntrinsic *MI, RuntimeCallInserter &RTCI); - Value *memToShadow(Value *Shadow, IRBuilder<> &IRB, - uint32_t AddressSpace = kSpirOffloadPrivateAS); + Value *memToShadow( + Value *Shadow, IRBuilder<> &IRB, + uint32_t AddressSpace = SanitizerCommonUtils::kSpirOffloadPrivateAS); bool suppressInstrumentationSiteForDebug(int &Instrumented); bool instrumentFunction(Function &F, const TargetLibraryInfo *TLI); bool maybeInsertAsanInitAtFunctionEntry(Function &F); @@ -1403,8 +1397,9 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, SpirFixupKernels.emplace_back(&F); auto KernelName = F.getName(); - auto *KernelNameGV = GetOrCreateGlobalString( - M, "__asan_kernel", KernelName, kSpirOffloadConstantAS); + auto *KernelNameGV = + GetOrCreateGlobalString(M, "__asan_kernel", KernelName, + SanitizerCommonUtils::kSpirOffloadConstantAS); SpirKernelsMetadata.emplace_back(ConstantStruct::get( StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy), ConstantInt::get(IntptrTy, KernelName.size()))); @@ -1439,7 +1434,8 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, } // New argument: uintptr_t as(1)*, which is allocated in shared USM buffer - Types.push_back(llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS)); + Types.push_back(llvm::PointerType::get( + IntptrTy, SanitizerCommonUtils::kSpirOffloadGlobalAS)); FunctionType *NewFTy = FunctionType::get(F->getReturnType(), Types, false); @@ -1497,9 +1493,9 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, FixupMetadata("kernel_arg_exclusive_ptr", ConstantAsMetadata::get(Builder.getFalse())); - FixupMetadata( - "kernel_arg_addr_space", - ConstantAsMetadata::get(Builder.getInt32(kSpirOffloadGlobalAS))); + FixupMetadata("kernel_arg_addr_space", + ConstantAsMetadata::get(Builder.getInt32( + SanitizerCommonUtils::kSpirOffloadGlobalAS))); FixupMetadata("kernel_arg_access_qual", MDString::get(M.getContext(), "read_write")); FixupMetadata("kernel_arg_type", MDString::get(M.getContext(), "void*")); @@ -1680,15 +1676,16 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, auto &Loc = InsertBefore->getDebugLoc(); // SPIR constant address space - PointerType *ConstASPtrTy = - llvm::PointerType::get(Type::getInt8Ty(C), kSpirOffloadConstantAS); + PointerType *ConstASPtrTy = llvm::PointerType::get( + Type::getInt8Ty(C), SanitizerCommonUtils::kSpirOffloadConstantAS); // File & Line if (Loc) { llvm::SmallString<128> Source = Loc->getDirectory(); sys::path::append(Source, Loc->getFilename()); - auto *FileNameGV = GetOrCreateGlobalString(*M, "__asan_file", Source, - kSpirOffloadConstantAS); + auto *FileNameGV = + GetOrCreateGlobalString(*M, "__asan_file", Source, + SanitizerCommonUtils::kSpirOffloadConstantAS); Args.push_back(ConstantExpr::getPointerCast(FileNameGV, ConstASPtrTy)); Args.push_back(ConstantInt::get(Type::getInt32Ty(C), Loc.getLine())); } else { @@ -1698,8 +1695,9 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, // Function auto FuncName = InsertBefore->getFunction()->getName(); - auto *FuncNameGV = GetOrCreateGlobalString( - *M, "__asan_func", demangle(FuncName), kSpirOffloadConstantAS); + auto *FuncNameGV = + GetOrCreateGlobalString(*M, "__asan_func", demangle(FuncName), + SanitizerCommonUtils::kSpirOffloadConstantAS); Args.push_back(ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy)); } @@ -1735,7 +1733,8 @@ bool AddressSanitizer::instrumentSyclDynamicLocalMemory( SmallVector LocalArgs; for (auto &Arg : F.args()) { Type *PtrTy = dyn_cast(Arg.getType()->getScalarType()); - if (PtrTy && PtrTy->getPointerAddressSpace() == kSpirOffloadLocalAS) + if (PtrTy && PtrTy->getPointerAddressSpace() == + SanitizerCommonUtils::kSpirOffloadLocalAS) LocalArgs.push_back(&Arg); } @@ -1780,8 +1779,8 @@ void AddressSanitizer::instrumentInitAsanLaunchInfo( // FIXME: if the initial value of "__AsanLaunchInfo" is zero, we'll not need // this step initializeCallbacks(TLI); - IRB.CreateStore(ConstantPointerNull::get( - llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS)), + IRB.CreateStore(ConstantPointerNull::get(llvm::PointerType::get( + IntptrTy, SanitizerCommonUtils::kSpirOffloadGlobalAS)), AsanLaunchInfo); } @@ -2802,7 +2801,8 @@ void ModuleAddressSanitizer::instrumentDeviceGlobal(IRBuilder<> &IRB) { SmallVector GlobalsToRemove; SmallVector DeviceGlobalMetadata; - Type *IntptrTy = M.getDataLayout().getIntPtrType(*C, kSpirOffloadGlobalAS); + Type *IntptrTy = M.getDataLayout().getIntPtrType( + *C, SanitizerCommonUtils::kSpirOffloadGlobalAS); // Device global meta data is described by a structure // size_t device_global_size @@ -3560,8 +3560,8 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { // char* func // ) if (TargetTriple.isSPIROrSPIRV()) { - auto *Int8PtrTy = - llvm::PointerType::get(Type::getInt8Ty(*C), kSpirOffloadConstantAS); + auto *Int8PtrTy = llvm::PointerType::get( + Type::getInt8Ty(*C), SanitizerCommonUtils::kSpirOffloadConstantAS); Args1.push_back(Int8PtrTy); // file Args1.push_back(Type::getInt32Ty(*C)); // line @@ -3662,11 +3662,16 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { AsanLaunchInfo = M.getOrInsertGlobal( "__AsanLaunchInfo", - llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS), [&] { + llvm::PointerType::get(IntptrTy, + SanitizerCommonUtils::kSpirOffloadGlobalAS), + [&] { return new GlobalVariable( - M, llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS), false, - GlobalVariable::ExternalLinkage, nullptr, "__AsanLaunchInfo", - nullptr, GlobalVariable::NotThreadLocal, kSpirOffloadLocalAS); + M, + llvm::PointerType::get( + IntptrTy, SanitizerCommonUtils::kSpirOffloadGlobalAS), + false, GlobalVariable::ExternalLinkage, nullptr, + "__AsanLaunchInfo", nullptr, GlobalVariable::NotThreadLocal, + SanitizerCommonUtils::kSpirOffloadLocalAS); }); AsanMemToShadow = M.getOrInsertFunction(kAsanMemToShadow, IntptrTy, @@ -4451,8 +4456,8 @@ void FunctionStackPoisoner::processStaticAllocas() { const auto &ShadowAfterScope = GetShadowBytesAfterScope(SVD, L); // Poison the stack red zones at the entry. - Value *ShadowBase = - ASan.memToShadow(LocalStackBase, IRB, kSpirOffloadPrivateAS); + Value *ShadowBase = ASan.memToShadow( + LocalStackBase, IRB, SanitizerCommonUtils::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, From 62acbce8ba6821f32d4934fed7bd72befe91ba04 Mon Sep 17 00:00:00 2001 From: haonanya1 Date: Sun, 23 Mar 2025 23:44:22 -0700 Subject: [PATCH 3/6] Update lit --- .../Instrumentation/SanitizerCommonUtils.h | 4 -- .../Instrumentation/AddressSanitizer.cpp | 50 ++++++++++++++++++- .../Instrumentation/MemorySanitizer.cpp | 41 ++++++++++++++- .../Instrumentation/SanitizerCommonUtils.cpp | 40 --------------- .../SPIRV/ignore_target_ext_type.ll | 49 ++++++------------ 5 files changed, 102 insertions(+), 82 deletions(-) diff --git a/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h b/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h index 011bd449b0629..947ccfbfadcfb 100644 --- a/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h +++ b/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h @@ -29,10 +29,6 @@ constexpr unsigned kSpirOffloadGenericAS = 4; TargetExtType *getTargetExtType(Type *Ty); bool isJointMatrixAccess(Value *V); -bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I, - bool InstrumentLocalPtr, - bool InstrumentPrivatePtr); - } // namespace SanitizerCommonUtils } // namespace llvm diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index d3b805b2a0fd2..845ed7845da44 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -1668,6 +1668,53 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false"); } +static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) { + // Skip SPIR-V built-in varibles + auto *OrigValue = Addr->stripInBoundsOffsets(); + if (OrigValue->getName().starts_with("__spirv_BuiltIn")) + return true; + + GlobalVariable *GV = dyn_cast(OrigValue); + if (GV && isUnsupportedDeviceGlobal(GV)) + return true; + + // Ignore load/store for target ext type since we can't know exactly what size + // it is. + if (auto *SI = dyn_cast(Inst)) + if (SanitizerCommonUtils::getTargetExtType( + SI->getValueOperand()->getType()) || + SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) + return true; + + if (auto *LI = dyn_cast(Inst)) + if (SanitizerCommonUtils::getTargetExtType(Inst->getType()) || + SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) + return true; + + Type *PtrTy = cast(Addr->getType()->getScalarType()); + switch (PtrTy->getPointerAddressSpace()) { + case SanitizerCommonUtils::kSpirOffloadPrivateAS: { + if (!ClSpirOffloadPrivates) + return true; + // Skip kernel arguments + return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL && + isa(Addr); + } + case SanitizerCommonUtils::kSpirOffloadGlobalAS: { + return !ClSpirOffloadGlobals; + } + case SanitizerCommonUtils::kSpirOffloadLocalAS: { + if (!ClSpirOffloadLocals) + return true; + return Addr->getName().starts_with("__Asan"); + } + case SanitizerCommonUtils::kSpirOffloadGenericAS: { + return !ClSpirOffloadGenerics; + } + } + return true; +} + void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, Value *Addr, SmallVectorImpl &Args) { @@ -1836,8 +1883,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) { bool AddressSanitizer::ignoreAccess(Instruction *Inst, Value *Ptr) { // SPIR has its own rules to filter the instrument accesses if (TargetTriple.isSPIROrSPIRV()) { - if (SanitizerCommonUtils::isUnsupportedSPIRAccess( - Ptr, Inst, ClSpirOffloadLocals, ClSpirOffloadPrivates)) + if (isUnsupportedSPIRAccess(Ptr, Inst)) return true; } else { // Instrument accesses from different address spaces only for AMDGPU. diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 8e4086a7148dc..1f15550898d55 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -1735,6 +1735,44 @@ static unsigned TypeSizeToSizeIndex(TypeSize TS) { return Log2_32_Ceil((TypeSizeFixed + 7) / 8); } +static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) { + if (isa(Addr) && + cast(Addr)->getMetadata(LLVMContext::MD_nosanitize)) { + return true; + } + + // Skip SPIR-V built-in varibles + auto *OrigValue = Addr->stripInBoundsOffsets(); + assert(OrigValue != nullptr); + if (OrigValue->getName().starts_with("__spirv_BuiltIn")) + return true; + + // Ignore load/store for target ext type since we can't know exactly what size + // it is. + if (auto *SI = dyn_cast(I)) + if (SanitizerCommonUtils::getTargetExtType( + SI->getValueOperand()->getType()) || + SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) + return true; + + if (auto *LI = dyn_cast(I)) + if (SanitizerCommonUtils::getTargetExtType(I->getType()) || + SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) + return true; + + Type *PtrTy = cast(Addr->getType()->getScalarType()); + switch (PtrTy->getPointerAddressSpace()) { + case kSpirOffloadPrivateAS: + return !ClSpirOffloadPrivates; + case kSpirOffloadLocalAS: + return !ClSpirOffloadLocals; + case kSpirOffloadGenericAS: + return false; + } + + return false; +} + static void setNoSanitizedMetadataSPIR(Instruction &I) { const Value *Addr = nullptr; if (const auto *LI = dyn_cast(&I)) @@ -1786,8 +1824,7 @@ static void setNoSanitizedMetadataSPIR(Instruction &I) { } } - if (Addr && SanitizerCommonUtils::isUnsupportedSPIRAccess( - Addr, &I, ClSpirOffloadLocals, ClSpirOffloadPrivates)) + if (Addr && isUnsupportedSPIRAccess(Addr, &I)) I.setNoSanitizeMetadata(); } diff --git a/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp b/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp index 358ad2a2404b9..64fcfb153cf15 100644 --- a/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp +++ b/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp @@ -59,45 +59,5 @@ bool isJointMatrixAccess(Value *V) { } return false; } - -bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I, - bool InstrumentLocalPtr, - bool InstrumentPrivatePtr) { - if (isa(Addr) && - cast(Addr)->getMetadata(LLVMContext::MD_nosanitize)) { - return true; - } - - // Skip SPIR-V built-in varibles - auto *OrigValue = Addr->stripInBoundsOffsets(); - assert(OrigValue != nullptr); - if (OrigValue->getName().starts_with("__spirv_BuiltIn")) - return true; - - // Ignore load/store for target ext type since we can't know exactly what size - // it is. - if (auto *SI = dyn_cast(I)) - if (getTargetExtType(SI->getValueOperand()->getType()) || - isJointMatrixAccess(SI->getPointerOperand())) - return true; - - if (auto *LI = dyn_cast(I)) - if (getTargetExtType(I->getType()) || - isJointMatrixAccess(LI->getPointerOperand())) - return true; - - Type *PtrTy = cast(Addr->getType()->getScalarType()); - switch (PtrTy->getPointerAddressSpace()) { - case kSpirOffloadPrivateAS: - return !InstrumentPrivatePtr; - case kSpirOffloadLocalAS: - return !InstrumentLocalPtr; - case kSpirOffloadGenericAS: - return false; - } - - return false; -} - } // namespace SanitizerCommonUtils } // namespace llvm diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll index 845589c161314..abdf450a97571 100644 --- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll @@ -3,42 +3,23 @@ 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" -%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1) } +%"class.sycl::_V1::ext::oneapi::bfloat16" = type { i16 } +%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.CooperativeMatrixKHR", i16, 3, 16, 32, 0) } -; Function Attrs: sanitize_address -define spir_kernel void @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE() #0 { -entry: -; CHECK-LABEL-DAG: @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE -; CHECK-NOT: MyAlloc - %a = alloca [2 x %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix"], i32 0, align 8 - br label %for.cond10.i - -for.cond10.i: ; preds = %for.cond10.i, %entry - %0 = load target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1), ptr null, align 8 - store target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) zeroinitializer, ptr null, align 8 -; CHECK-NOT: call void @asan_load -; CHECK-NOT: call void @asan_store - br label %for.cond10.i -} +; CHECK-LABEL: @test +; CHECK-NOT: call i64 @__msan_get_shadow +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) -; Function Attrs: sanitize_address -define spir_kernel void @AccessChain() #0 { +define weak_odr dso_local spir_kernel void @test() { entry: -; CHECK-LABEL-DAG: @AccessChain - %a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8 - %0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0 - %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) - %1 = getelementptr inbounds { i16 }, ptr %call.i35, i64 0, i32 0 -; CHECK-NOT: call void @__asan_load -; CHECK-NOT: call void @__asan_store - %2 = load i16, ptr %1, align 4 - %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) - %3 = getelementptr inbounds { i16 }, ptr %call.i42, i64 0, i32 0 - store i16 %2, ptr %3, align 4 + %sub_a.i = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8 + %element.i = alloca %"class.sycl::_V1::ext::oneapi::bfloat16", align 2 + %0 = getelementptr inbounds { i16 }, ptr %element.i, i64 0, i32 0 + %spvm.i = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %sub_a.i, i64 0, i32 0 + %addrcast = addrspacecast ptr %spvm.i to ptr addrspace(4) + %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) + %gep = getelementptr inbounds nuw { i16 }, ptr addrspace(4) %call.i67, i64 0, i32 0 + %val = load i16, ptr %0, align 2 + store i16 %val, ptr addrspace(4) %gep, align 2 ret void } - -declare spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr, i64) - -attributes #0 = { sanitize_address } - From 7cac17b48da7e76527529f4d337a69e44a3865a9 Mon Sep 17 00:00:00 2001 From: haonanya1 Date: Mon, 24 Mar 2025 22:43:04 -0700 Subject: [PATCH 4/6] Change file name --- ...monUtils.h => SPIRVSanitizerCommonUtils.h} | 14 ++-- .../Instrumentation/AddressSanitizer.cpp | 68 ++++++++++--------- .../Transforms/Instrumentation/CMakeLists.txt | 2 +- .../Instrumentation/MemorySanitizer.cpp | 10 +-- ...tils.cpp => SPIRVSanitizerCommonUtils.cpp} | 10 +-- 5 files changed, 53 insertions(+), 51 deletions(-) rename llvm/include/llvm/Transforms/Instrumentation/{SanitizerCommonUtils.h => SPIRVSanitizerCommonUtils.h} (68%) rename llvm/lib/Transforms/Instrumentation/{SanitizerCommonUtils.cpp => SPIRVSanitizerCommonUtils.cpp} (87%) diff --git a/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h b/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h similarity index 68% rename from llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h rename to llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h index 947ccfbfadcfb..3a7827255fbfd 100644 --- a/llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h +++ b/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h @@ -1,4 +1,4 @@ -//===- SanitizerCommonUtils.h - Sanitizer commnon utils ---------*- C++ -*-===// +//===- SPIRVSanitizerCommonUtils.h - Commnon utils --------------*- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,11 +6,11 @@ // //===----------------------------------------------------------------------===// // -// This file declares common infrastructure for Sanitizer. +// This file declares common infrastructure for SPIRV Sanitizer. // //===----------------------------------------------------------------------===// -#ifndef LLVM_TRANSFORMS_INSTRUMENTATION_SANITIZERCOMMONUTILS_H -#define LLVM_TRANSFORMS_INSTRUMENTATION_SANITIZERCOMMONUTILS_H +#ifndef LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H +#define LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H #include "llvm/IR/DerivedTypes.h" #include "llvm/IR/Instructions.h" @@ -18,7 +18,7 @@ #include "llvm/IR/Value.h" namespace llvm { -namespace SanitizerCommonUtils { +namespace SPIRVSanitizerCommonUtils { // Spir memory address space constexpr unsigned kSpirOffloadPrivateAS = 0; @@ -29,7 +29,7 @@ constexpr unsigned kSpirOffloadGenericAS = 4; TargetExtType *getTargetExtType(Type *Ty); bool isJointMatrixAccess(Value *V); -} // namespace SanitizerCommonUtils +} // namespace SPIRVSanitizerCommonUtils } // namespace llvm -#endif // LLVM_TRANSFORMS_INSTRUMENTATION_SANITIZERCOMMONUTILS_H +#endif // LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 845ed7845da44..a1a5dd509f894 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -74,7 +74,7 @@ #include "llvm/TargetParser/Triple.h" #include "llvm/Transforms/Instrumentation/AddressSanitizerCommon.h" #include "llvm/Transforms/Instrumentation/AddressSanitizerOptions.h" -#include "llvm/Transforms/Instrumentation/SanitizerCommonUtils.h" +#include "llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h" #include "llvm/Transforms/Utils/ASanStackFrameLayout.h" #include "llvm/Transforms/Utils/BasicBlockUtils.h" #include "llvm/Transforms/Utils/Instrumentation.h" @@ -867,7 +867,7 @@ struct AddressSanitizer { void instrumentMemIntrinsic(MemIntrinsic *MI, RuntimeCallInserter &RTCI); Value *memToShadow( Value *Shadow, IRBuilder<> &IRB, - uint32_t AddressSpace = SanitizerCommonUtils::kSpirOffloadPrivateAS); + uint32_t AddressSpace = SPIRVSanitizerCommonUtils::kSpirOffloadPrivateAS); bool suppressInstrumentationSiteForDebug(int &Instrumented); bool instrumentFunction(Function &F, const TargetLibraryInfo *TLI); bool maybeInsertAsanInitAtFunctionEntry(Function &F); @@ -1397,9 +1397,9 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, SpirFixupKernels.emplace_back(&F); auto KernelName = F.getName(); - auto *KernelNameGV = - GetOrCreateGlobalString(M, "__asan_kernel", KernelName, - SanitizerCommonUtils::kSpirOffloadConstantAS); + auto *KernelNameGV = GetOrCreateGlobalString( + M, "__asan_kernel", KernelName, + SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); SpirKernelsMetadata.emplace_back(ConstantStruct::get( StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy), ConstantInt::get(IntptrTy, KernelName.size()))); @@ -1435,7 +1435,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, // New argument: uintptr_t as(1)*, which is allocated in shared USM buffer Types.push_back(llvm::PointerType::get( - IntptrTy, SanitizerCommonUtils::kSpirOffloadGlobalAS)); + IntptrTy, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS)); FunctionType *NewFTy = FunctionType::get(F->getReturnType(), Types, false); @@ -1495,7 +1495,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, FixupMetadata("kernel_arg_addr_space", ConstantAsMetadata::get(Builder.getInt32( - SanitizerCommonUtils::kSpirOffloadGlobalAS))); + SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS))); FixupMetadata("kernel_arg_access_qual", MDString::get(M.getContext(), "read_write")); FixupMetadata("kernel_arg_type", MDString::get(M.getContext(), "void*")); @@ -1681,34 +1681,34 @@ static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) { // Ignore load/store for target ext type since we can't know exactly what size // it is. if (auto *SI = dyn_cast(Inst)) - if (SanitizerCommonUtils::getTargetExtType( + if (SPIRVSanitizerCommonUtils::getTargetExtType( SI->getValueOperand()->getType()) || - SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) + SPIRVSanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) return true; if (auto *LI = dyn_cast(Inst)) - if (SanitizerCommonUtils::getTargetExtType(Inst->getType()) || - SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) + if (SPIRVSanitizerCommonUtils::getTargetExtType(Inst->getType()) || + SPIRVSanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) return true; Type *PtrTy = cast(Addr->getType()->getScalarType()); switch (PtrTy->getPointerAddressSpace()) { - case SanitizerCommonUtils::kSpirOffloadPrivateAS: { + case SPIRVSanitizerCommonUtils::kSpirOffloadPrivateAS: { if (!ClSpirOffloadPrivates) return true; // Skip kernel arguments return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL && isa(Addr); } - case SanitizerCommonUtils::kSpirOffloadGlobalAS: { + case SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS: { return !ClSpirOffloadGlobals; } - case SanitizerCommonUtils::kSpirOffloadLocalAS: { + case SPIRVSanitizerCommonUtils::kSpirOffloadLocalAS: { if (!ClSpirOffloadLocals) return true; return Addr->getName().starts_with("__Asan"); } - case SanitizerCommonUtils::kSpirOffloadGenericAS: { + case SPIRVSanitizerCommonUtils::kSpirOffloadGenericAS: { return !ClSpirOffloadGenerics; } } @@ -1724,15 +1724,15 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, // SPIR constant address space PointerType *ConstASPtrTy = llvm::PointerType::get( - Type::getInt8Ty(C), SanitizerCommonUtils::kSpirOffloadConstantAS); + Type::getInt8Ty(C), SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); // File & Line if (Loc) { llvm::SmallString<128> Source = Loc->getDirectory(); sys::path::append(Source, Loc->getFilename()); - auto *FileNameGV = - GetOrCreateGlobalString(*M, "__asan_file", Source, - SanitizerCommonUtils::kSpirOffloadConstantAS); + auto *FileNameGV = GetOrCreateGlobalString( + *M, "__asan_file", Source, + SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); Args.push_back(ConstantExpr::getPointerCast(FileNameGV, ConstASPtrTy)); Args.push_back(ConstantInt::get(Type::getInt32Ty(C), Loc.getLine())); } else { @@ -1742,9 +1742,9 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, // Function auto FuncName = InsertBefore->getFunction()->getName(); - auto *FuncNameGV = - GetOrCreateGlobalString(*M, "__asan_func", demangle(FuncName), - SanitizerCommonUtils::kSpirOffloadConstantAS); + auto *FuncNameGV = GetOrCreateGlobalString( + *M, "__asan_func", demangle(FuncName), + SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); Args.push_back(ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy)); } @@ -1781,7 +1781,7 @@ bool AddressSanitizer::instrumentSyclDynamicLocalMemory( for (auto &Arg : F.args()) { Type *PtrTy = dyn_cast(Arg.getType()->getScalarType()); if (PtrTy && PtrTy->getPointerAddressSpace() == - SanitizerCommonUtils::kSpirOffloadLocalAS) + SPIRVSanitizerCommonUtils::kSpirOffloadLocalAS) LocalArgs.push_back(&Arg); } @@ -1826,9 +1826,10 @@ void AddressSanitizer::instrumentInitAsanLaunchInfo( // FIXME: if the initial value of "__AsanLaunchInfo" is zero, we'll not need // this step initializeCallbacks(TLI); - IRB.CreateStore(ConstantPointerNull::get(llvm::PointerType::get( - IntptrTy, SanitizerCommonUtils::kSpirOffloadGlobalAS)), - AsanLaunchInfo); + IRB.CreateStore( + ConstantPointerNull::get(llvm::PointerType::get( + IntptrTy, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS)), + AsanLaunchInfo); } // Instrument memset/memmove/memcpy @@ -1874,7 +1875,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) { !(SSGI && SSGI->isSafe(AI)) && // ignore alloc contains target ext type since we can't know exactly what // size it is. - !SanitizerCommonUtils::getTargetExtType(AI.getAllocatedType())); + !SPIRVSanitizerCommonUtils::getTargetExtType(AI.getAllocatedType())); It->second = IsInteresting; return IsInteresting; @@ -2848,7 +2849,7 @@ void ModuleAddressSanitizer::instrumentDeviceGlobal(IRBuilder<> &IRB) { SmallVector DeviceGlobalMetadata; Type *IntptrTy = M.getDataLayout().getIntPtrType( - *C, SanitizerCommonUtils::kSpirOffloadGlobalAS); + *C, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS); // Device global meta data is described by a structure // size_t device_global_size @@ -3607,7 +3608,8 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { // ) if (TargetTriple.isSPIROrSPIRV()) { auto *Int8PtrTy = llvm::PointerType::get( - Type::getInt8Ty(*C), SanitizerCommonUtils::kSpirOffloadConstantAS); + Type::getInt8Ty(*C), + SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); Args1.push_back(Int8PtrTy); // file Args1.push_back(Type::getInt32Ty(*C)); // line @@ -3709,15 +3711,15 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { AsanLaunchInfo = M.getOrInsertGlobal( "__AsanLaunchInfo", llvm::PointerType::get(IntptrTy, - SanitizerCommonUtils::kSpirOffloadGlobalAS), + SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS), [&] { return new GlobalVariable( M, llvm::PointerType::get( - IntptrTy, SanitizerCommonUtils::kSpirOffloadGlobalAS), + IntptrTy, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS), false, GlobalVariable::ExternalLinkage, nullptr, "__AsanLaunchInfo", nullptr, GlobalVariable::NotThreadLocal, - SanitizerCommonUtils::kSpirOffloadLocalAS); + SPIRVSanitizerCommonUtils::kSpirOffloadLocalAS); }); AsanMemToShadow = M.getOrInsertFunction(kAsanMemToShadow, IntptrTy, @@ -4503,7 +4505,7 @@ void FunctionStackPoisoner::processStaticAllocas() { // Poison the stack red zones at the entry. Value *ShadowBase = ASan.memToShadow( - LocalStackBase, IRB, SanitizerCommonUtils::kSpirOffloadPrivateAS); + LocalStackBase, IRB, SPIRVSanitizerCommonUtils::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, diff --git a/llvm/lib/Transforms/Instrumentation/CMakeLists.txt b/llvm/lib/Transforms/Instrumentation/CMakeLists.txt index b42b349e4be4f..8a3af2bfd75ef 100644 --- a/llvm/lib/Transforms/Instrumentation/CMakeLists.txt +++ b/llvm/lib/Transforms/Instrumentation/CMakeLists.txt @@ -26,7 +26,7 @@ add_llvm_component_library(LLVMInstrumentation TypeSanitizer.cpp HWAddressSanitizer.cpp RealtimeSanitizer.cpp - SanitizerCommonUtils.cpp + SPIRVSanitizerCommonUtils.cpp ADDITIONAL_HEADER_DIRS ${LLVM_MAIN_INCLUDE_DIR}/llvm/Transforms diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 1f15550898d55..638cd51907151 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -200,7 +200,7 @@ #include "llvm/Support/Path.h" #include "llvm/Support/raw_ostream.h" #include "llvm/TargetParser/Triple.h" -#include "llvm/Transforms/Instrumentation/SanitizerCommonUtils.h" +#include "llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h" #include "llvm/Transforms/Utils/BasicBlockUtils.h" #include "llvm/Transforms/Utils/Instrumentation.h" #include "llvm/Transforms/Utils/Local.h" @@ -1750,14 +1750,14 @@ static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) { // Ignore load/store for target ext type since we can't know exactly what size // it is. if (auto *SI = dyn_cast(I)) - if (SanitizerCommonUtils::getTargetExtType( + if (SPIRVSanitizerCommonUtils::getTargetExtType( SI->getValueOperand()->getType()) || - SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) + SPIRVSanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) return true; if (auto *LI = dyn_cast(I)) - if (SanitizerCommonUtils::getTargetExtType(I->getType()) || - SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) + if (SPIRVSanitizerCommonUtils::getTargetExtType(I->getType()) || + SPIRVSanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) return true; Type *PtrTy = cast(Addr->getType()->getScalarType()); diff --git a/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp b/llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp similarity index 87% rename from llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp rename to llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp index 64fcfb153cf15..4d3a2e5ce9fbf 100644 --- a/llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp +++ b/llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp @@ -1,4 +1,4 @@ -//===- SanitizerCommonUtils.cpp- Sanitizer commnon utils------------------===// +//===- SPIRVSanitizerCommonUtils.cpp- SPIRV Sanitizer commnon utils ------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,16 +6,16 @@ // //===----------------------------------------------------------------------===// // -// This file declares common infrastructure for Sanitizer. +// This file defines common infrastructure for SPIRV Sanitizer. // //===----------------------------------------------------------------------===// -#include "llvm/Transforms/Instrumentation/SanitizerCommonUtils.h" +#include "llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h" using namespace llvm; namespace llvm { -namespace SanitizerCommonUtils { +namespace SPIRVSanitizerCommonUtils { TargetExtType *getTargetExtType(Type *Ty) { if (auto *TargetTy = dyn_cast(Ty)) @@ -59,5 +59,5 @@ bool isJointMatrixAccess(Value *V) { } return false; } -} // namespace SanitizerCommonUtils +} // namespace SPIRVSanitizerCommonUtils } // namespace llvm From 92978d42793d8110fe62e7574711298c4620dbb3 Mon Sep 17 00:00:00 2001 From: haonanya1 Date: Wed, 26 Mar 2025 01:42:33 -0700 Subject: [PATCH 5/6] Remove SPIRVSanitizerCommonUtils namespace --- .../SPIRVSanitizerCommonUtils.h | 4 - .../Instrumentation/AddressSanitizer.cpp | 83 ++++++++----------- .../Instrumentation/MemorySanitizer.cpp | 16 +--- .../SPIRVSanitizerCommonUtils.cpp | 4 +- 4 files changed, 39 insertions(+), 68 deletions(-) diff --git a/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h b/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h index 3a7827255fbfd..f6edabcc3bbf3 100644 --- a/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h +++ b/llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h @@ -13,13 +13,10 @@ #define LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H #include "llvm/IR/DerivedTypes.h" -#include "llvm/IR/Instructions.h" #include "llvm/IR/Type.h" #include "llvm/IR/Value.h" namespace llvm { -namespace SPIRVSanitizerCommonUtils { - // Spir memory address space constexpr unsigned kSpirOffloadPrivateAS = 0; constexpr unsigned kSpirOffloadGlobalAS = 1; @@ -29,7 +26,6 @@ constexpr unsigned kSpirOffloadGenericAS = 4; TargetExtType *getTargetExtType(Type *Ty); bool isJointMatrixAccess(Value *V); -} // namespace SPIRVSanitizerCommonUtils } // namespace llvm #endif // LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index a1a5dd509f894..2679b480b1a43 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -865,9 +865,8 @@ struct AddressSanitizer { Value *SizeArgument, uint32_t Exp, RuntimeCallInserter &RTCI); void instrumentMemIntrinsic(MemIntrinsic *MI, RuntimeCallInserter &RTCI); - Value *memToShadow( - Value *Shadow, IRBuilder<> &IRB, - uint32_t AddressSpace = SPIRVSanitizerCommonUtils::kSpirOffloadPrivateAS); + 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); @@ -1398,8 +1397,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, auto KernelName = F.getName(); auto *KernelNameGV = GetOrCreateGlobalString( - M, "__asan_kernel", KernelName, - SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); + M, "__asan_kernel", KernelName, kSpirOffloadConstantAS); SpirKernelsMetadata.emplace_back(ConstantStruct::get( StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy), ConstantInt::get(IntptrTy, KernelName.size()))); @@ -1434,8 +1432,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, } // New argument: uintptr_t as(1)*, which is allocated in shared USM buffer - Types.push_back(llvm::PointerType::get( - IntptrTy, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS)); + Types.push_back(llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS)); FunctionType *NewFTy = FunctionType::get(F->getReturnType(), Types, false); @@ -1493,9 +1490,9 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM, FixupMetadata("kernel_arg_exclusive_ptr", ConstantAsMetadata::get(Builder.getFalse())); - FixupMetadata("kernel_arg_addr_space", - ConstantAsMetadata::get(Builder.getInt32( - SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS))); + FixupMetadata( + "kernel_arg_addr_space", + ConstantAsMetadata::get(Builder.getInt32(kSpirOffloadGlobalAS))); FixupMetadata("kernel_arg_access_qual", MDString::get(M.getContext(), "read_write")); FixupMetadata("kernel_arg_type", MDString::get(M.getContext(), "void*")); @@ -1681,34 +1678,33 @@ static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) { // Ignore load/store for target ext type since we can't know exactly what size // it is. if (auto *SI = dyn_cast(Inst)) - if (SPIRVSanitizerCommonUtils::getTargetExtType( - SI->getValueOperand()->getType()) || - SPIRVSanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) + if (getTargetExtType(SI->getValueOperand()->getType()) || + isJointMatrixAccess(SI->getPointerOperand())) return true; if (auto *LI = dyn_cast(Inst)) - if (SPIRVSanitizerCommonUtils::getTargetExtType(Inst->getType()) || - SPIRVSanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) + if (getTargetExtType(Inst->getType()) || + isJointMatrixAccess(LI->getPointerOperand())) return true; Type *PtrTy = cast(Addr->getType()->getScalarType()); switch (PtrTy->getPointerAddressSpace()) { - case SPIRVSanitizerCommonUtils::kSpirOffloadPrivateAS: { + case kSpirOffloadPrivateAS: { if (!ClSpirOffloadPrivates) return true; // Skip kernel arguments return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL && isa(Addr); } - case SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS: { + case kSpirOffloadGlobalAS: { return !ClSpirOffloadGlobals; } - case SPIRVSanitizerCommonUtils::kSpirOffloadLocalAS: { + case kSpirOffloadLocalAS: { if (!ClSpirOffloadLocals) return true; return Addr->getName().starts_with("__Asan"); } - case SPIRVSanitizerCommonUtils::kSpirOffloadGenericAS: { + case kSpirOffloadGenericAS: { return !ClSpirOffloadGenerics; } } @@ -1723,16 +1719,15 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, auto &Loc = InsertBefore->getDebugLoc(); // SPIR constant address space - PointerType *ConstASPtrTy = llvm::PointerType::get( - Type::getInt8Ty(C), SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); + PointerType *ConstASPtrTy = + llvm::PointerType::get(Type::getInt8Ty(C), kSpirOffloadConstantAS); // File & Line if (Loc) { llvm::SmallString<128> Source = Loc->getDirectory(); sys::path::append(Source, Loc->getFilename()); - auto *FileNameGV = GetOrCreateGlobalString( - *M, "__asan_file", Source, - SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); + auto *FileNameGV = GetOrCreateGlobalString(*M, "__asan_file", Source, + kSpirOffloadConstantAS); Args.push_back(ConstantExpr::getPointerCast(FileNameGV, ConstASPtrTy)); Args.push_back(ConstantInt::get(Type::getInt32Ty(C), Loc.getLine())); } else { @@ -1743,8 +1738,7 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, // Function auto FuncName = InsertBefore->getFunction()->getName(); auto *FuncNameGV = GetOrCreateGlobalString( - *M, "__asan_func", demangle(FuncName), - SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); + *M, "__asan_func", demangle(FuncName), kSpirOffloadConstantAS); Args.push_back(ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy)); } @@ -1780,8 +1774,7 @@ bool AddressSanitizer::instrumentSyclDynamicLocalMemory( SmallVector LocalArgs; for (auto &Arg : F.args()) { Type *PtrTy = dyn_cast(Arg.getType()->getScalarType()); - if (PtrTy && PtrTy->getPointerAddressSpace() == - SPIRVSanitizerCommonUtils::kSpirOffloadLocalAS) + if (PtrTy && PtrTy->getPointerAddressSpace() == kSpirOffloadLocalAS) LocalArgs.push_back(&Arg); } @@ -1826,10 +1819,9 @@ void AddressSanitizer::instrumentInitAsanLaunchInfo( // FIXME: if the initial value of "__AsanLaunchInfo" is zero, we'll not need // this step initializeCallbacks(TLI); - IRB.CreateStore( - ConstantPointerNull::get(llvm::PointerType::get( - IntptrTy, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS)), - AsanLaunchInfo); + IRB.CreateStore(ConstantPointerNull::get( + llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS)), + AsanLaunchInfo); } // Instrument memset/memmove/memcpy @@ -1875,7 +1867,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) { !(SSGI && SSGI->isSafe(AI)) && // ignore alloc contains target ext type since we can't know exactly what // size it is. - !SPIRVSanitizerCommonUtils::getTargetExtType(AI.getAllocatedType())); + !getTargetExtType(AI.getAllocatedType())); It->second = IsInteresting; return IsInteresting; @@ -2848,8 +2840,7 @@ void ModuleAddressSanitizer::instrumentDeviceGlobal(IRBuilder<> &IRB) { SmallVector GlobalsToRemove; SmallVector DeviceGlobalMetadata; - Type *IntptrTy = M.getDataLayout().getIntPtrType( - *C, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS); + Type *IntptrTy = M.getDataLayout().getIntPtrType(*C, kSpirOffloadGlobalAS); // Device global meta data is described by a structure // size_t device_global_size @@ -3607,9 +3598,8 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { // char* func // ) if (TargetTriple.isSPIROrSPIRV()) { - auto *Int8PtrTy = llvm::PointerType::get( - Type::getInt8Ty(*C), - SPIRVSanitizerCommonUtils::kSpirOffloadConstantAS); + auto *Int8PtrTy = + llvm::PointerType::get(Type::getInt8Ty(*C), kSpirOffloadConstantAS); Args1.push_back(Int8PtrTy); // file Args1.push_back(Type::getInt32Ty(*C)); // line @@ -3710,16 +3700,11 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { AsanLaunchInfo = M.getOrInsertGlobal( "__AsanLaunchInfo", - llvm::PointerType::get(IntptrTy, - SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS), - [&] { + llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS), [&] { return new GlobalVariable( - M, - llvm::PointerType::get( - IntptrTy, SPIRVSanitizerCommonUtils::kSpirOffloadGlobalAS), - false, GlobalVariable::ExternalLinkage, nullptr, - "__AsanLaunchInfo", nullptr, GlobalVariable::NotThreadLocal, - SPIRVSanitizerCommonUtils::kSpirOffloadLocalAS); + M, llvm::PointerType::get(IntptrTy, kSpirOffloadGlobalAS), false, + GlobalVariable::ExternalLinkage, nullptr, "__AsanLaunchInfo", + nullptr, GlobalVariable::NotThreadLocal, kSpirOffloadLocalAS); }); AsanMemToShadow = M.getOrInsertFunction(kAsanMemToShadow, IntptrTy, @@ -4504,8 +4489,8 @@ void FunctionStackPoisoner::processStaticAllocas() { const auto &ShadowAfterScope = GetShadowBytesAfterScope(SVD, L); // Poison the stack red zones at the entry. - Value *ShadowBase = ASan.memToShadow( - LocalStackBase, IRB, SPIRVSanitizerCommonUtils::kSpirOffloadPrivateAS); + 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, diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 638cd51907151..9cd29df752ebb 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -589,13 +589,6 @@ static const PlatformMemoryMapParams Intel_SPIR_MemoryMapParams = { &Intel_SPIR64_MemoryMapParams, }; -// Spir memory address space -static constexpr unsigned kSpirOffloadPrivateAS = 0; -static constexpr unsigned kSpirOffloadGlobalAS = 1; -static constexpr unsigned kSpirOffloadConstantAS = 2; -static constexpr unsigned kSpirOffloadLocalAS = 3; -static constexpr unsigned kSpirOffloadGenericAS = 4; - namespace { class MemorySanitizerOnSpirv; @@ -1750,14 +1743,13 @@ static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) { // Ignore load/store for target ext type since we can't know exactly what size // it is. if (auto *SI = dyn_cast(I)) - if (SPIRVSanitizerCommonUtils::getTargetExtType( - SI->getValueOperand()->getType()) || - SPIRVSanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand())) + if (getTargetExtType(SI->getValueOperand()->getType()) || + isJointMatrixAccess(SI->getPointerOperand())) return true; if (auto *LI = dyn_cast(I)) - if (SPIRVSanitizerCommonUtils::getTargetExtType(I->getType()) || - SPIRVSanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand())) + if (getTargetExtType(I->getType()) || + isJointMatrixAccess(LI->getPointerOperand())) return true; Type *PtrTy = cast(Addr->getType()->getScalarType()); diff --git a/llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp b/llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp index 4d3a2e5ce9fbf..f08d931b96375 100644 --- a/llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp +++ b/llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp @@ -11,12 +11,11 @@ //===----------------------------------------------------------------------===// #include "llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h" +#include "llvm/IR/Instructions.h" using namespace llvm; namespace llvm { -namespace SPIRVSanitizerCommonUtils { - TargetExtType *getTargetExtType(Type *Ty) { if (auto *TargetTy = dyn_cast(Ty)) return TargetTy; @@ -59,5 +58,4 @@ bool isJointMatrixAccess(Value *V) { } return false; } -} // namespace SPIRVSanitizerCommonUtils } // namespace llvm From d8b55068555a5203ba6fd9c77e39ed50977ea640 Mon Sep 17 00:00:00 2001 From: "Yang, Haonan" Date: Thu, 17 Apr 2025 08:12:26 +0200 Subject: [PATCH 6/6] Add CODEOWNERS --- .github/CODEOWNERS | 2 ++ 1 file changed, 2 insertions(+) diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index e94ede9d2edc8..a8db78e42c403 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -202,9 +202,11 @@ llvm/include/llvm/Transforms/Instrumentation/AddressSanitizer.h @intel/dpcpp-san llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerCommon.h @intel/dpcpp-sanitizers-review llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerOptions.h @intel/dpcpp-sanitizers-review llvm/include/llvm/Transforms/Instrumentation/MemorySanitizer.h @intel/dpcpp-sanitizers-review +llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h @intel/dpcpp-sanitizers-review llvm/include/llvm/Transforms/Instrumentation/ThreadSanitizer.h @intel/dpcpp-sanitizers-review llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @intel/dpcpp-sanitizers-review llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @intel/dpcpp-sanitizers-review +llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp @intel/dpcpp-sanitizers-review llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp @intel/dpcpp-sanitizers-review llvm/test/Instrumentation/AddressSanitizer/ @intel/dpcpp-sanitizers-review llvm/test/Instrumentation/MemorySanitizer/ @intel/dpcpp-sanitizers-review