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 ↗

(On Diff #47355)

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

1251 ↗

(On Diff #47355)

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 ↗

(On Diff #47356)

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 ↗

(On Diff #47356)

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 ↗

(On Diff #47356)

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

1255 ↗

(On Diff #47356)

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

1255 ↗

(On Diff #47368)

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

llvm/trunk/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.
+    In some parallel execution models, there exist operations that cannot be
+    made control-dependent on any additional values.  We call such operations
+    ``convergent``, and mark them with this function attribute.
+
+    For example, the intrinsic ``llvm.cuda.syncthreads`` is ``convergent``, so
+    calls to this intrinsic cannot be made control-dependent on additional
+    values.  Other functions may also be marked as convergent; this prevents
+    the same optimization on those functions.
+
+    The optimizer may remove the ``convergent`` attribute when it can prove
+    that the function does not execute any convergent operations.
 ``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``