[clang] [llvm] [AMDGPU] Add OpenCL-specific fence address space masks (PR #78572)

[Matt Arsenault via llvm-commits]

================
@@ -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