Expose llvm fence instruction as clang intrinsic

[saiislam]

Summary

Expose llvm fence instruction as clang builtin for AMDGPU target

__builtin_amdgcn_fence(unsigned int memoryOrdering, const char *syncScope)

The first argument of this builtin is one of the memory-ordering specifiers
ATOMIC_ACQUIRE, ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST
following C++11 memory model semantics. This is mapped to corresponding
LLVM atomic memory ordering for the fence instruction using LLVM atomic C
ABI. The second argument is an AMDGPU-specific synchronization scope
defined as string.

[JonChesterfield]

clang/lib/Sema/SemaChecking.cpp

1885

isIntegerType will return true for signed integers as well as unsigned. It seems reasonable to call this with a signed integer type (e.g. '2'), so perhaps the references to unsigned should be dropped from the code and error message

clang/test/CodeGenOpenCL/atomic-ops.cl

291 ↗

(On Diff #249340)

I'm hoping this intrinsic will be usable from C or C++, as well as from OpenCL - please could you add a non-opencl codegen test.

It doesn't need to check all the cases again, just enough to show that the intrinsic and arguments are available (they're spelled like __ATOMIC_SEQ_CST, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES outside of opencl)

[JonChesterfield]

clang/lib/CodeGen/CGBuiltin.cpp

3734

Interesting, fence can't be relaxed

‘fence’ instructions take an ordering argument which defines what synchronizes-with edges they add. They can only be given acquire, release, acq_rel, and seq_cst orderings.

clang/lib/Sema/SemaChecking.cpp

1894

This should reject 'relaxed' - I think that's currently accepted by sema then silently thrown away by codegen

clang/test/SemaOpenCL/atomic-ops.cl

198 ↗

(On Diff #249340)

A happy case too please, e.g. to show that it accepts a couple of integers. Looks like __builtin_memory_fence(2, 2); but without an expected-error comment

[JonChesterfield]

clang/include/clang/Basic/Builtins.def

1516 ↗

(On Diff #249340)

BUILTIN(__builtin_memory_fence, "vii", "n")?

The other fence intrinsics (e.g. __c11_atomic_thread_fence) take signed integers and rely on the built in type checking, which seems reasonable here too

[JonChesterfield]

@jdoerfert this is one of the two intrinsics needed to drop the .ll source from the amdgcn deviceRTL. The other is atomic_inc.

[Anastasia]

clang/lib/Sema/SemaChecking.cpp

1885

I think we should accept implicit type conversion rules from the language...

[sameerds]

The commit summary needs improvement. The syntax is not really necessary there, but instead this needs a better explanation of how this builtin fits in with the overall scheme of language-specific and target-specific details of an atomic operation. For example, is this meant only for OpenCL? Does it work with CUDA? Or HIP? What is the behaviour for scope in C++?

clang/include/clang/Basic/DiagnosticSemaKinds.td

7860–7863 ↗

(On Diff #249340)

Just above this addition, atomic op seems to emit a warning for invalid memory order. Should that be the case with fence too?

clang/lib/CodeGen/CGBuiltin.cpp

3728

The proposed builtin does not claim to be an OpenCL builtin, so it's probably not correct to simply assume the OpenCL model. Should the model be chosen based on the source language specified?

[JonChesterfield]

In D75917#1916160, @sameerds wrote:

how this builtin fits in with the overall scheme of language-specific and target-specific details of an atomic operation. For example, is this meant only for OpenCL? Does it work with CUDA? Or HIP? What is the behaviour for scope in C++?

Identical to the fence instruction. Which is assumed well thought through already, given it's an IR instruction.

As far as I can tell, fence composes sensibly with other IR then generates the right thing at the back end. So it looks fit for purpose, just not currently available from clang.

clang/lib/CodeGen/CGBuiltin.cpp

3728

The only values for AtomicScopeModelKind are none and OpenCL.

[sameerds]

In D75917#1916664, @JonChesterfield wrote:

In D75917#1916160, @sameerds wrote:

how this builtin fits in with the overall scheme of language-specific and target-specific details of an atomic operation. For example, is this meant only for OpenCL? Does it work with CUDA? Or HIP? What is the behaviour for scope in C++?

Identical to the fence instruction. Which is assumed well thought through already, given it's an IR instruction.

As far as I can tell, fence composes sensibly with other IR then generates the right thing at the back end. So it looks fit for purpose, just not currently available from clang.

Well, there is a problem: The LangRef says that scopes are target-defined. This change says that scopes are defined by the high-level language and further assumes that OpenCL scopes make sense in all languages. Besides conflicting with the LangRef, this not seem to work with C++, which has no scopes and nor with CUDA or HIP, whose scopes are not represented in any AtomicScopeModel.

[JonChesterfield]

In D75917#1916972, @sameerds wrote:

Well, there is a problem: The LangRef says that scopes are target-defined. This change says that scopes are defined by the high-level language and further assumes that OpenCL scopes make sense in all languages. Besides conflicting with the LangRef, this not seem to work with C++, which has no scopes and nor with CUDA or HIP, whose scopes are not represented in any AtomicScopeModel.

I don't follow. IR has a fence instruction. This builtin maps directly to it, passing whatever integer arguments were given to the intrinsic along unchanged. It's exactly as valid, or invalid, as said fence instruction.

Are you objecting to passing enums in the test cases instead of raw integers?

Or is your issue with the fence instruction itself?

[sameerds]

In D75917#1917296, @JonChesterfield wrote:

In D75917#1916972, @sameerds wrote:

Well, there is a problem: The LangRef says that scopes are target-defined. This change says that scopes are defined by the high-level language and further assumes that OpenCL scopes make sense in all languages. Besides conflicting with the LangRef, this not seem to work with C++, which has no scopes and nor with CUDA or HIP, whose scopes are not represented in any AtomicScopeModel.

I don't follow. IR has a fence instruction. This builtin maps directly to it, passing whatever integer arguments were given to the intrinsic along unchanged. It's exactly as valid, or invalid, as said fence instruction.

Is it really? The scope argument of the IR fence is a target-specific string:
http://llvm.org/docs/LangRef.html#syncscope

The change that I see here is assuming a numerical argument, and also assuming that the numbers used must conform to the OpenCL enum. That would certainly make the builtin quite different from the IR fence.

[JonChesterfield]

In D75917#1925617, @sameerds wrote:

Is it really? The scope argument of the IR fence is a target-specific string:
http://llvm.org/docs/LangRef.html#syncscope

The change that I see here is assuming a numerical argument, and also assuming that the numbers used must conform to the OpenCL enum. That would certainly make the builtin quite different from the IR fence.

I think I follow. The syncscope takes a string, therefore the builtin that maps onto fence should also take a string for that parameter? That's fine by me. Will help if a new non-opencl syncscope is introduced later.

[sameerds]

In D75917#1925700, @JonChesterfield wrote:

I think I follow. The syncscope takes a string, therefore the builtin that maps onto fence should also take a string for that parameter? That's fine by me. Will help if a new non-opencl syncscope is introduced later.

Right. To be precise, it is a target-specific string, and should not be processed as if it was an OpenCL scope. The builtin should allow anything that the IR fence would allow in a .ll file created for the specified target.

[JonChesterfield]

Can this be revived? Changing the enum to a string still sounds good to me

[b-sumner]

Please go ahead and update to a string for the scope.

[saiislam]

Removed OpenCL specific dependencies

Now it takes target-specific sync scope as an string.

[sameerds]

clang/include/clang/Basic/Builtins.def

1583 ↗

(On Diff #255173)

This should be moved to be near line 786, where atomic builtins are listed under the comment "// GCC does not support these, they are a Clang extension."

clang/lib/CodeGen/CGBuiltin.cpp

14145

There should no mention of any high-level language here. The correct enum to validate against is llvm::AtomicOrdering from llvm/Support/AtomicOrdering.h, and not the C ABI or any other language ABI.

14166

This seems to be creating a new ID for any arbitrary string passed as sync scope. This should be validated against LLVMContext::getSyncScopeNames().

clang/test/CodeGenHIP/builtin_memory_fence.cpp

5 ↗

(On Diff #255173)

There should be a line that tries to do:

__builtin_memory_fence(__ATOMIC_SEQ_CST, "foobar");

9 ↗

(On Diff #255173)

Orderings like __ATOMIC_SEQ_CST are defined for C/C++ memory models. They should not be used with the new builtin because this new builtin does not follow any specific language model. For user convenience, the right thing to do is to introduce new tokens in the Clang preprocessor, similar to the __ATOMIC_* tokens. The convenient shortcut is to just tell the user to supply numerical values by looking at the LLVM source code.

From llvm/Support/AtomicOrdering.h, note how the numerical value for __ATOMIC_SEQ_CST is 5, but the numerical value for the LLVM SequentiallyConsistent ordering is 7. The numerical value 5 refers to the LLVM ordering "release". So, if the implementation were correct, this line should result in the following unexpected LLVM IR:

fence syncscope("workgroup") release

[saiislam]

clang/lib/CodeGen/CGBuiltin.cpp

14145

Even though this builtin is supposed to be language-independent, here intention was to provide interface in terms of well known standard C11/C++11 enums for memory order (__ATOMIC_ACQUIRE, etc.), so that user of the builtin don't have to remember and modify their code. The builtin internally maps it as per the expectation of fence instruction.

14166

As the FE is not aware about specific target and implementation of sync scope for that target, getSyncScopeNames() here returns llvm'sdefault sync scopes, which only supports singlethreaded and system as valid scopes. Validity checking of memory scope string is being intentionally left for the later stages which deal with the generated IR.

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

As you pointed out, the range of acquire to sequentiallly consistent memory orders for llvm::AtomicOrdering is [4, 7], while for llvm::AtomicOrderingCABI is [2, 5]. Enums of C ABI was taken to ensure easy of use for the users who are familiar with C/C++ standard memory model. It allows them to use macros like __ATOMIC_ACQUIRE etc.
Clang CodeGen of the builtin internally maps C ABI ordering to llvm atomic ordering.

[JonChesterfield]

clang/lib/Sema/SemaChecking.cpp

1888

I think I'd write this as a switch over the enum instead of a ranged compare.

It'll codegen to the same thing, but we'll get warnings if more values are introduced to the enum and things will keep working (here, anyway) if the values are reordered.

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

What language, implemented in clang, do you have in mind that reusing the existing __ATOMIC_* macros would be incorrect for?

[saiislam]

Changes:

1. Moved builtin definition with rest of atomic builtins
2. Updated validity checking of memory order using exact mathches instead of range matching
3. Added a sucessful test case which passes arbitrary string scope
4. Corrected formatting

[sameerds]

clang/lib/CodeGen/CGBuiltin.cpp

14166

That's pretty strange. At this point, Clang should know what the target is, and it should have a chance to update the list of sync scopes somewhere. @arsenm, do you see a way around this?

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

I think we agreed that this builtin exposes the LLVM fence exactly. That would mean it takes arguments defined by LLVM. If you are implementing something different from that, then it first needs to be specified properly. Perhaps you could say that this is a C ABI compatible builtin, that happens to take target specific scopes? That should cover OpenCL whose scope enum is designed to be compatible with C.

Whatever it is that you are trying to implement here, it definitely does not expose a raw LLVM fence.

[JonChesterfield]

clang/lib/CodeGen/CGBuiltin.cpp

14166

There is already sufficient IR level checking for the string at the instruction level. Warning in clang as well could be a nicer user experience, but that seems low priority to me.

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

The llvm fence, in text form, uses a symbol for the memory scope. Not an enum.

This symbol is set using these macros for the existing atomic builtins. Using an implementation detail of clang instead is strictly worse, by layering and by precedent.

ABI is not involved here. Nor is OpenCl.

[sameerds]

clang/lib/CodeGen/CGBuiltin.cpp

14166

If there is some checking happening anywhere, then that needs to be demonstrated in a testcase where the input high-level program passes an illegal string as the scope argument.

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

The __ATOMIC_* symbols in Clang quite literally represent the C/C++ ABI. See the details in AtomicOrdering.h and InitPreprocessor.cpp. I am not opposed to specifying that the builtin expects these symbols, but then it is incorrect to say that the builtin exposes the raw LLVM builtin. It is a C-ABI-compatible builtin that happens to take target-specific scope as a string argument. And that would also make it an overload of the already existing builting __atomic_fence().

[JonChesterfield]

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

I don't know what you mean by "raw", but am guessing you're asking for documentation for the intrinsic. Said documentation should indeed be added for this builtin - it'll probably be in a tablegen file.

[JonChesterfield]

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

I will try to stop using builtin and intrinsic as synonyms.

[JonChesterfield]

clang/test/CodeGenHIP/builtin_memory_fence.cpp

1 ↗

(On Diff #255288)

Codegen test should be under CodeGen and/or CodeGenCXX

[sameerds]

clang/test/CodeGenHIP/builtin_memory_fence.cpp

9 ↗

(On Diff #255173)

Right. It's actually an LLVM instruction, not even an intrinsic. I am guilty of using the wrong term quite often, but usually the context helps. I think the following is still needed:

1. A testcase that exercises the builtin with an invalid string argument for scope.
2. An update to the change description describing what is being introduced here. It is more than a direct mapping from builtin to instruction. The C ABI is involved.
3. An update to the Clang documentation describing this new builtin: https://clang.llvm.org/docs/LanguageExtensions.html#builtin-functions

[b-sumner]

In addition to predefining __ATOMIC_RELAXED, etc., clang also predefines __OPENCL_MEMORY_SCOPE_WORK_ITEM and friends. So it doesn't really seem unreasonable for clang to also predefine its known syncscopes, and to require the argument to be one of those integers.

[saiislam]

1. Moved test case to clang/test/CodeGenCXX.
2. Added a failing test case with invalid sync scope, which gets detected by implementation of fence instruction.
3. Updated the change description of the builtin.
4. Updated the clang documentation describing mapping of C++ memory-ordering to LLVM memory-ordering.

[JonChesterfield]

The tests look good, but I can't see the implementation in this diff. Maybe a file missing from the patch? Can be hard to tell with phabricator, the error may be at my end.

[saiislam]

Changed the builtin to be AMDGCN-specific

It is named as __builtin_amdgcn_fence(order, scope)

[saiislam]

Removed stale commented code

[JonChesterfield]

Amdgcn specific is fine by me. Hopefully that unblocks this.

[arsenm]

clang/test/CodeGenCXX/builtin-amdgcn-fence-failure.cpp

5 ↗

(On Diff #259239)

Does this really depend on C++? Can it use OpenCL like the other builtin tests?This also belongs in a Sema* test directory since it's checking an error

[JonChesterfield]

clang/test/CodeGenCXX/builtin-amdgcn-fence-failure.cpp

5 ↗

(On Diff #259239)

Making it opencl-only would force some of the openmp runtime to be written in opencl, which is not presently the case. Currently that library is written in a dialect of hip, but there's a plan to implement it in openmp instead.

I'd much rather this builtin work from any language, instead of tying it to opencl, as that means one can use it from openmp target regions.

[b-sumner]

clang/test/CodeGenCXX/builtin-amdgcn-fence-failure.cpp

5 ↗

(On Diff #259239)

I thought the question was about this test itself. The test being in CodeGenOpenCL doesn't affect whether other languages can use the builtin. Why not put this test in CodeGenOpenCL alongside all of the other builtins-amdgcn-*.cl ?

[saiislam]

Moved builtin-amdgcn-fence-failure.cpp from clang/test/CodeGenCXX/
to clang/test/Sema/ since it is checking an error.

[saiislam]

clang/test/CodeGenCXX/builtin-amdgcn-fence-failure.cpp

5 ↗

(On Diff #259239)

This test is basically relying on implementation of sync scope in the AMDGCN backend to check for validity, and is not testing the code generation capability. Doesn't it make more sense in Sema directory?

Won't putting it in SemaOpenCL or SemaOpenCLCXX indicate some kind of relation with OpenCL?

[sameerds]

Can you please submit a squashed single commit for review against the master branch? All the recent commits seem to be relative to previous commits, and I am having trouble looking at the change as a whole.

The commit description needs some fixes:

1. Remove use of Title Case, in places like "using Fence instruction" and "LLVM Atomic Memory Ordering" and "as a String". They are simply common nouns, not proper nouns. When in doubt, look at the LangRef to see how these words are used as common nouns.
2. Don't mention that enum values are okay too. I am sure they will always be okay, but it's better to encourage the use of symbols.
3. Don't mention HSA even if the AMDGPU user manual does so.
4. In fact the last two sentences in the description are not strictly necessary ... they are the logical outcome of the scopes being target-specific strings.
5. Note that the general practice in documentation is to say "AMDGPU" which covers the LLVM Target, while "amdgcn" is only used in the code because it is one of multiple architectures in the AMDGPU backend.

clang/docs/LanguageExtensions.rst

2458 ↗

(On Diff #259298)

We probably don't need to document the new builtin here. Clearly, we have not documented any other AMDGPU builtin, and there is no need to start now. If necessary, we can take that up as a separate task covering all builtins.

[saiislam]

Removed documentation from clang doc. Squashed all changes into a single commit.

[saiislam]

Added check and test for sync scope to be a constant literal.

[saiislam]

clang/lib/Sema/SemaChecking.cpp

3070

@sameerds here is the check for sync scope to be constant literal

[sameerds]

Thanks @saiislam ... this looks much better!

Two nitpicks, that must be fixed. But it is okay if you directly submit after fixing them.

1. The change description should use "const char *" in the signature and not "String".
2. Can you please add a test that passes an integer constant as the scope? I am assuming that the signature check will complain that it is not a string.

[saiislam]

Updated description and added a failing test case for integer scope.

clang/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
 BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n")
+BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
 
 // FIXME: Need to disallow constant address space.
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
 BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
 BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
 BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")

clang/include/clang/Sema/Sema.h

private:
   bool CheckMipsBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckMipsBuiltinCpu(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckMipsBuiltinArgument(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckSystemZBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckX86BuiltinRoundingOrSAE(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckX86BuiltinGatherScatterScale(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckX86BuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
   bool CheckPPCBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
+  bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
 
   bool SemaBuiltinVAStart(unsigned BuiltinID, CallExpr *TheCall);
   bool SemaBuiltinVAStartARMMicrosoft(CallExpr *Call);
   bool SemaBuiltinUnorderedCompare(CallExpr *TheCall);
   bool SemaBuiltinFPClassification(CallExpr *TheCall, unsigned NumArgs);
   bool SemaBuiltinVSX(CallExpr *TheCall);
   bool SemaBuiltinOSLogFormat(CallExpr *TheCall);
 

clang/lib/CodeGen/CGBuiltin.cpp

 #include "clang/AST/Attr.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/OSLog.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/CodeGen/CGFunctionInfo.h"
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/ADT/StringExtras.h"
+#include "llvm/Analysis/ValueTracking.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/IntrinsicsAArch64.h"
 #include "llvm/IR/IntrinsicsAMDGPU.h"
 #include "llvm/IR/IntrinsicsARM.h"
 #include "llvm/IR/IntrinsicsBPF.h"
 #include "llvm/IR/IntrinsicsHexagon.h"
case Builtin::BI__builtin_coro_param:
     return EmitCoroutineIntrinsic(E, Intrinsic::coro_param);
 
   // OpenCL v2.0 s6.13.16.2, Built-in pipe read and write functions
   case Builtin::BIread_pipe:
   case Builtin::BIwrite_pipe: {
     Value *Arg0 = EmitScalarExpr(E->getArg(0)),
           *Arg1 = EmitScalarExpr(E->getArg(1));
     CGOpenCLRuntime OpenCLRT(CGM);
     Value *PacketSize = OpenCLRT.getPipeElemSize(E->getArg(0));
     Value *PacketAlign = OpenCLRT.getPipeElemAlign(E->getArg(0));
 
     // Type of the generic packet parameter.
     unsigned GenericAS =
         getContext().getTargetAddressSpace(LangAS::opencl_generic);
     llvm::Type *I8PTy = llvm::PointerType::get(
         llvm::Type::getInt8Ty(getLLVMContext()), GenericAS);
 
     // Testing which overloaded version we should generate the call for.
     if (2U == E->getNumArgs()) {
       const char *Name = (BuiltinID == Builtin::BIread_pipe) ? "__read_pipe_2"
                                                              : "__write_pipe_2";
       // Creating a generic function type to be able to call with any builtin or
       // user defined type.
case AMDGPU::BI__builtin_r600_read_tidig_z:
     return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0, 1024);
   case AMDGPU::BI__builtin_amdgcn_alignbit: {
     llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
     llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
     llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
     Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());
     return Builder.CreateCall(F, { Src0, Src1, Src2 });
   }
+
+  case AMDGPU::BI__builtin_amdgcn_fence: {
+    llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;
+    llvm::SyncScope::ID SSID;
+    Value *Order = EmitScalarExpr(E->getArg(0));
+    Value *Scope = EmitScalarExpr(E->getArg(1));
+
+    if (isa<llvm::ConstantInt>(Order)) {
+      int ord = cast<llvm::ConstantInt>(Order)->getZExtValue();
+
+      // Map C11/C++11 memory ordering to LLVM memory ordering
+      switch (static_cast<llvm::AtomicOrderingCABI>(ord)) {
+      case llvm::AtomicOrderingCABI::acquire:
+        AO = llvm::AtomicOrdering::Acquire;
+        break;
+      case llvm::AtomicOrderingCABI::release:
+        AO = llvm::AtomicOrdering::Release;
+        break;
+      case llvm::AtomicOrderingCABI::acq_rel:
+        AO = llvm::AtomicOrdering::AcquireRelease;
+        break;
+      case llvm::AtomicOrderingCABI::seq_cst:
+        AO = llvm::AtomicOrdering::SequentiallyConsistent;
+        break;
+      case llvm::AtomicOrderingCABI::consume: // not supported by LLVM fence
+      case llvm::AtomicOrderingCABI::relaxed: // not supported by LLVM fence
+        break;
+      }
+
+      StringRef scp;
+      llvm::getConstantStringInfo(Scope, scp);
+      SSID = getLLVMContext().getOrInsertSyncScopeID(scp);
+
+      return Builder.CreateFence(AO, SSID);
+    }
+    LLVM_FALLTHROUGH;
+  }
   default:
     return nullptr;
   }
 }
 
 /// Handle a SystemZ function in which the final argument is a pointer
 /// to an int that receives the post-instruction CC value.  At the LLVM level
 /// this is represented as a function that returns a {result, cc} pair.

clang/lib/Sema/SemaChecking.cpp

#include "clang/Basic/Builtins.def"
   if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
     switch (Context.getTargetInfo().getTriple().getArch()) {
       case llvm::Triple::arm:
       case llvm::Triple::armeb:
       case llvm::Triple::thumb:
       case llvm::Triple::thumbeb:
         if (CheckARMBuiltinFunctionCall(BuiltinID, TheCall))
           return ExprError();
         break;
       case llvm::Triple::aarch64:
       case llvm::Triple::aarch64_32:
       case llvm::Triple::aarch64_be:
         if (CheckAArch64BuiltinFunctionCall(BuiltinID, TheCall))
           return ExprError();
         break;
       case llvm::Triple::bpfeb:
       case llvm::Triple::bpfel:
         if (CheckBPFBuiltinFunctionCall(BuiltinID, TheCall))
           return ExprError();
         break;
       case llvm::Triple::hexagon:
         if (CheckHexagonBuiltinFunctionCall(BuiltinID, TheCall))
           return ExprError();
         break;
       case llvm::Triple::mips:
       case llvm::Triple::mipsel:
switch (Context.getTargetInfo().getTriple().getArch()) {
           return ExprError();
         break;
       case llvm::Triple::ppc:
       case llvm::Triple::ppc64:
       case llvm::Triple::ppc64le:
         if (CheckPPCBuiltinFunctionCall(BuiltinID, TheCall))
           return ExprError();
         break;
+      case llvm::Triple::amdgcn:
+        if (CheckAMDGCNBuiltinFunctionCall(BuiltinID, TheCall))
+          return ExprError();
+        break;
       default:
         break;
     }
   }
 
   return TheCallResult;
 }
 
case PPC::BI__builtin_unpack_vector_int128:
     return SemaVSXCheck(TheCall) ||
            SemaBuiltinConstantArgRange(TheCall, 1, 0, 1);
   case PPC::BI__builtin_pack_vector_int128:
     return SemaVSXCheck(TheCall);
   }
   return SemaBuiltinConstantArgRange(TheCall, i, l, u);
 }
 
+bool Sema::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
+                                          CallExpr *TheCall) {
+  switch (BuiltinID) {
+  case AMDGPU::BI__builtin_amdgcn_fence: {
+    ExprResult Arg = TheCall->getArg(0);
+    auto ArgExpr = Arg.get();
+    Expr::EvalResult ArgResult;
+
+    if (!ArgExpr->EvaluateAsInt(ArgResult, Context))
+      return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
+             << ArgExpr->getType();
+    int ord = ArgResult.Val.getInt().getZExtValue();
+
+    // Check valididty of memory ordering as per C11 / C++11's memody model.
+    switch (static_cast<llvm::AtomicOrderingCABI>(ord)) {
+    case llvm::AtomicOrderingCABI::acquire:
+    case llvm::AtomicOrderingCABI::release:
+    case llvm::AtomicOrderingCABI::acq_rel:
+    case llvm::AtomicOrderingCABI::seq_cst:
+      break;
+    default: {
+      return Diag(ArgExpr->getBeginLoc(),
+                  diag::warn_atomic_op_has_invalid_memory_order)
+             << ArgExpr->getSourceRange();
+    }
+    }
+
+    Arg = TheCall->getArg(1);
+    ArgExpr = Arg.get();
+    Expr::EvalResult ArgResult1;
+    // Check that sync scope is a constant literal
+    if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, Expr::EvaluateForCodeGen,
+                                         Context))
+      return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
+             << ArgExpr->getType();
+  } break;
+  }
+  return false;
+}
+
 bool Sema::CheckSystemZBuiltinFunctionCall(unsigned BuiltinID,
                                            CallExpr *TheCall) {
   if (BuiltinID == SystemZ::BI__builtin_tabort) {
     Expr *Arg = TheCall->getArg(0);
     llvm::APSInt AbortCode(32);
     if (Arg->isIntegerConstantExpr(AbortCode, Context) &&
         AbortCode.getSExtValue() >= 0 && AbortCode.getSExtValue() < 256)
       return Diag(Arg->getBeginLoc(), diag::err_systemz_invalid_tabort_code)

clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
+// RUN:   -triple=amdgcn-amd-amdhsa  | opt -S | FileCheck %s
+
+void test_memory_fence_success() {
+  // CHECK-LABEL: test_memory_fence_success
+
+  // CHECK: fence syncscope("workgroup") seq_cst
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+
+  // CHECK: fence syncscope("agent") acquire
+  __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
+
+  // CHECK: fence seq_cst
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+
+  // CHECK: fence syncscope("agent") acq_rel
+  __builtin_amdgcn_fence(4, "agent");
+
+  // CHECK: fence syncscope("workgroup") release
+  __builtin_amdgcn_fence(3, "workgroup");
+}

clang/test/Sema/builtin-amdgcn-fence-failure.cpp

+// REQUIRES: amdgpu-registered-target
+// RUN: not %clang_cc1 %s -S \
+// RUN:   -triple=amdgcn-amd-amdhsa 2>&1 | FileCheck %s
+
+void test_amdgcn_fence_failure() {
+
+  // CHECK: error: Unsupported atomic synchronization scope
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "foobar");
+}

clang/test/SemaOpenCL/builtins-amdgcn-error.cl

void test_ds_fminf(local float *out, float src, int a) {
   *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
 }
 
 void test_ds_fmaxf(local float *out, float src, int a) {
   *out = __builtin_amdgcn_ds_fmaxf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
   *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
   *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
 }
+
+void test_fence() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST + 1, "workgroup"); // expected-warning {{memory order argument to atomic operation is invalid}}
+  __builtin_amdgcn_fence(__ATOMIC_ACQUIRE - 1, "workgroup"); // expected-warning {{memory order argument to atomic operation is invalid}}
+  __builtin_amdgcn_fence(4); // expected-error {{too few arguments to function call, expected 2}}
+  __builtin_amdgcn_fence(4, 4, 4); // expected-error {{too many arguments to function call, expected 2}}
+  __builtin_amdgcn_fence(3.14, ""); // expected-warning {{implicit conversion from 'double' to 'unsigned int' changes value from 3.14 to 3}}
+  __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, 5); // expected-warning {{incompatible integer to pointer conversion passing 'int' to parameter of type 'const char *'}}
+  const char ptr[] = "workgroup";
+  __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, ptr); // expected-error {{expression is not a string literal}}
+}