[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).
> > +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, but this attribute is not
correctly propagated. See:

https://reviews.llvm.org/D33774

Best,
Tobias