[polly] r304074 - Allow side-effect free function calls in valid affine SCEVs
Michael Kruse via llvm-commits
llvm-commits at lists.llvm.org
Mon May 29 15:16:12 PDT 2017
2017-05-27 17:18 GMT+02:00 Tobias Grosser via llvm-commits
<llvm-commits at lists.llvm.org>:
> Author: grosser
> Date: Sat May 27 10:18:46 2017
> New Revision: 304074
>
> URL: http://llvm.org/viewvc/llvm-project?rev=304074&view=rev
> Log:
> Allow side-effect free function calls in valid affine SCEVs
>
> Side-effect free function calls with only constant parameters can be easily
> re-generated and consequently do not prevent us from modeling a SCEV. This
> change allows array subscripts to reference function calls such as
> 'get_global_id()' as used in OpenCL.
>
> We use the function name plus the constant operands to name the parameter. This
> is possible as the function name is required and is not dropped in release
> builds the same way names of llvm::Values are dropped. We also provide more
> readable names for common OpenCL functions, to make it easy to understand the
> polyhedral model we generate.
>
> Added:
> polly/trunk/test/ScopInfo/constant_functions_as_unknowns.ll
> Modified:
> polly/trunk/include/polly/Support/SCEVValidator.h
> polly/trunk/lib/Analysis/ScopInfo.cpp
> polly/trunk/lib/Support/SCEVValidator.cpp
>
> Modified: polly/trunk/include/polly/Support/SCEVValidator.h
> URL: http://llvm.org/viewvc/llvm-project/polly/trunk/include/polly/Support/SCEVValidator.h?rev=304074&r1=304073&r2=304074&view=diff
> ==============================================================================
> --- polly/trunk/include/polly/Support/SCEVValidator.h (original)
> +++ polly/trunk/include/polly/Support/SCEVValidator.h Sat May 27 10:18:46 2017
> @@ -23,9 +23,19 @@ class ScalarEvolution;
> class Value;
> class Loop;
> class LoadInst;
> +class CallInst;
> } // namespace llvm
>
> namespace polly {
> +
> +/// Check if a call is side-effect free and has only constant arguments.
> +///
> +/// Such calls can be re-generated easily, so we do not need to model them
> +/// as scalar dependences.
> +///
> +/// @param Call The call to check.
> +bool isConstCall(llvm::CallInst *Call);
> +
> /// Find the loops referenced from a SCEV expression.
> ///
> /// @param Expr The SCEV expression to scan for loops.
>
> Modified: polly/trunk/lib/Analysis/ScopInfo.cpp
> URL: http://llvm.org/viewvc/llvm-project/polly/trunk/lib/Analysis/ScopInfo.cpp?rev=304074&r1=304073&r2=304074&view=diff
> ==============================================================================
> --- polly/trunk/lib/Analysis/ScopInfo.cpp (original)
> +++ polly/trunk/lib/Analysis/ScopInfo.cpp Sat May 27 10:18:46 2017
> @@ -1998,16 +1998,50 @@ const SCEV *Scop::getRepresentingInvaria
> return SCEVSensitiveParameterRewriter::rewrite(S, *SE, InvEquivClassVMap);
> }
>
> +// This table of function names is used to translate parameter names in more
> +// human-readable names. This makes it easier to interpret Polly analysis
> +// results.
> +StringMap<std::string> KnownNames = {
> + {"_Z13get_global_idj", "global_id"},
> + {"_Z12get_local_idj", "local_id"},
> + {"_Z15get_global_sizej", "global_size"},
> + {"_Z14get_local_sizej", "local_size"},
> + {"_Z12get_work_dimv", "work_dim"},
> + {"_Z17get_global_offsetj", "global_offset"},
> + {"_Z12get_group_idj", "group_id"},
> + {"_Z14get_num_groupsj", "num_groups"},
> +};
This is platform dependent (Itanium name mangling).
I also don't see what it has to do with side-effect free function calls.
> +static std::string getCallParamName(CallInst *Call) {
> + std::string Result;
> + raw_string_ostream OS(Result);
> + std::string Name = Call->getCalledFunction()->getName();
> +
> + auto Iterator = KnownNames.find(Name);
> + if (Iterator != KnownNames.end())
> + Name = "__" + KnownNames[Name];
> + OS << Name;
> + for (auto &Operand : Call->arg_operands()) {
> + ConstantInt *Op = cast<ConstantInt>(&Operand);
> + OS << "_" << Op->getValue();
> + }
> + OS.flush();
> + return Result;
> +}
> +
> void Scop::createParameterId(const SCEV *Parameter) {
> assert(Parameters.count(Parameter));
> assert(!ParameterIds.count(Parameter));
>
> std::string ParameterName = "p_" + std::to_string(getNumParams() - 1);
>
> - if (UseInstructionNames) {
> - if (const SCEVUnknown *ValueParameter = dyn_cast<SCEVUnknown>(Parameter)) {
> - Value *Val = ValueParameter->getValue();
> + if (const SCEVUnknown *ValueParameter = dyn_cast<SCEVUnknown>(Parameter)) {
> + Value *Val = ValueParameter->getValue();
> + CallInst *Call = dyn_cast<CallInst>(Val);
>
> + if (Call && isConstCall(Call)) {
> + ParameterName = getCallParamName(Call);
> + } else if (UseInstructionNames) {
> // If this parameter references a specific Value and this value has a name
> // we use this name as it is likely to be unique and more useful than just
> // a number.
>
> Modified: polly/trunk/lib/Support/SCEVValidator.cpp
> URL: http://llvm.org/viewvc/llvm-project/polly/trunk/lib/Support/SCEVValidator.cpp?rev=304074&r1=304073&r2=304074&view=diff
> ==============================================================================
> --- polly/trunk/lib/Support/SCEVValidator.cpp (original)
> +++ polly/trunk/lib/Support/SCEVValidator.cpp Sat May 27 10:18:46 2017
> @@ -117,6 +117,17 @@ raw_ostream &operator<<(raw_ostream &OS,
> return OS;
> }
>
> +bool polly::isConstCall(llvm::CallInst *Call) {
> + if (Call->mayReadOrWriteMemory())
> + return false;
> +
> + for (auto &Operand : Call->arg_operands())
> + if (!isa<ConstantInt>(&Operand))
> + return false;
> +
> + return true;
> +}
mayReadOrWriteMemory() is very different from "side-effect free" or
even "const call". mayReadOrWriteMemory() for instance allows malloc()
to be called (and I believe in principle sytem calls, e.g. WriteFile).
isSafeToSpeculativelyExecute() seems to be more what you want.
> +
> /// Check if a SCEV is valid in a SCoP.
> struct SCEVValidator
> : public SCEVVisitor<SCEVValidator, class ValidatorResult> {
> @@ -306,6 +317,17 @@ public:
> return ValidatorResult(SCEVType::PARAM, S);
> }
>
> + ValidatorResult visitCallInstruction(Instruction *I, const SCEV *S) {
> + assert(I->getOpcode() == Instruction::Call && "Call instruction expected");
> +
> + auto Call = cast<CallInst>(I);
> +
> + if (!isConstCall(Call))
> + return ValidatorResult(SCEVType::INVALID, S);
> +
> + return ValidatorResult(SCEVType::PARAM, S);
> + }
> +
> ValidatorResult visitLoadInstruction(Instruction *I, const SCEV *S) {
> if (R->contains(I) && ILS) {
> ILS->insert(cast<LoadInst>(I));
> @@ -396,6 +418,8 @@ public:
> return visitSDivInstruction(I, Expr);
> case Instruction::SRem:
> return visitSRemInstruction(I, Expr);
> + case Instruction::Call:
> + return visitCallInstruction(I, Expr);
> default:
> return visitGenericInst(I, Expr);
> }
> @@ -420,6 +444,11 @@ public:
> if (auto Unknown = dyn_cast<SCEVUnknown>(S)) {
> Instruction *Inst = dyn_cast<Instruction>(Unknown->getValue());
>
> + CallInst *Call = dyn_cast<CallInst>(Unknown->getValue());
> +
> + if (Call && isConstCall(Call))
> + return false;
> +
> // Return true when Inst is defined inside the region R.
> if (!Inst || !R->contains(Inst))
> return true;
>
> Added: polly/trunk/test/ScopInfo/constant_functions_as_unknowns.ll
> URL: http://llvm.org/viewvc/llvm-project/polly/trunk/test/ScopInfo/constant_functions_as_unknowns.ll?rev=304074&view=auto
> ==============================================================================
> --- polly/trunk/test/ScopInfo/constant_functions_as_unknowns.ll (added)
> +++ polly/trunk/test/ScopInfo/constant_functions_as_unknowns.ll Sat May 27 10:18:46 2017
> @@ -0,0 +1,67 @@
> +; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
> +
> +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
> +target triple = "x86_64-unknown-linux-gnu"
> +
> +; CHECK: Context:
> +; CHECK-NEXT: [__global_id_0] -> { : -9223372036854775808 <= __global_id_0 <= 9223372036854775807 }
> +; CHECK-NEXT: Assumed Context:
> +; CHECK-NEXT: [__global_id_0] -> { : }
> +; CHECK-NEXT: Invalid Context:
> +; CHECK-NEXT: [__global_id_0] -> { : 1 = 0 }
> +; CHECK-NEXT: p0: %__global_id_0
> +; CHECK-NEXT: Arrays {
> +; CHECK-NEXT: i64 MemRef_A[*]; // Element size 8
> +; CHECK-NEXT: }
> +; CHECK-NEXT: Arrays (Bounds as pw_affs) {
> +; CHECK-NEXT: i64 MemRef_A[*]; // Element size 8
> +; CHECK-NEXT: }
> +; CHECK-NEXT: Alias Groups (0):
> +; CHECK-NEXT: n/a
> +; CHECK-NEXT: Statements {
> +; CHECK-NEXT: Stmt_bb
> +; CHECK-NEXT: Domain :=
> +; CHECK-NEXT: [__global_id_0] -> { Stmt_bb[] };
> +; CHECK-NEXT: Schedule :=
> +; CHECK-NEXT: [__global_id_0] -> { Stmt_bb[] -> [] };
> +; CHECK-NEXT: MustWriteAccess := [Reduction Type: NONE] [Scalar: 0]
> +; CHECK-NEXT: [__global_id_0] -> { Stmt_bb[] -> MemRef_A[__global_id_0] };
> +; CHECK-NEXT: }
> +
> +define void @globalid(i64* nocapture %A) local_unnamed_addr #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 {
> +entry:
> + br label %next
> +
> +next:
> + br i1 true, label %bb, label %exit
> +
> +bb:
> + %__global_id_0 = tail call i64 @_Z13get_global_idj(i32 0) #2
> + %arrayidx = getelementptr inbounds i64, i64* %A, i64 %__global_id_0
> + store i64 0, i64* %arrayidx, align 8, !tbaa !6
> + br label %exit
> +
> +exit:
> + ret void
> +}
> +
> +; Function Attrs: nounwind readnone
> +declare i64 @_Z13get_global_idj(i32) local_unnamed_addr #1
> +
> +attributes #0 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-__global_id_0s"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
> +attributes #1 = { nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-__global_id_0s"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
> +attributes #2 = { nounwind readnone }
> +
> +!llvm.module.flags = !{!0}
> +!llvm.ident = !{!1}
> +
> +!0 = !{i32 1, !"wchar_size", i32 4}
> +!1 = !{!"clang version 5.0.0 (trunk 303846) (llvm/trunk 303834)"}
> +!2 = !{i32 1}
> +!3 = !{!"none"}
> +!4 = !{!"long*"}
> +!5 = !{!""}
> +!6 = !{!7, !7, i64 0}
> +!7 = !{!"long", !8, i64 0}
> +!8 = !{!"omnipotent char", !9, i64 0}
> +!9 = !{!"Simple C/C++ TBAA"}
>
>
> _______________________________________________
> llvm-commits mailing list
> llvm-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits
More information about the llvm-commits
mailing list