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

[Sameer Sahasrabuddhe via cfe-commits]

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