================ @@ -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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
