Synchronization Functions
=========================

Currently, only several of the SYCL synchronization operations are supported. For synchronization of all threads in the same thread block, a helper function called ``numba_dpex.barrier()`` is provided. This function implements the same pattern as barriers in traditional
multi-threaded programming: invoking the function forces a thread to wait until
all threads in the block reach the barrier, at which point it returns control
to all its callers.

``numba_dpex.barrier()`` supports two memory fence options:

- ``numba_dpex.CLK_GLOBAL_MEM_FENCE``: The barrier function will queue a memory
  fence to ensure correct ordering of memory operations to global memory. Using
  the option can be useful when work-items, for example, write to buffer or
  image objects and then want to read the updated data. Passing no arguments to ``numba_dpex.barrier()`` is equivalent to setting the global memory fence option. For example,

  .. literalinclude:: ../../../numba_dpex/examples/barrier.py
   :pyobject: no_arg_barrier_support

- ``numba_dpex.CLK_LOCAL_MEM_FENCE``: The barrier function will either flush
  any variables stored in local memory or queue a memory fence to ensure
  correct ordering of memory operations to local memory. For example,

.. literalinclude:: ../../../numba_dpex/examples/barrier.py
   :pyobject: local_memory


.. note::

    The ``numba_dpex.barrier()`` function is semantically equivalent to
    ``numba.cuda.syncthreads``.