From 623bf14bf8cf3cef2da2d78cc53b857156fc4aed Mon Sep 17 00:00:00 2001 From: Wenju He Date: Thu, 1 Aug 2024 02:53:44 -0700 Subject: [PATCH] [sycl-post-link] Fix spec constant pattern match for DeviceSanitizer (#14740) Adjust spec constant pattern match for base alloca + offset case in device sanitizer. Address sanitizer merges static allocas into a large layout base alloca and original alloca is replaced with base + offset. --- llvm/include/llvm/SYCLLowerIR/SpecConstants.h | 4 +-- llvm/lib/Passes/PassBuilder.cpp | 1 + llvm/lib/Passes/PassRegistry.def | 1 + llvm/lib/SYCLLowerIR/SpecConstants.cpp | 17 ++++++---- .../literal-address-alloca-asan.ll | 33 +++++++++++++++++++ 5 files changed, 48 insertions(+), 8 deletions(-) create mode 100644 llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll diff --git a/llvm/include/llvm/SYCLLowerIR/SpecConstants.h b/llvm/include/llvm/SYCLLowerIR/SpecConstants.h index bbd0213158d46..8bf8bdf894d07 100644 --- a/llvm/include/llvm/SYCLLowerIR/SpecConstants.h +++ b/llvm/include/llvm/SYCLLowerIR/SpecConstants.h @@ -59,7 +59,7 @@ class SpecConstantsPass : public PassInfoMixin { enum class HandlingMode { default_values, emulation, native }; public: - SpecConstantsPass(HandlingMode Mode) : Mode(Mode) {} + SpecConstantsPass(HandlingMode Mode = HandlingMode::emulation) : Mode(Mode) {} PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM); // Searches given module for occurrences of specialization constant-specific @@ -73,7 +73,7 @@ class SpecConstantsPass : public PassInfoMixin { std::vector &DefaultValues); private: - HandlingMode Mode = HandlingMode::emulation; + HandlingMode Mode; }; bool checkModuleContainsSpecConsts(const Module &M); diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index 5bbaea52085e3..4c09bd60a8e65 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/llvm/lib/Passes/PassBuilder.cpp @@ -131,6 +131,7 @@ #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h" #include "llvm/SYCLLowerIR/SYCLPropagateJointMatrixUsage.h" #include "llvm/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.h" +#include "llvm/SYCLLowerIR/SpecConstants.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index f306e77b43afe..9f4297d0522da 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/llvm/lib/Passes/PassRegistry.def @@ -164,6 +164,7 @@ MODULE_PASS("lower-slm-reservation-calls", ESIMDLowerSLMReservationCalls()) MODULE_PASS("record-sycl-aspect-names", RecordSYCLAspectNamesPass()) MODULE_PASS("sycl-virtual-functions-analysis", SYCLVirtualFunctionsAnalysisPass()) +MODULE_PASS("spec-constants", SpecConstantsPass()) #undef MODULE_PASS #ifndef MODULE_PASS_WITH_PARAMS diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp index 58f5a0d54b26e..4f43a22e95fd9 100644 --- a/llvm/lib/SYCLLowerIR/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -20,6 +20,7 @@ #include "llvm/IR/Instructions.h" #include "llvm/IR/IntrinsicInst.h" #include "llvm/IR/Operator.h" +#include "llvm/IR/PatternMatch.h" #include "llvm/TargetParser/Triple.h" #include @@ -101,12 +102,16 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, // so that %1 is trivially known to be the address of the @.str literal. Value *TmpPtr = L->getPointerOperand(); - AssertRelease((isa(TmpPtr) && - isa(cast(TmpPtr) - ->getPointerOperand() - ->stripPointerCasts())) || - isa(TmpPtr), - "unexpected instruction type"); + auto ValueIsAlloca = [](Value *V) { + if (auto *ASC = dyn_cast(V)) + V = ASC->getPointerOperand()->stripPointerCasts(); + using namespace PatternMatch; + Value *X; + if (match(V, m_IntToPtr(m_Add(m_PtrToInt(m_Value(X)), m_ConstantInt())))) + V = X; + return isa(V); + }; + AssertRelease(ValueIsAlloca(TmpPtr), "unexpected instruction type"); // find the store of the literal address into TmpPtr StoreInst *Store = nullptr; diff --git a/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll b/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll new file mode 100644 index 0000000000000..1b904abfa0f3a --- /dev/null +++ b/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll @@ -0,0 +1,33 @@ +; RUN: opt -passes=spec-constants %s -S -o - | FileCheck %s + +; Check there is no assert error when literal address is loaded from an alloca +; with offset. + +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" + +%"class.sycl::_V1::specialization_id" = type { i32 } + +@_ZL9test_id_1 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i32 42 } +@__usid_str = constant [36 x i8] c"uide7faddc6b4d2fe92____ZL9test_id_1\00" + +define spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_14kernel_handlerEE_clES4_(ptr addrspace(4) %this1.i7) { +entry: + %MyAlloca = alloca i8, i64 224, align 32 + %0 = ptrtoint ptr %MyAlloca to i64 + %1 = add i64 %0, 96 + %2 = inttoptr i64 %1 to ptr + %SymbolicID.ascast.i = addrspacecast ptr %2 to ptr addrspace(4) + store ptr addrspace(4) addrspacecast (ptr @__usid_str to ptr addrspace(4)), ptr addrspace(4) %SymbolicID.ascast.i, align 8 + %3 = load ptr addrspace(4), ptr addrspace(4) %SymbolicID.ascast.i, align 8 + %4 = load ptr addrspace(4), ptr addrspace(4) %this1.i7, align 8 + +; CHECK-NOT: call spir_func noundef i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_( +; CHECK: %conv = sitofp i32 %load to double + + %call.i8 = call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(ptr addrspace(4) %3, ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL9test_id_1 to ptr addrspace(4)), ptr addrspace(4) %4) + %conv = sitofp i32 %call.i8 to double + ret void +} + +declare spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4))