Update document about convergent attribute.

[jlebar]

Summary

Be more explicit that 'convergent' means that you either are or may
invoke an inherently convergent operation, such as
llvm.cuda.syncthreads. Where the compiler can prove that you do not
invoke such an operation, it's free to remove the attribute and treat
you as not-convergent.

[mehdi_amini]

LGTM (but please wait for at least another pair of eyes)

[resistor]

LGTM.

[jlebar]

I reread this patch after explaining this change IRL today, and it now feels
unnecessarily vague. With appologies for the churn, here's an attempt at being
more precise.

Please let me know what you think.

[jlebar]

Fix typo.

[mehdi_amini]

docs/LangRef.rst

1248

I've never read "intrinsically convergent" before, what does "intrinsically" adds here?

1251

I don't understand why the "caller" is the one relying on convergent behavior from the callee, I think it is the opposite and the "callee" is relying on convergent behavior of the call site.

[jlebar]

docs/LangRef.rst

1248

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.

1251

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?

[jingyue]

Just a reminder if you haven't done that already, double-check how the web page looks like before you commit.

docs/LangRef.rst

1248

Italicizing a definition is more common than monospacing it, so suggest intrinsically convergent instead of intrinsically convergent.

1255

Nit: I don't know how this empty line affects the web display, but I didn't see other attributes (such as ssq) have an empty line at the end.

[mehdi_amini]

Here is a straw-man proposal since I don't get the need for distinction between intrinsics and regular functions:

In these execution models, certain operations have the property that they
cannot be made control-dependent on any additional values.  We call these
operations ``convergent``.  For example, the ``llvm.cuda.syncthreads``
intrinsic is convergent, and so functions which (transitively) invoke it
must also be convergent. Any call to a function with the convergent 
attribute is considered a ``convergent operation``.

The optimizer is safely allowed to remove the convergent attribute from a 
function if it can be proven that it'll never execute a convergent operation.

[jlebar]

Reworked per joker.eph's suggestion.

[jlebar]

Update again after IRC conversation with joker.eph.

[mehdi_amini]

LGTM.

(I think the last sentence of the second paragraph unless the optimizer can prove that they never execute a convergent operation. is redundant with the third paragraph)

[jlebar]

docs/LangRef.rst

1257

Removed the line and checked the HTML output, thank you.

docs/LangRef.rst

``builtin``
     direct calls to functions that are declared with the ``nobuiltin``
     attribute.
 ``cold``
     This attribute indicates that this function is rarely called. When
     computing edge weights, basic blocks post-dominated by a cold
     function call are also considered to be cold; and, thus, given low
     weight.
 ``convergent``
-    This attribute indicates that the callee is dependent on a convergent
-    thread execution pattern under certain parallel execution models.
-    Transformations that are execution model agnostic may not make the execution
-    of a convergent operation control dependent on any additional values.
+    This attribute indicates that the optimizer should preserve convergent
+    behavior produced by the callee.
+
+    In some parallel execution models, there exist operations that are useful
+    only if the optimizer does not make them control-dependent on any
+    additional values.  We call these operations (e.g. the
+    ``llvm.cuda.syncthreads`` intrinsic) ``intrinsically convergent``.
+
+    The ``convergent`` attribute indicates that the caller may rely on
+    convergent behavior within of the callee.  Unless it can be proved that the
+    callee does not transitively invoke an intrinsically convergent operation,
+    a convergent callee must not be made control-dependent on any additional
+    values.
+
 ``inaccessiblememonly``
     This attribute indicates that the function may only access memory that
     is not accessible by the module being compiled. This is a weaker form
     of ``readnone``.
 ``inaccessiblemem_or_argmemonly``
     This attribute indicates that the function may only access memory that is
     either not accessible by the module being compiled, or is pointed to
     by its pointer arguments. This is a weaker form of  ``argmemonly``
 ``inlinehint``
     This attribute indicates that the source code contained a hint that