[polly] r304074 - Allow side-effect free function calls in valid affine SCEVs

[Tobias Grosser via llvm-commits]

On Tue, May 30, 2017, at 12:16 AM, Michael Kruse via llvm-commits wrote:
> 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).

Right. It is mostly for me (and others) to see more readable output, but
it does not carry semantics. So it does not hurt a lot in case it won't
apply on other platforms. Also, are there platforms for OpenCL which do
not use Itanium name mangeling?
 
> I also don't see what it has to do with side-effect free function calls.

OK. I should not have sneaked this in as part of this commit. Will be
more careful in the future.
 
> > +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.

This is a very good catch and I clearly should have pushed this patch
for reviews first. Sorry. The problem is that these functions are
annotated as readnone, but not as speculatable. The clang opencl headers
use "pure" to annotate these functions. I will post a message to cfe-dev
to understand if this implies speculatable.

Best,
Tobias