================ @@ -4403,6 +4404,60 @@ 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. + +* ``unsigned`` address space mask +* ``unsigned`` atomic ordering, e.g. ``__ATOMIC_ACQUIRE`` +* ``const char *`` synchronization scope, e.g. ``workgroup`` + +The address space mask is a bitmask and each bit corresponds +to an OpenCL memory type: + +* ``0x1`` (first bit) is for local memory. +* ``0x2`` (second bit) is for global memory. +* ``0x4`` (third bit) is for image memory. ---------------- arsenm wrote:
Should either get named enums or string-like treatment like the sync scope 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
