.. MIT License Copyright (c) 2024-2026 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. .. meta:: :description: API reference for cuda::barrier, a multi-phase thread coordination mechanism with thread scope support, transaction counting, and extensions for shared memory barriers in libhipcxx. :keywords: libhipcxx, ROCm, HIP, C++, barrier, thread scope, shared memory, arrive, transaction count, mbarrier .. _libcudacxx-extended-api-synchronization-barrier: ``cuda::barrier`` ================= This page documents ``cuda::barrier``, a multi-phase thread coordination mechanism with thread scope support, transaction counting, and extensions for shared memory barriers. .. toctree:: :hidden: :maxdepth: 1 barrier/init barrier/barrier_native_handle barrier/barrier_arrive_tx barrier/barrier_expect_tx Defined in header ````: .. code:: cuda template class cuda::barrier; The class template ``cuda::barrier`` is an extended form of `cuda::std::barrier `_ that takes an additional :ref:`cuda::thread_scope ` argument. If ``!(scope == thread_block_scope && __isShared(this))``, then the semantics are the same as `cuda::std::barrier `_, otherwise, see below. The ``cuda::barrier`` class template extends ``cuda::std::barrier`` with the following additional operations: .. list-table:: :widths: 25 75 :header-rows: 0 * - :ref:`cuda::barrier::init ` - Initialize a ``cuda::barrier``. * - :ref:`cuda::device::barrier_native_handle ` - Get the native handle to a ``cuda::barrier``. * - :ref:`cuda::device::barrier_arrive_tx ` - Arrive on a ``cuda::barrier`` with transaction count update. * - :ref:`cuda::device::barrier_expect_tx ` - Update transaction count of ``cuda::barrier``. If ``scope == thread_scope_block && __isShared(this)``, then the semantics of `[thread.barrier.class] `_ of ISO/IEC IS 14882 (the C++ Standard) are modified as follows: .. pull-quote:: A barrier is a thread coordination mechanism whose lifetime consists of a sequence of barrier phases, where each phase allows at most an expected number of threads to block until the expected number of threads **and the expected number of transaction-based asynchronous operations** arrive at the barrier. Each *barrier phase* consists of the following steps: 1. The *expected count* is decremented by each call to ``arrive``,\ ``arrive_and_drop``\ **, or cuda::device::barrier_arrive_tx**. 2. **The transaction count is incremented by each call to cuda::device::barrier_arrive_tx and decremented by the completion of transaction-based asynchronous operations such as cuda::memcpy_async_tx.** 3. Exactly once after **both** the *expected count* **and the transaction count** reach zero, a thread executes the *completion step* during its call to ``arrive``, ``arrive_and_drop``, ``cuda::device::barrier_arrive_tx``, or ``wait``, except that it is implementation-defined whether the step executes if no thread calls ``wait``. 4. When the completion step finishes, the *expected count* is reset to what was specified by the ``expected`` argument to the constructor, possibly adjusted by calls to ``arrive_and_drop``, and the next phase starts. Concurrent invocations of the member functions of barrier **and the non-member barrier APIs in cuda::device**, other than its destructor, do not introduce data races. The member functions ``arrive`` and ``arrive_and_drop``, **and the non-member function cuda::device::barrier_arrive_tx**, execute atomically. .. rubric:: NVCC ``__shared__`` Initialization Warnings When using libhipcxx with NVCC, a ``__shared__`` ``cuda::barrier`` will lead to the following warning because ``__shared__`` variables are not initialized: .. code:: bash warning: dynamic initialization is not supported for a function-scope static __shared__ variable within a __device__/__global__ function It can be silenced using ``#pragma nv_diag_suppress static_var_with_dynamic_init``. To properly initialize a ``__shared__`` ``cuda::barrier``, use the :ref:`cuda::barrier::init ` friend function. .. rubric:: Concurrency Restrictions An object of type ``cuda::barrier`` or ``cuda::std::barrier`` shall not be accessed concurrently by CPU and GPU threads unless: - it is in unified memory and the `concurrentManagedAccess property `_ is 1, or - it is in CPU memory and the `hostNativeAtomicSupported property `_ is 1. Note: for objects of scopes other than ``cuda::thread_scope_system``, this is a data race, and is therefore also prohibited regardless of memory characteristics. Under CUDA Compute Capability 8 (Ampere) or above, when an object of type ``cuda::barrier`` is placed in ``__shared__`` memory, the member function ``arrive`` performs a reduction of the arrival count among `coalesced threads `_ followed by the arrival operation in one thread. Programs shall ensure that this transformation would not introduce errors, for example relative to the requirements of `thread.barrier.class paragraph 12 `_ of ISO/IEC IS 14882 (the C++ Standard). Under CUDA Compute Capability 6 (Pascal) or prior, an object of type ``cuda::barrier`` or ``cuda::std::barrier`` may not be used. .. rubric:: Shared memory barriers with transaction count In addition to the arrival count, a ``cuda::barrier`` object located in shared memory supports a `tx-count `_, which is used for tracking the completion of some asynchronous memory operations or transactions. The tx-count tracks the number of asynchronous transactions, in units specified by the asynchronous memory operation (typically bytes) that are outstanding and have not yet completed. This capability is exposed, starting with the Hopper architecture (CUDA Compute Capability 9). The tx-count of ``cuda::barrier`` must be set to the total amount of asynchronous memory operations, in units as specified by the asynchronous operations, to be tracked by the current phase. This can be achieved with the :ref:`cuda::device::barrier_arrive_tx ` function call. Upon completion of each of the asynchronous operations, the tx-count of the ``cuda::barrier`` will be updated and thus progress the ``cuda::barrier`` towards the completion of the current phase. This may complete the current phase. .. rubric:: Implementation-Defined Behavior For each :ref:`cuda::thread_scope ` ``S`` and ``CompletionFunction`` ``F``, the value of ``cuda::barrier::max()`` is as follows: .. list-table:: :widths: 25 25 50 :header-rows: 0 * - :ref:`cuda::thread_scope ` ``S`` - ``CompletionFunction`` ``F`` - ``barrier::max()`` * - ``cuda::thread_scope_block`` - Default or user-provided - ``(1 << 20) - 1`` * - *Not* ``cuda::thread_scope_block`` - Default - ``cuda::std::numeric_limits::max()`` * - *Not* ``cuda::thread_scope_block`` - User-provided - ``cuda::std::numeric_limits::max()`` .. rubric:: Example .. code:: cuda #include __global__ void example_kernel() { // This barrier is suitable for all threads in the system. cuda::barrier a(10); // This barrier has the same type as the previous one (`a`). cuda::std::barrier<> b(10); // This barrier is suitable for all threads on the current processor (e.g. GPU). cuda::barrier c(10); // This barrier is suitable for all threads in the same thread block. cuda::barrier d(10); } `See it on Godbolt `_