Skip to content

Commit f6a3cd5

Browse files
authored
[clang] noconvergent does not affect calls to convergent functions (#132701)
When placed on a function, the ``clang::noconvergent`` attribute ensures that the function is not assumed to be convergent. But the same attribute has no effect on function calls. A call is convergent if the callee is convergent. This is based on the fact that in LLVM, a call always inherits all the attributes of the callee. Only ``convergent`` is an attribute in LLVM IR, and there is no equivalent of ``clang::noconvergent``.
1 parent 9ee950b commit f6a3cd5

File tree

1 file changed

+10
-10
lines changed

1 file changed

+10
-10
lines changed

clang/include/clang/Basic/AttrDocs.td

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1703,17 +1703,16 @@ Sample usage:
17031703
def NoConvergentDocs : Documentation {
17041704
let Category = DocCatFunction;
17051705
let Content = [{
1706-
This attribute prevents a function from being treated as convergent, which
1707-
means that optimizations can only move calls to that function to
1708-
control-equivalent blocks. If a statement is marked as ``noconvergent`` and
1709-
contains calls, it also prevents those calls from being treated as convergent.
1710-
In other words, those calls are not restricted to only being moved to
1711-
control-equivalent blocks.
1706+
This attribute prevents a function from being treated as convergent; when a
1707+
function is marked ``noconvergent``, calls to that function are not
1708+
automatically assumed to be convergent, unless such calls are explicitly marked
1709+
as ``convergent``. If a statement is marked as ``noconvergent``, any calls to
1710+
inline ``asm`` in that statement are no longer treated as convergent.
17121711

17131712
In languages following SPMD/SIMT programming model, e.g., CUDA/HIP, function
1714-
declarations and calls are treated as convergent by default for correctness.
1715-
This ``noconvergent`` attribute is helpful for developers to prevent them from
1716-
being treated as convergent when it's safe.
1713+
declarations and inline asm calls are treated as convergent by default for
1714+
correctness. This ``noconvergent`` attribute is helpful for developers to
1715+
prevent them from being treated as convergent when it's safe.
17171716

17181717
.. code-block:: c
17191718

@@ -1722,7 +1721,8 @@ being treated as convergent when it's safe.
17221721

17231722
__device__ int example(void) {
17241723
float x;
1725-
[[clang::noconvergent]] x = bar(x);
1724+
[[clang::noconvergent]] x = bar(x); // no effect on convergence
1725+
[[clang::noconvergent]] { asm volatile ("nop"); } // the asm call is non-convergent
17261726
}
17271727

17281728
}];

0 commit comments

Comments
 (0)