[llvm] [NVPTX] Add elect.sync Intrinsic (PR #104780)
Durgadoss R via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 19 11:15:18 PDT 2024
================
@@ -251,10 +251,34 @@ Overview:
The '``@llvm.nvvm.barrier0()``' intrinsic emits a PTX ``bar.sync 0``
instruction, equivalent to the ``__syncthreads()`` call in CUDA.
+Electing a thread
+-----------------
+
+'``llvm.nvvm.elect.sync``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+ declare {i32, i1} @llvm.nvvm.elect.sync(i32 %membermask)
+
+Overview:
+"""""""""
+
+The '``@llvm.nvvm.elect.sync``' intrinsic generates the ``elect.sync``
+PTX instruction, which elects one predicated active leader thread from
+a set of threads specified by ``membermask``. The behavior is undefined
+if the executing thread is not in ``membermask``. The laneid of the
+elected thread is captured in the i32 return value. The i1 return
+value is set to ``True`` for the leader thread and ``False`` for all
+the other threads. For more information, refer PTX ISA
----------------
durga4github wrote:
Sure, updated the docs.
https://github.com/llvm/llvm-project/pull/104780
More information about the llvm-commits
mailing list