Page MenuHomePhabricator

Update document about convergent attribute.

Authored by jlebar on Feb 8 2016, 6:13 PM.



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.

Diff Detail

Event Timeline

jlebar updated this revision to Diff 47281.Feb 8 2016, 6:13 PM
jlebar retitled this revision from to Update document about convergent attribute..
jlebar updated this object.
jlebar added reviewers: hfinkel, resistor, chandlerc, arsenm.
jlebar added subscribers: llvm-commits, jingyue, mehdi_amini and 2 others.

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

resistor accepted this revision.Feb 8 2016, 9:04 PM
resistor edited edge metadata.


This revision is now accepted and ready to land.Feb 8 2016, 9:04 PM
jlebar updated this revision to Diff 47355.Feb 9 2016, 1:17 PM
jlebar edited edge metadata.

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 updated this revision to Diff 47356.Feb 9 2016, 1:18 PM

Fix typo.

mehdi_amini added inline comments.Feb 9 2016, 1:20 PM

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


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 added inline comments.Feb 9 2016, 1:30 PM

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() {

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.


If we have

void sync() { __syncthreads(); }

void bar() {

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

jingyue accepted this revision.Feb 9 2016, 1:34 PM
jingyue added a reviewer: jingyue.

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


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


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.

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 updated this revision to Diff 47364.Feb 9 2016, 2:05 PM
jlebar edited edge metadata.

Reworked per joker.eph's suggestion.

jlebar updated this revision to Diff 47368.Feb 9 2016, 2:18 PM

Update again after IRC conversation with joker.eph.

mehdi_amini accepted this revision.Feb 9 2016, 2:20 PM
mehdi_amini added a reviewer: mehdi_amini.


(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 marked an inline comment as done.Feb 9 2016, 2:20 PM
jlebar added inline comments.

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

This revision was automatically updated to reflect the committed changes.
jlebar marked an inline comment as done.