Add convergent-removing bits to FunctionAttrs pass.

[jlebar]

Summary

Remove the convergent attribute on any functions which provably do not
contain or invoke any convergent functions.

After this change, we'll be able to modify clang to conservatively add
'convergent' to all functions when compiling CUDA.

[mehdi_amini]

lib/Transforms/IPO/FunctionAttrs.cpp

950 ↗

(On Diff #47282)

could also by a call to any_of?

954 ↗

(On Diff #47282)

Callee->isIntrinsic() should be redundant here (I mean !Callee->isIntrinsic() && Callee->isConvergent() should be impossible after the llvm::any_of loop above right?

[jingyue]

lib/Transforms/IPO/FunctionAttrs.cpp

934 ↗

(On Diff #47282)

Please comment on what it returns.

936 ↗

(On Diff #47282)

SCC may contain multiple functions. In that case, you need to consider all of them.

949 ↗

(On Diff #47282)

You can iterate through the call sites with CallGraphNode::iterator (http://llvm.org/docs/doxygen/html/CallGraph_8h_source.html#l00186).

[mehdi_amini]

lib/Transforms/IPO/FunctionAttrs.cpp

936 ↗

(On Diff #47282)

That's a good point: any reason not to take the SCCNodeSet and iterate over like addNonNullAttrs for instance?

[jlebar]

Address jingyue, joker.eph's comments.

[jlebar]

Thank you for the reviews. I've updated the patch, please have a look when you have a chance.

lib/Transforms/IPO/FunctionAttrs.cpp

936 ↗

(On Diff #47282)

SCC may contain multiple functions. In that case, you need to consider all of them.

Thanks, this was quite broken.

any reason not to take the SCCNodeSet?

SCCNodeSet is the set of strongly-connected nodes -- aiui a set of mutually-recursive functions, excluding any OptimizeNone or external functions that may be present in the set.

Two problems with using this rather than SCC itself. First, I think we do want to consider OptimizeNone functions. At least, if Foo and Bar are convergent and mutually recursive and Bar is OptimizeNone, then Foo should remain convergent. (The original code didn't handle OptimizeNone correctly either; I've fixed this.)

Second, if we use this, we don't get the actual control-flow graph, which is contained in SCC. There appears to be nontrivial logic to e.g. devirtualize calls in there, so I don't think we want to replicate all that code when we iterate over the function looking for calls to convergent intrinsics.

949 ↗

(On Diff #47282)

I may be misunderstanding, but I think that's what we're doing in the loop above when we do llvm::any_of(*CGN) -- we're iterating over CallGraphNode::begin()/end(). The issue is that, as the CallGraphNode is constructed in CGPassManager::RefreshCallGraph, does not contain calls to intrinsic functions.

954 ↗

(On Diff #47282)

You're right, it wasn't necessary. However, it now is, because now we're ignoring calls to functions in the same SCC.

[jingyue]

lib/Transforms/IPO/FunctionAttrs.cpp

949 ↗

(On Diff #47352)

Looks good now.

[mehdi_amini]

LGTM.

lib/Transforms/IPO/FunctionAttrs.cpp

973 ↗

(On Diff #47352)

Suggestion: outline the loop body in a helper function canRemoveConvergent(CallGraphNode *CGN)

The loop becomes then: if(!llvm::all_of(SCC, canRemoveConvergent)) return false;

[jlebar]

lib/Transforms/IPO/FunctionAttrs.cpp

973 ↗

(On Diff #47352)

I'm actually not wild about this because it's a bit misleading. Consider

void foo() { bar(); }
void bar() { __syncthreads(); foo(); }

canRemoveConvergent(foo) would return true, even though, because we cannot remove convergent from bar and foo and bar are in the same SCC, we cannot in fact remove convergent from foo.

[mehdi_amini]

lib/Transforms/IPO/FunctionAttrs.cpp

973 ↗

(On Diff #47352)

To be clear: I was suggesting a pure "syntactic" change no functional change intended. I agree that the naming I proposed isn't well suited though (considering your example)

llvm/trunk/include/llvm/IR/Function.h

public:
   /// @brief Determine if the call is convergent.
   bool isConvergent() const {
     return AttributeSets.hasAttribute(AttributeSet::FunctionIndex,
                                       Attribute::Convergent);
   }
   void setConvergent() {
     addFnAttr(Attribute::Convergent);
   }
+  void setNotConvergent() {
+    removeFnAttr(Attribute::Convergent);
+  }
 
   /// Determine if the function is known not to recurse, directly or
   /// indirectly.
   bool doesNotRecurse() const {
     return AttributeSets.hasAttribute(AttributeSet::FunctionIndex,
                                       Attribute::NoRecurse);
   }
   void setDoesNotRecurse() {

llvm/trunk/lib/Transforms/IPO/FunctionAttrs.cpp

for (Function *F : SCCNodes) {
       ++NumNonNullReturn;
       MadeChange = true;
     }
   }
 
   return MadeChange;
 }
 
+/// Removes convergent attributes where we can prove that none of the SCC's
+/// callees are themselves convergent.  Returns true if successful at removing
+/// the attribute.
+static bool removeConvergentAttrs(const CallGraphSCC &SCC,
+                                  const SCCNodeSet &SCCNodes) {
+  // Determines whether a function can be made non-convergent, ignoring all
+  // other functions in SCC.  (A function can *actually* be made non-convergent
+  // only if all functions in its SCC can be made convergent.)
+  auto CanRemoveConvergent = [&] (CallGraphNode *CGN) {
+    Function *F = CGN->getFunction();
+    if (!F) return false;
+
+    if (!F->isConvergent()) return true;
+
+    // Can't remove convergent from declarations.
+    if (F->isDeclaration()) return false;
+
+    // Don't remove convergent from optnone functions.
+    if (F->hasFnAttribute(Attribute::OptimizeNone))
+      return false;
+
+    // Can't remove convergent if any of F's callees -- ignoring functions in the
+    // SCC itself -- are convergent.
+    if (llvm::any_of(*CGN, [&](const CallGraphNode::CallRecord &CR) {
+          Function *F = CR.second->getFunction();
+          return SCCNodes.count(F) == 0 && (!F || F->isConvergent());
+        }))
+      return false;
+
+    // CGN doesn't contain calls to intrinsics, so iterate over all of F's
+    // callsites, looking for any calls to convergent intrinsics.  If we find one,
+    // F must remain marked as convergent.
+    auto IsConvergentIntrinsicCall = [](Instruction &I) {
+      CallSite CS(cast<Value>(&I));
+      if (!CS)
+        return false;
+      Function *Callee = CS.getCalledFunction();
+      return Callee && Callee->isIntrinsic() && Callee->isConvergent();
+    };
+    return !llvm::any_of(*F, [=](BasicBlock &BB) {
+      return llvm::any_of(BB, IsConvergentIntrinsicCall);
+    });
+  };
+
+  // We can remove the convergent attr from functions in the SCC if they all can
+  // be made non-convergent (because they call only non-convergent functions,
+  // other than each other).
+  if (!llvm::all_of(SCC, CanRemoveConvergent)) return false;
+
+  // If we got here, all of the SCC's callees are non-convergent, and none of
+  // the optnone functions in the SCC are marked as convergent.  Therefore all
+  // of the SCC's functions can be marked as non-convergent.
+  for (CallGraphNode *CGN : SCC)
+    if (Function *F = CGN->getFunction()) {
+      if (F->isConvergent())
+        DEBUG(dbgs() << "Removing convergent attr from " << F->getName()
+                     << "\n");
+      F->setNotConvergent();
+    }
+  return true;
+}
+
 static bool setDoesNotRecurse(Function &F) {
   if (F.doesNotRecurse())
     return false;
   F.setDoesNotRecurse();
   ++NumNoRecurse;
   return true;
 }
 
bool PostOrderFunctionAttrs::runOnSCC(CallGraphSCC &SCC) {
   Changed |= addReadAttrs(SCCNodes, AARGetter);
   Changed |= addArgumentAttrs(SCCNodes);
 
   // If we have no external nodes participating in the SCC, we can deduce some
   // more precise attributes as well.
   if (!ExternalNode) {
     Changed |= addNoAliasAttrs(SCCNodes);
     Changed |= addNonNullAttrs(SCCNodes, *TLI);
+    Changed |= removeConvergentAttrs(SCC, SCCNodes);
   }
 
   Changed |= addNoRecurseAttrs(SCC);
   return Changed;
 }
 
 namespace {
 /// A pass to do RPO deduction and propagation of function attributes.

llvm/trunk/test/Transforms/FunctionAttrs/convergent.ll

+; RUN: opt < %s -basicaa -functionattrs -rpo-functionattrs -S | FileCheck %s
+
+; CHECK: Function Attrs
+; CHECK-NOT: convergent
+; CHECK-NEXT: define i32 @nonleaf()
+define i32 @nonleaf() convergent {
+  %a = call i32 @leaf()
+  ret i32 %a
+}
+
+; CHECK: Function Attrs
+; CHECK-NOT: convergent
+; CHECK-NEXT: define i32 @leaf()
+define i32 @leaf() convergent {
+  ret i32 0
+}
+
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: declare i32 @k()
+declare i32 @k() convergent
+
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: define i32 @extern()
+define i32 @extern() convergent {
+  %a = call i32 @k()
+  ret i32 %a
+}
+
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: define i32 @call_extern()
+define i32 @call_extern() convergent {
+  %a = call i32 @extern()
+  ret i32 %a
+}
+
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: declare void @llvm.cuda.syncthreads()
+declare void @llvm.cuda.syncthreads() convergent
+
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: define i32 @intrinsic()
+define i32 @intrinsic() convergent {
+  call void @llvm.cuda.syncthreads()
+  ret i32 0
+}
+
+@xyz = global i32 ()* null
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: define i32 @functionptr()
+define i32 @functionptr() convergent {
+  %1 = load i32 ()*, i32 ()** @xyz
+  %2 = call i32 %1()
+  ret i32 %2
+}
+
+; CHECK: Function Attrs
+; CHECK-NOT: convergent
+; CHECK-NEXT: define i32 @recursive1()
+define i32 @recursive1() convergent {
+  %a = call i32 @recursive2()
+  ret i32 %a
+}
+
+; CHECK: Function Attrs
+; CHECK-NOT: convergent
+; CHECK-NEXT: define i32 @recursive2()
+define i32 @recursive2() convergent {
+  %a = call i32 @recursive1()
+  ret i32 %a
+}
+
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: define i32 @noopt()
+define i32 @noopt() convergent optnone noinline {
+  %a = call i32 @noopt_friend()
+  ret i32 0
+}
+
+; A function which is mutually-recursive with a convergent, optnone function
+; shouldn't have its convergent attribute stripped.
+; CHECK: Function Attrs
+; CHECK-SAME: convergent
+; CHECK-NEXT: define i32 @noopt_friend()
+define i32 @noopt_friend() convergent {
+  %a = call i32 @noopt()
+  ret i32 0
+}