Author: Sameer Sahasrabuddhe Date: 2025-03-25T10:44:08+05:30 New Revision: f6a3cd54bc063efe522cc3df82fcadc86ac5a354
URL: https://github.com/llvm/llvm-project/commit/f6a3cd54bc063efe522cc3df82fcadc86ac5a354 DIFF: https://github.com/llvm/llvm-project/commit/f6a3cd54bc063efe522cc3df82fcadc86ac5a354.diff LOG: [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``. Added: Modified: clang/include/clang/Basic/AttrDocs.td Removed: ################################################################################ diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 34e7ff9612859..a52ece467ec70 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1703,17 +1703,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 @@ -1722,7 +1721,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
