Skip to content

Commit

Permalink
Merge pull request #141 from frasercrmck/fix-assumption-cache
Browse files Browse the repository at this point in the history
[vecz] Fix invalid dangling AssumptionCaches
  • Loading branch information
frasercrmck authored Sep 21, 2023
2 parents 5b6a50d + dc30e9f commit 93298ff
Show file tree
Hide file tree
Showing 6 changed files with 72 additions and 9 deletions.
10 changes: 6 additions & 4 deletions modules/compiler/vecz/source/analysis/stride_analysis.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -49,8 +49,9 @@ OffsetInfo &StrideAnalysisResult::analyze(Value *V) {
}

StrideAnalysisResult::StrideAnalysisResult(llvm::Function &f,
UniformValueResult &uvr)
: F(f), UVR(uvr), assumptions(F) {
UniformValueResult &uvr,
AssumptionCache &AC)
: F(f), UVR(uvr), AC(AC) {
for (auto &BB : F) {
for (auto &I : BB) {
if (!UVR.isVarying(&I)) {
Expand Down Expand Up @@ -83,6 +84,7 @@ Value *StrideAnalysisResult::buildMemoryStride(IRBuilder<> &B, llvm::Value *Ptr,

StrideAnalysisResult StrideAnalysis::run(llvm::Function &F,
llvm::FunctionAnalysisManager &AM) {
UniformValueResult &UVR = AM.getResult<UniformValueAnalysis>(F);
return Result(F, UVR);
auto &AC = AM.getResult<AssumptionAnalysis>(F);
auto &UVR = AM.getResult<UniformValueAnalysis>(F);
return Result(F, UVR, AC);
}
Original file line number Diff line number Diff line change
Expand Up @@ -46,9 +46,10 @@ class StrideAnalysisResult {
/// @brief The Uniform Value Result to use during analysis
UniformValueResult &UVR;
/// @brief AssumptionCache for computing live bits of uniform values
llvm::AssumptionCache assumptions;
llvm::AssumptionCache &AC;

StrideAnalysisResult(llvm::Function &f, UniformValueResult &uvr);
StrideAnalysisResult(llvm::Function &f, UniformValueResult &uvr,
llvm::AssumptionCache &AC);

/// @brief generate stride `ConstantInt`s or `Instruction`s for all analyzed
/// values.
Expand Down
4 changes: 2 additions & 2 deletions modules/compiler/vecz/source/offset_info.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -176,8 +176,8 @@ OffsetInfo &OffsetInfo::analyze(Value *Offset, StrideAnalysisResult &SAR) {

// If we have a uniform value here we don't need to analyse any further.
if (!SAR.UVR.isVarying(Ins)) {
auto const &KB = computeKnownBits(Ins, SAR.F.getParent()->getDataLayout(),
0, &SAR.assumptions);
auto const &KB =
computeKnownBits(Ins, SAR.F.getParent()->getDataLayout(), 0, &SAR.AC);
auto const bitWidth = OffsetTy->getIntegerBitWidth();

// We are interested in the bits that are not known to be zero.
Expand Down
2 changes: 1 addition & 1 deletion modules/compiler/vecz/source/pass.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -166,7 +166,7 @@ PreservedAnalyses RunVeczPass::run(Module &M, ModuleAnalysisManager &MAM) {
// If we fail to vectorize a function, we still cloned and then
// deleted it which affects internal addresses. The module has changed
// and we can't cache any analyses.
Mach.getFAM().invalidate(*VectorizedFn, llvm::PreservedAnalyses::all());
Mach.getFAM().invalidate(*VectorizedFn, llvm::PreservedAnalyses::none());
// Remove the partially-vectorized function if something went wrong.
Ctx.clearActiveVU(VectorizedFn);
VU->setVectorizedFunction(nullptr);
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,44 @@
; Copyright (C) Codeplay Software Limited
;
; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
; Exceptions; you may not use this file except in compliance with the License.
; You may obtain a copy of the License at
;
; https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
;
; Unless required by applicable law or agreed to in writing, software
; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
; License for the specific language governing permissions and limitations
; under the License.
;
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

; Just check that we correctly clean up the assumption cache when vectorizing
; this function.:
; RUN: veczc -k foo -w 2 -S < %s
; RUN: not veczc -k foo -w 2 -vecz-scalable -S < %s

target triple = "spir64-unknown-unknown"
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"

define spir_kernel void @foo(ptr addrspace(1) nocapture readonly %_arg_v_acc) #0 {
entry:
%v4 = tail call i64 @__mux_get_global_id(i32 0) #2
%v5 = tail call i64 @__mux_get_global_offset(i32 0) #2
%v6 = sub i64 %v4, %v5
%v7 = icmp ult i64 %v6, 2147483648
tail call void @llvm.assume(i1 %v7)
%arrayidx.i.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_v_acc, i64 %v6
%v8 = load i32, ptr addrspace(1) %arrayidx.i.i, align 4
ret void
}

declare void @llvm.assume(i1 noundef) #1

declare i64 @__mux_get_global_id(i32) #2
declare i64 @__mux_get_global_offset(i32) #2

attributes #0 = { convergent nounwind "mux-kernel"="entry-point" "mux-orig-fn"="foo" }
attributes #1 = { mustprogress nocallback nofree nosync nounwind willreturn inaccessiblememonly }
attributes #2 = { alwaysinline norecurse nounwind readonly }
Original file line number Diff line number Diff line change
@@ -1,3 +1,19 @@
; Copyright (C) Codeplay Software Limited
;
; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
; Exceptions; you may not use this file except in compliance with the License.
; You may obtain a copy of the License at
;
; https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
;
; Unless required by applicable law or agreed to in writing, software
; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
; License for the specific language governing permissions and limitations
; under the License.
;
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

; RUN: not veczc -k noduplicate:4,8 -S < %s 2>&1 | FileCheck %s

target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
Expand Down

0 comments on commit 93298ff

Please sign in to comment.