================ @@ -199,21 +199,58 @@ map in the following way to CUDA builtins: Barriers -------- -'``llvm.nvvm.barrier0``' -^^^^^^^^^^^^^^^^^^^^^^^^^^^ +'``llvm.nvvm.barrier.cta.*``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ Syntax: """"""" .. code-block:: llvm - declare void @llvm.nvvm.barrier0() + declare void @llvm.nvvm.barrier.cta.sync(i32 %id, i32 %n) + declare void @llvm.nvvm.barrier.cta.sync.all(i32 %id) + declare void @llvm.nvvm.barrier.cta.arrive(i32 %id, i32 %n) + + declare void @llvm.nvvm.barrier.cta.sync.aligned(i32 %id, i32 %n) + declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %id) + declare void @llvm.nvvm.barrier.cta.arrive.aligned(i32 %id, i32 %n) Overview: """"""""" -The '``@llvm.nvvm.barrier0()``' intrinsic emits a PTX ``bar.sync 0`` -instruction, equivalent to the ``__syncthreads()`` call in CUDA. +The '``@llvm.nvvm.barrier.cta.*``' family of intrinsics perform barrier +synchronization and communication within a CTA. They can be used by the threads +within the CTA for synchronization and communication. + +Semantics: +"""""""""" + +Operand %id specifies a logical barrier resource and must fall within the range +0 through 15. When present, operand %n specifies the number of threads +participating in the barrier. When specifying a thread count, the value must be +a multiple of the warp size. With the '``@llvm.nvvm.barrier.cta.sync.*``' +variants, the '``.all``' suffix indicates that all threads in the CTA should +participate in the barrier and the %n operand is not present. + +All forms of the '``@llvm.nvvm.barrier.cta.*``' intrinsic cause the executing +thread to wait for all non-exited threads from its warp and then marks the +warp's arrival at the barrier. In addition to signaling its arrival at the +barrier, the '``@llvm.nvvm.barrier.cta.sync.*``' intrinsics cause the executing +thread to wait for non-exited threads of all other warps participating in the +barrier to arrive. On the other hand, the '``@llvm.nvvm.barrier.cta.arrive.*``' +intrinsic does not cause the executing thread to wait for threads of other +participating warps. + +When a barrier completes, the waiting threads are restarted without delay, +and the barrier is reinitialized so that it can be immediately reused. + +The '``@llvm.nvvm.barrier.cta.*``' intrinsic has an optional '``.aligned``' +modifier to indicate textual alignment of the barrier. When specified, it +indicates that all threads in the CTA will execute the same +'``@llvm.nvvm.barrier.cta.*``' instruction. In conditionally executed code, an +aligned '``@llvm.nvvm.barrier.cta.*``' instruction should only be used if it is +known that all threads in the CTA evaluate the condition identically, otherwise +behavior is undefined. ---------------- AlexMaclean wrote:
I think the PTX instruction used to lower the intrinsic is an implementation detail. It should not be necessary to link to it if the semantics of the intrinsic are fully defined. https://github.com/llvm/llvm-project/pull/140615 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
