[clang] [llvm] [AMDGPU] Add amdgpu-as MMRA for fences (PR #78572)

[Sameer Sahasrabuddhe via llvm-commits]

================
@@ -4408,6 +4409,42 @@ Target-Specific Extensions
 
 Clang supports some language features conditionally on some targets.
 
+AMDGPU Language Extensions
+--------------------------
+
+__builtin_amdgcn_fence
+^^^^^^^^^^^^^^^^^^^^^^
+
+``__builtin_amdgcn_fence`` emits a fence.
+
+* ``unsigned`` atomic ordering, e.g. ``__ATOMIC_ACQUIRE``
+* ``const char *`` synchronization scope, e.g. ``workgroup``
+* Zero or more ``const char *`` address spaces names.
+
+The address spaces arguments must be string literals with known values, such as:
+
+* ``"local"``
+* ``"global"``
+* ``"image"``
+
+If one or more address space name are provided, the code generator will attempt
+to emit potentially faster instructions that only fence those address spaces.
----------------
ssahasra wrote:

This use of "fence" as a verb seems a bit too informal. Reword it to say "instructions that order access to at least those address spaces"? (Note the addition of "at least" to signify a lower bound)

https://github.com/llvm/llvm-project/pull/78572