================ @@ -4408,6 +4409,54 @@ Target-Specific Extensions Clang supports some language features conditionally on some targets. +AMDGPU Language Extensions +-------------------------- + +__builtin_amdgcn_fence +^^^^^^^^^^^^^^^^^^^^^^ + +``__builtin_amdgcn_fence`` emits a fence for all address spaces +and takes the following arguments: + +* ``unsigned`` atomic ordering, e.g. ``__ATOMIC_ACQUIRE`` +* ``const char *`` synchronization scope, e.g. ``workgroup`` + +.. code-block:: c++ + + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup"); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent"); + +__builtin_amdgcn_masked_fence +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +``__builtin_amdgcn_masked_fence`` emits a fence for one or more address +spaces and takes the following arguments: + +* ``unsigned`` atomic ordering, e.g. ``__ATOMIC_ACQUIRE`` +* ``const char *`` synchronization scope, e.g. ``workgroup`` +* Zero or more ``const char *`` address spaces. + +The address spaces arguments must be string literals with known values, such as: + +* ``"local"`` +* ``"global"`` +* ``"image"`` + +If there are no address spaces specified, this fence behaves like ---------------- ssahasra wrote:
To answer comments elsewhere, this documentation makes it really obvious that the two fences are really just one. https://github.com/llvm/llvm-project/pull/78572 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
