[PATCH] D17012: Update document about convergent attribute.

[Justin Lebar via llvm-commits]

jlebar added inline comments.

================
Comment at: docs/LangRef.rst:1248
@@ +1247,3 @@
+    additional values.  We call these operations (e.g. the
+    ``llvm.cuda.syncthreads`` intrinsic) ``intrinsically convergent``.
+
----------------
I'm trying to get across that there exist functions that we will never remove `convergent` from, because

 a) they actually generate convergent behavior, and
 b) the only reason you'd call one of these functions is if you want this behavior.

These "intrinsically convergent" functions are different than e.g. foo in

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

foo() generates convergent behavior, but it may or may not be the case that the only reason you'd call foo() is if you want this behavior.  This patch says, we'll preserve the convergent behavior of foo if foo is marked as convergent and it may transitively invoke an intrinsically convergent function.  If the latter isn't true, there's no convergent behavior to preserve, so we can ignore the attribute.

It sounds like this isn't clear from the language here -- if you can help me with where you stumbled, I can try to rephrase.

================
Comment at: docs/LangRef.rst:1251
@@ +1250,3 @@
+    The ``convergent`` attribute indicates that the caller may rely on
+    convergent behavior within the callee.  Unless it can be proved that the
+    callee does not transitively invoke an intrinsically convergent operation,
----------------
If we have

  void sync() { __syncthreads(); }

  void bar() {
    do_stuff1();
    sync();
    do_stuff2();
  }

Then `bar` relies on the convergent behavior of `sync`.  I think `bar` is the caller and `sync` is the callee?


http://reviews.llvm.org/D17012