https://github.com/ssahasra created
https://github.com/llvm/llvm-project/pull/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``.
>From 6a390c2cf3ec17587ded6d359f48abbbe976a5ac Mon Sep 17 00:00:00 2001
From: Sameer Sahasrabuddhe <sameer.sahasrabud...@amd.com>
Date: Mon, 24 Mar 2025 14:55:23 +0530
Subject: [PATCH] [clang] ``noconvergent`` does not affect calls to convergent
functions
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``.
---
clang/include/clang/Basic/AttrDocs.td | 20 ++++++++++----------
1 file changed, 10 insertions(+), 10 deletions(-)
diff --git a/clang/include/clang/Basic/AttrDocs.td
b/clang/include/clang/Basic/AttrDocs.td
index a8b588169725a..98d29838dce42 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -1658,17 +1658,16 @@ Sample usage:
def NoConvergentDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
-This attribute prevents a function from being treated as convergent, which
-means that optimizations can only move calls to that function to
-control-equivalent blocks. If a statement is marked as ``noconvergent`` and
-contains calls, it also prevents those calls from being treated as convergent.
-In other words, those calls are not restricted to only being moved to
-control-equivalent blocks.
+This attribute prevents a function from being treated as convergent; when a
+function is marked ``noconvergent``, calls to that function are not
+automatically assumed to be convergent, unless such calls are explicitly marked
+as ``convergent``. If a statement is marked as ``noconvergent``, any calls to
+inline ``asm`` in that statement are no longer treated as convergent.
In languages following SPMD/SIMT programming model, e.g., CUDA/HIP, function
-declarations and calls are treated as convergent by default for correctness.
-This ``noconvergent`` attribute is helpful for developers to prevent them from
-being treated as convergent when it's safe.
+declarations and inline asm calls are treated as convergent by default for
+correctness. This ``noconvergent`` attribute is helpful for developers to
+prevent them from being treated as convergent when it's safe.
.. code-block:: c
@@ -1677,7 +1676,8 @@ being treated as convergent when it's safe.
__device__ int example(void) {
float x;
- [[clang::noconvergent]] x = bar(x);
+ [[clang::noconvergent]] x = bar(x); // no effect on convergence
+ [[clang::noconvergent]] { asm volatile ("nop"); } // the asm call is
non-convergent
}
}];
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
