[polly] r304074 - Allow side-effect free function calls in valid affine SCEVs
Tobias Grosser via llvm-commits
llvm-commits at lists.llvm.org
Sat May 27 08:18:46 PDT 2017
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"},
+};
+
+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;
+}
+
/// 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"}
More information about the llvm-commits
mailing list