Skip to content

[sycl-post-link] Fix spec constant pattern match for DeviceSanitizer #14740

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Aug 1, 2024
Merged

[sycl-post-link] Fix spec constant pattern match for DeviceSanitizer #14740

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions llvm/include/llvm/SYCLLowerIR/SpecConstants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ class SpecConstantsPass : public PassInfoMixin<SpecConstantsPass> {
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
Expand All @@ -73,7 +73,7 @@ class SpecConstantsPass : public PassInfoMixin<SpecConstantsPass> {
std::vector<char> &DefaultValues);

private:
HandlingMode Mode = HandlingMode::emulation;
HandlingMode Mode;
};

bool checkModuleContainsSpecConsts(const Module &M);
Expand Down
1 change: 1 addition & 0 deletions llvm/lib/Passes/PassBuilder.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
1 change: 1 addition & 0 deletions llvm/lib/Passes/PassRegistry.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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())
Copy link
Contributor

@AlexeySachkov AlexeySachkov Aug 1, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks! Most of pass tests use incorrect approach with testing it through sycl-post-link, this is a way better direction

#undef MODULE_PASS

#ifndef MODULE_PASS_WITH_PARAMS
Expand Down
17 changes: 11 additions & 6 deletions llvm/lib/SYCLLowerIR/SpecConstants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <vector>
Expand Down Expand Up @@ -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<AddrSpaceCastInst>(TmpPtr) &&
isa<AllocaInst>(cast<AddrSpaceCastInst>(TmpPtr)
->getPointerOperand()
->stripPointerCasts())) ||
isa<AllocaInst>(TmpPtr),
"unexpected instruction type");
auto ValueIsAlloca = [](Value *V) {
if (auto *ASC = dyn_cast<AddrSpaceCastInst>(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<AllocaInst>(V);
};
AssertRelease(ValueIsAlloca(TmpPtr), "unexpected instruction type");

// find the store of the literal address into TmpPtr
StoreInst *Store = nullptr;
Expand Down
33 changes: 33 additions & 0 deletions llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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))