Index: clang/include/clang/AST/ASTContext.h
===================================================================
--- clang/include/clang/AST/ASTContext.h
+++ clang/include/clang/AST/ASTContext.h
@@ -39,6 +39,7 @@
 #include "clang/Basic/SanitizerBlacklist.h"
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Basic/Specifiers.h"
+#include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/XRayLists.h"
 #include "llvm/ADT/APSInt.h"
 #include "llvm/ADT/ArrayRef.h"
@@ -671,7 +672,9 @@
   /// Returns empty type if there is no appropriate target types.
   QualType getRealTypeForBitwidth(unsigned DestWidth, bool ExplicitIEEE) const;
 
-  bool AtomicUsesUnsupportedLibcall(const AtomicExpr *E) const;
+  TargetInfo::AtomicSupportKind
+  getTargetAtomicSupport(TargetInfo::AtomicOperationKind Op,
+                         QualType AtomicTy) const;
 
   const LangOptions& getLangOpts() const { return LangOpts; }
 
Index: clang/include/clang/AST/Expr.h
===================================================================
--- clang/include/clang/AST/Expr.h
+++ clang/include/clang/AST/Expr.h
@@ -26,6 +26,7 @@
 #include "clang/Basic/CharInfo.h"
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/SyncScope.h"
+#include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/TypeTraits.h"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/APSInt.h"
@@ -6275,6 +6276,10 @@
   std::unique_ptr<AtomicScopeModel> getScopeModel() const {
     return getScopeModel(getOp());
   }
+
+  /// Get the target atomic operation kind which can be used to query target
+  /// atomic support.
+  TargetInfo::AtomicOperationKind getTargetAtomicOp() const;
 };
 
 /// TypoExpr - Internal placeholder for expressions where typo correction
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8156,6 +8156,9 @@
 def err_atomic_op_needs_trivial_copy : Error<
   "address argument to atomic operation must be a pointer to a "
   "trivially-copyable type (%0 invalid)">;
+def err_atomic_op_needs_atomic_int_ptr_or_fp : Error<
+  "address argument to atomic operation must be a pointer to %select{|atomic }0"
+  "integer, pointer or supported floating point type (%1 invalid)">;
 def err_atomic_op_needs_atomic_int_or_ptr : Error<
   "address argument to atomic operation must be a pointer to %select{|atomic }0"
   "integer or pointer (%1 invalid)">;
@@ -8180,9 +8183,9 @@
   "__builtin_mul_overflow does not support signed _ExtInt operands of more "
   "than %0 bits">;
 
-def err_atomic_load_store_uses_lib : Error<
-  "atomic %select{load|store}0 requires runtime support that is not "
-  "available for this target">;
+def err_atomic_op_unsupported : Error<
+  "atomic %select{init|C11 load/store|load/store|add/sub|min/max|logic op|exchange|compare/exchange}0"
+  " of %1 type requires runtime support that is not available for this target">;
 
 def err_nontemporal_builtin_must_be_pointer : Error<
   "address argument to nontemporal builtin must be a pointer (%0 invalid)">;
Index: clang/include/clang/Basic/TargetInfo.h
===================================================================
--- clang/include/clang/Basic/TargetInfo.h
+++ clang/include/clang/Basic/TargetInfo.h
@@ -686,15 +686,6 @@
   /// Set the maximum inline or promote width lock-free atomic operation
   /// for the given target.
   virtual void setMaxAtomicWidth() {}
-  /// Returns true if the given target supports lock-free atomic
-  /// operations at the specified width and alignment.
-  virtual bool hasBuiltinAtomic(uint64_t AtomicSizeInBits,
-                                uint64_t AlignmentInBits) const {
-    return AtomicSizeInBits <= AlignmentInBits &&
-           AtomicSizeInBits <= getMaxAtomicInlineWidth() &&
-           (AtomicSizeInBits <= getCharWidth() ||
-            llvm::isPowerOf2_64(AtomicSizeInBits / getCharWidth()));
-  }
 
   /// Return the maximum vector alignment supported for the given target.
   unsigned getMaxVectorAlign() const { return MaxVectorAlign; }
@@ -1489,6 +1480,41 @@
   /// Whether target allows debuginfo types for decl only variables.
   virtual bool allowDebugInfoForExternalVar() const { return false; }
 
+  /// Abstraction of source level atomic operations.
+  enum class AtomicOperationKind {
+    Init,
+    C11LoadStore,
+    LoadStore,
+    AddSub,
+    MinMax,
+    LogicOp,
+    Xchg,
+    CmpXchg,
+  };
+
+  /// What is emitted in LLVM IR by clang for the atomic operation:
+  /// LockFree - LLVM atomic instructions
+  /// InlineWithLock - LLVM instructions but not lock free
+  /// Library - call of library functions
+  /// Unsupported - diagnostics
+  enum class AtomicSupportKind {
+    LockFree,
+    InlineWithLock,
+    Library,
+    Unsupported,
+  };
+
+  /// Support of floating point atomic add/sub operations by the target.
+  virtual AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const;
+
+  /// Support of atomic operations by the target. If \p FS is Bogus, the atomic
+  /// type is not a floating point type.
+  virtual AtomicSupportKind
+  getAtomicSupport(AtomicOperationKind Op, uint64_t AtomicWidthInBits,
+                   uint64_t AlignmentInBits,
+                   const llvm::fltSemantics &FS = llvm::APFloat::Bogus()) const;
+
 protected:
   /// Copy type and layout related info.
   void copyAuxTarget(const TargetInfo *Aux);
Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ clang/lib/AST/ASTContext.cpp
@@ -11092,22 +11092,21 @@
   return New;
 }
 
-bool ASTContext::AtomicUsesUnsupportedLibcall(const AtomicExpr *E) const {
-  const llvm::Triple &T = getTargetInfo().getTriple();
-  if (!T.isOSDarwin())
-    return false;
-
-  if (!(T.isiOS() && T.isOSVersionLT(7)) &&
-      !(T.isMacOSX() && T.isOSVersionLT(10, 9)))
-    return false;
-
-  QualType AtomicTy = E->getPtr()->getType()->getPointeeType();
-  CharUnits sizeChars = getTypeSizeInChars(AtomicTy);
-  uint64_t Size = sizeChars.getQuantity();
-  CharUnits alignChars = getTypeAlignInChars(AtomicTy);
-  unsigned Align = alignChars.getQuantity();
-  unsigned MaxInlineWidthInBits = getTargetInfo().getMaxAtomicInlineWidth();
-  return (Size != Align || toBits(sizeChars) > MaxInlineWidthInBits);
+TargetInfo::AtomicSupportKind
+ASTContext::getTargetAtomicSupport(TargetInfo::AtomicOperationKind TargetOp,
+                                   QualType AtomicTy) const {
+  AtomicTy = AtomicTy.getCanonicalType();
+  auto ValTy = AtomicTy;
+  if (ValTy->isAtomicType())
+    ValTy = ValTy->getAs<AtomicType>()->getValueType();
+  auto AtomicTI = getTypeInfo(AtomicTy);
+  uint64_t AtomicWidthInBits = AtomicTI.Width;
+  uint64_t AtomicAlignInBits = AtomicTI.Align;
+  const llvm::fltSemantics &FS = ValTy->isRealFloatingType()
+                                     ? getFloatTypeSemantics(ValTy)
+                                     : llvm::APFloat::Bogus();
+  return getTargetInfo().getAtomicSupport(TargetOp, AtomicWidthInBits,
+                                          AtomicAlignInBits, FS);
 }
 
 bool
Index: clang/lib/AST/Expr.cpp
===================================================================
--- clang/lib/AST/Expr.cpp
+++ clang/lib/AST/Expr.cpp
@@ -4606,6 +4606,76 @@
   return T;
 }
 
+TargetInfo::AtomicOperationKind AtomicExpr::getTargetAtomicOp() const {
+  switch (getOp()) {
+  case AtomicExpr::AO__c11_atomic_init:
+  case AtomicExpr::AO__opencl_atomic_init:
+    return TargetInfo::AtomicOperationKind::Init;
+
+  case AtomicExpr::AO__c11_atomic_load:
+  case AtomicExpr::AO__opencl_atomic_load:
+  case AtomicExpr::AO__c11_atomic_store:
+  case AtomicExpr::AO__opencl_atomic_store:
+    return TargetInfo::AtomicOperationKind::C11LoadStore;
+
+  case AtomicExpr::AO__atomic_load_n:
+  case AtomicExpr::AO__atomic_load:
+  case AtomicExpr::AO__atomic_store:
+  case AtomicExpr::AO__atomic_store_n:
+    return TargetInfo::AtomicOperationKind::LoadStore;
+
+  case AtomicExpr::AO__c11_atomic_fetch_add:
+  case AtomicExpr::AO__c11_atomic_fetch_sub:
+  case AtomicExpr::AO__opencl_atomic_fetch_add:
+  case AtomicExpr::AO__opencl_atomic_fetch_sub:
+  case AtomicExpr::AO__atomic_fetch_add:
+  case AtomicExpr::AO__atomic_fetch_sub:
+  case AtomicExpr::AO__atomic_add_fetch:
+  case AtomicExpr::AO__atomic_sub_fetch:
+    return TargetInfo::AtomicOperationKind::AddSub;
+
+  case AtomicExpr::AO__c11_atomic_fetch_and:
+  case AtomicExpr::AO__c11_atomic_fetch_or:
+  case AtomicExpr::AO__c11_atomic_fetch_xor:
+  case AtomicExpr::AO__opencl_atomic_fetch_and:
+  case AtomicExpr::AO__opencl_atomic_fetch_or:
+  case AtomicExpr::AO__opencl_atomic_fetch_xor:
+  case AtomicExpr::AO__atomic_fetch_and:
+  case AtomicExpr::AO__atomic_fetch_or:
+  case AtomicExpr::AO__atomic_fetch_xor:
+  case AtomicExpr::AO__atomic_fetch_nand:
+  case AtomicExpr::AO__atomic_and_fetch:
+  case AtomicExpr::AO__atomic_or_fetch:
+  case AtomicExpr::AO__atomic_xor_fetch:
+  case AtomicExpr::AO__atomic_nand_fetch:
+    return TargetInfo::AtomicOperationKind::LogicOp;
+
+  case AtomicExpr::AO__c11_atomic_fetch_min:
+  case AtomicExpr::AO__c11_atomic_fetch_max:
+  case AtomicExpr::AO__opencl_atomic_fetch_min:
+  case AtomicExpr::AO__opencl_atomic_fetch_max:
+  case AtomicExpr::AO__atomic_min_fetch:
+  case AtomicExpr::AO__atomic_max_fetch:
+  case AtomicExpr::AO__atomic_fetch_min:
+  case AtomicExpr::AO__atomic_fetch_max:
+    return TargetInfo::AtomicOperationKind::MinMax;
+
+  case AtomicExpr::AO__c11_atomic_exchange:
+  case AtomicExpr::AO__opencl_atomic_exchange:
+  case AtomicExpr::AO__atomic_exchange_n:
+  case AtomicExpr::AO__atomic_exchange:
+    return TargetInfo::AtomicOperationKind::Xchg;
+
+  case AtomicExpr::AO__c11_atomic_compare_exchange_strong:
+  case AtomicExpr::AO__c11_atomic_compare_exchange_weak:
+  case AtomicExpr::AO__opencl_atomic_compare_exchange_strong:
+  case AtomicExpr::AO__opencl_atomic_compare_exchange_weak:
+  case AtomicExpr::AO__atomic_compare_exchange:
+  case AtomicExpr::AO__atomic_compare_exchange_n:
+    return TargetInfo::AtomicOperationKind::CmpXchg;
+  }
+}
+
 QualType OMPArraySectionExpr::getBaseOriginalType(const Expr *Base) {
   unsigned ArraySectionCount = 0;
   while (auto *OASE = dyn_cast<OMPArraySectionExpr>(Base->IgnoreParens())) {
Index: clang/lib/Basic/TargetInfo.cpp
===================================================================
--- clang/lib/Basic/TargetInfo.cpp
+++ clang/lib/Basic/TargetInfo.cpp
@@ -844,3 +844,23 @@
   auto *Src = static_cast<const TransferrableTargetInfo*>(Aux);
   *Target = *Src;
 }
+
+TargetInfo::AtomicSupportKind
+TargetInfo::getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const {
+  return AtomicSupportKind::Unsupported;
+}
+
+TargetInfo::AtomicSupportKind
+TargetInfo::getAtomicSupport(AtomicOperationKind Op, uint64_t AtomicWidthInBits,
+                             uint64_t AlignmentInBits,
+                             const llvm::fltSemantics &FS) const {
+  if (&FS != &llvm::APFloat::Bogus() && Op == AtomicOperationKind::AddSub)
+    return getFPAtomicAddSubSupport(FS);
+
+  return AtomicWidthInBits <= AlignmentInBits &&
+                 AtomicWidthInBits <= getMaxAtomicInlineWidth() &&
+                 (AtomicWidthInBits <= getCharWidth() ||
+                  llvm::isPowerOf2_64(AtomicWidthInBits / getCharWidth()))
+             ? AtomicSupportKind::LockFree
+             : AtomicSupportKind::Library;
+}
Index: clang/lib/Basic/Targets/AArch64.h
===================================================================
--- clang/lib/Basic/Targets/AArch64.h
+++ clang/lib/Basic/Targets/AArch64.h
@@ -136,6 +136,17 @@
   bool hasInt128Type() const override;
 
   bool hasExtIntType() const override { return true; }
+
+  AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const override {
+    switch (llvm::APFloat::SemanticsToEnum(FS)) {
+    case llvm::APFloat::S_IEEEsingle:
+    case llvm::APFloat::S_IEEEdouble:
+      return AtomicSupportKind::LockFree;
+    default:
+      return AtomicSupportKind::Unsupported;
+    }
+  }
 };
 
 class LLVM_LIBRARY_VISIBILITY AArch64leTargetInfo : public AArch64TargetInfo {
Index: clang/lib/Basic/Targets/AMDGPU.h
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.h
+++ clang/lib/Basic/Targets/AMDGPU.h
@@ -431,6 +431,28 @@
     return getCanonicalTargetID(getArchNameAMDGCN(GPUKind),
                                 OffloadArchFeatures);
   }
+
+  AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const override {
+    switch (llvm::APFloat::SemanticsToEnum(FS)) {
+    case llvm::APFloat::S_IEEEsingle:
+    case llvm::APFloat::S_IEEEdouble:
+      return AtomicSupportKind::LockFree;
+    default:
+      return AtomicSupportKind::Unsupported;
+    }
+  }
+
+  AtomicSupportKind
+  getAtomicSupport(AtomicOperationKind Op, uint64_t AtomicSizeInBits,
+                   uint64_t AlignmentInBits,
+                   const llvm::fltSemantics &FS) const override {
+    auto Res =
+        TargetInfo::getAtomicSupport(Op, AtomicSizeInBits, AlignmentInBits, FS);
+    if (Res == AtomicSupportKind::Library)
+      Res = AtomicSupportKind::Unsupported;
+    return Res;
+  }
 };
 
 } // namespace targets
Index: clang/lib/Basic/Targets/ARM.h
===================================================================
--- clang/lib/Basic/Targets/ARM.h
+++ clang/lib/Basic/Targets/ARM.h
@@ -188,6 +188,16 @@
   bool hasExtIntType() const override { return true; }
   
   const char *getBFloat16Mangling() const override { return "u6__bf16"; };
+
+  AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const override {
+    switch (llvm::APFloat::SemanticsToEnum(FS)) {
+    case llvm::APFloat::S_IEEEsingle:
+      return AtomicSupportKind::LockFree;
+    default:
+      return AtomicSupportKind::Unsupported;
+    }
+  }
 };
 
 class LLVM_LIBRARY_VISIBILITY ARMleTargetInfo : public ARMTargetInfo {
Index: clang/lib/Basic/Targets/Hexagon.h
===================================================================
--- clang/lib/Basic/Targets/Hexagon.h
+++ clang/lib/Basic/Targets/Hexagon.h
@@ -140,6 +140,17 @@
   }
 
   bool hasExtIntType() const override { return true; }
+
+  AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const override {
+    switch (llvm::APFloat::SemanticsToEnum(FS)) {
+    case llvm::APFloat::S_IEEEsingle:
+    case llvm::APFloat::S_IEEEdouble:
+      return AtomicSupportKind::LockFree;
+    default:
+      return AtomicSupportKind::Unsupported;
+    }
+  }
 };
 } // namespace targets
 } // namespace clang
Index: clang/lib/Basic/Targets/Mips.h
===================================================================
--- clang/lib/Basic/Targets/Mips.h
+++ clang/lib/Basic/Targets/Mips.h
@@ -407,6 +407,17 @@
 
   bool validateTarget(DiagnosticsEngine &Diags) const override;
   bool hasExtIntType() const override { return true; }
+
+  AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const override {
+    switch (llvm::APFloat::SemanticsToEnum(FS)) {
+    case llvm::APFloat::S_IEEEsingle:
+    case llvm::APFloat::S_IEEEdouble:
+      return AtomicSupportKind::LockFree;
+    default:
+      return AtomicSupportKind::Unsupported;
+    }
+  }
 };
 } // namespace targets
 } // namespace clang
Index: clang/lib/Basic/Targets/OSTargets.h
===================================================================
--- clang/lib/Basic/Targets/OSTargets.h
+++ clang/lib/Basic/Targets/OSTargets.h
@@ -173,6 +173,22 @@
                            : TargetInfo::UnsignedLongLong)
                : TargetInfo::getLeastIntTypeByWidth(BitWidth, IsSigned);
   }
+
+  TargetInfo::AtomicSupportKind
+  getAtomicSupport(TargetInfo::AtomicOperationKind Op,
+                   uint64_t AtomicWidthInBits, uint64_t AlignmentInBits,
+                   const llvm::fltSemantics &FS) const override {
+    const llvm::Triple &T = this->getTriple();
+    if (Op == TargetInfo::AtomicOperationKind::C11LoadStore &&
+        ((T.isiOS() && T.isOSVersionLT(7)) ||
+         (T.isMacOSX() && T.isOSVersionLT(10, 9))) &&
+        (AtomicWidthInBits != AlignmentInBits ||
+         AtomicWidthInBits > this->getMaxAtomicInlineWidth())) {
+      return TargetInfo::AtomicSupportKind::Unsupported;
+    }
+    return TargetInfo::getAtomicSupport(Op, AtomicWidthInBits, AlignmentInBits,
+                                        FS);
+  }
 };
 
 // DragonFlyBSD Target
Index: clang/lib/Basic/Targets/X86.h
===================================================================
--- clang/lib/Basic/Targets/X86.h
+++ clang/lib/Basic/Targets/X86.h
@@ -374,6 +374,17 @@
   uint64_t getPointerAlignV(unsigned AddrSpace) const override {
     return getPointerWidthV(AddrSpace);
   }
+
+  AtomicSupportKind
+  getFPAtomicAddSubSupport(const llvm::fltSemantics &FS) const override {
+    switch (llvm::APFloat::SemanticsToEnum(FS)) {
+    case llvm::APFloat::S_IEEEsingle:
+    case llvm::APFloat::S_IEEEdouble:
+      return AtomicSupportKind::LockFree;
+    default:
+      return AtomicSupportKind::Unsupported;
+    }
+  }
 };
 
 // X86-32 generic target
Index: clang/lib/CodeGen/CGAtomic.cpp
===================================================================
--- clang/lib/CodeGen/CGAtomic.cpp
+++ clang/lib/CodeGen/CGAtomic.cpp
@@ -36,13 +36,12 @@
     CharUnits AtomicAlign;
     CharUnits ValueAlign;
     TypeEvaluationKind EvaluationKind;
-    bool UseLibcall;
     LValue LVal;
     CGBitFieldInfo BFI;
   public:
     AtomicInfo(CodeGenFunction &CGF, LValue &lvalue)
         : CGF(CGF), AtomicSizeInBits(0), ValueSizeInBits(0),
-          EvaluationKind(TEK_Scalar), UseLibcall(true) {
+          EvaluationKind(TEK_Scalar) {
       assert(!lvalue.isGlobalReg());
       ASTContext &C = CGF.getContext();
       if (lvalue.isSimple()) {
@@ -126,8 +125,6 @@
         AtomicAlign = ValueAlign = lvalue.getAlignment();
         LVal = lvalue;
       }
-      UseLibcall = !C.getTargetInfo().hasBuiltinAtomic(
-          AtomicSizeInBits, C.toBits(lvalue.getAlignment()));
     }
 
     QualType getAtomicType() const { return AtomicTy; }
@@ -136,7 +133,17 @@
     uint64_t getAtomicSizeInBits() const { return AtomicSizeInBits; }
     uint64_t getValueSizeInBits() const { return ValueSizeInBits; }
     TypeEvaluationKind getEvaluationKind() const { return EvaluationKind; }
-    bool shouldUseLibcall() const { return UseLibcall; }
+    bool shouldUseLibcall(TargetInfo::AtomicOperationKind Op) const {
+      const llvm::fltSemantics &FS =
+          ValueTy->isRealFloatingType()
+              ? CGF.getContext().getFloatTypeSemantics(ValueTy)
+              : llvm::APFloat::Bogus();
+      auto Support = CGF.getContext().getTargetInfo().getAtomicSupport(
+          Op, AtomicSizeInBits, CGF.getContext().toBits(LVal.getAlignment()),
+          FS);
+      assert(Support != TargetInfo::AtomicSupportKind::Unsupported);
+      return Support == TargetInfo::AtomicSupportKind::Library;
+    }
     const LValue &getAtomicLValue() const { return LVal; }
     llvm::Value *getAtomicPointer() const {
       if (LVal.isSimple())
@@ -602,21 +609,25 @@
     break;
 
   case AtomicExpr::AO__atomic_add_fetch:
-    PostOp = llvm::Instruction::Add;
+    PostOp = E->getValueType()->isFloatingType() ? llvm::Instruction::FAdd
+                                                 : llvm::Instruction::Add;
     LLVM_FALLTHROUGH;
   case AtomicExpr::AO__c11_atomic_fetch_add:
   case AtomicExpr::AO__opencl_atomic_fetch_add:
   case AtomicExpr::AO__atomic_fetch_add:
-    Op = llvm::AtomicRMWInst::Add;
+    Op = E->getValueType()->isFloatingType() ? llvm::AtomicRMWInst::FAdd
+                                             : llvm::AtomicRMWInst::Add;
     break;
 
   case AtomicExpr::AO__atomic_sub_fetch:
-    PostOp = llvm::Instruction::Sub;
+    PostOp = E->getValueType()->isFloatingType() ? llvm::Instruction::FSub
+                                                 : llvm::Instruction::Sub;
     LLVM_FALLTHROUGH;
   case AtomicExpr::AO__c11_atomic_fetch_sub:
   case AtomicExpr::AO__opencl_atomic_fetch_sub:
   case AtomicExpr::AO__atomic_fetch_sub:
-    Op = llvm::AtomicRMWInst::Sub;
+    Op = E->getValueType()->isFloatingType() ? llvm::AtomicRMWInst::FSub
+                                             : llvm::AtomicRMWInst::Sub;
     break;
 
   case AtomicExpr::AO__atomic_min_fetch:
@@ -813,6 +824,8 @@
   bool Oversized = getContext().toBits(TInfo.Width) > MaxInlineWidthInBits;
   bool Misaligned = (Ptr.getAlignment() % TInfo.Width) != 0;
   bool UseLibcall = Misaligned | Oversized;
+  bool ShouldCastToIntPtrTy = true;
+
   CharUnits MaxInlineWidth =
       getContext().toCharUnitsFromBits(MaxInlineWidthInBits);
 
@@ -892,11 +905,14 @@
       EmitStoreOfScalar(Val1Scalar, MakeAddrLValue(Temp, Val1Ty));
       break;
     }
-      LLVM_FALLTHROUGH;
+    LLVM_FALLTHROUGH;
   case AtomicExpr::AO__atomic_fetch_add:
   case AtomicExpr::AO__atomic_fetch_sub:
   case AtomicExpr::AO__atomic_add_fetch:
   case AtomicExpr::AO__atomic_sub_fetch:
+    ShouldCastToIntPtrTy = !MemTy->isFloatingType();
+    LLVM_FALLTHROUGH;
+
   case AtomicExpr::AO__c11_atomic_store:
   case AtomicExpr::AO__c11_atomic_exchange:
   case AtomicExpr::AO__opencl_atomic_store:
@@ -937,15 +953,23 @@
   LValue AtomicVal = MakeAddrLValue(Ptr, AtomicTy);
   AtomicInfo Atomics(*this, AtomicVal);
 
-  Ptr = Atomics.emitCastToAtomicIntPointer(Ptr);
-  if (Val1.isValid()) Val1 = Atomics.convertToAtomicIntPointer(Val1);
-  if (Val2.isValid()) Val2 = Atomics.convertToAtomicIntPointer(Val2);
-  if (Dest.isValid())
-    Dest = Atomics.emitCastToAtomicIntPointer(Dest);
-  else if (E->isCmpXChg())
+  if (ShouldCastToIntPtrTy) {
+    Ptr = Atomics.emitCastToAtomicIntPointer(Ptr);
+    if (Val1.isValid())
+      Val1 = Atomics.convertToAtomicIntPointer(Val1);
+    if (Val2.isValid())
+      Val2 = Atomics.convertToAtomicIntPointer(Val2);
+  }
+  if (Dest.isValid()) {
+    if (ShouldCastToIntPtrTy)
+      Dest = Atomics.emitCastToAtomicIntPointer(Dest);
+  } else if (E->isCmpXChg())
     Dest = CreateMemTemp(RValTy, "cmpxchg.bool");
-  else if (!RValTy->isVoidType())
-    Dest = Atomics.emitCastToAtomicIntPointer(Atomics.CreateTempAlloca());
+  else if (!RValTy->isVoidType()) {
+    Dest = Atomics.CreateTempAlloca();
+    if (ShouldCastToIntPtrTy)
+      Dest = Atomics.emitCastToAtomicIntPointer(Dest);
+  }
 
   // Use a library call.  See: http://gcc.gnu.org/wiki/Atomic/GCCMM/LIbrary .
   if (UseLibcall) {
@@ -1536,7 +1560,8 @@
   AtomicInfo AI(*this, LV);
   bool IsVolatile = LV.isVolatile() || hasVolatileMember(LV.getType());
   // An atomic is inline if we don't need to use a libcall.
-  bool AtomicIsInline = !AI.shouldUseLibcall();
+  bool AtomicIsInline =
+      !AI.shouldUseLibcall(TargetInfo::AtomicOperationKind::LoadStore);
   // MSVC doesn't seem to do this for types wider than a pointer.
   if (getContext().getTypeSize(LV.getType()) >
       getContext().getTypeSize(getContext().getIntPtrType()))
@@ -1561,7 +1586,7 @@
                                   bool AsValue, llvm::AtomicOrdering AO,
                                   bool IsVolatile) {
   // Check whether we should use a library call.
-  if (shouldUseLibcall()) {
+  if (shouldUseLibcall(TargetInfo::AtomicOperationKind::LoadStore)) {
     Address TempAddr = Address::invalid();
     if (LVal.isSimple() && !ResultSlot.isIgnored()) {
       assert(getEvaluationKind() == TEK_Aggregate);
@@ -1728,7 +1753,7 @@
     Failure = llvm::AtomicCmpXchgInst::getStrongestFailureOrdering(Success);
 
   // Check whether we should use a library call.
-  if (shouldUseLibcall()) {
+  if (shouldUseLibcall(TargetInfo::AtomicOperationKind::CmpXchg)) {
     // Produce a source address.
     Address ExpectedAddr = materializeRValue(Expected);
     Address DesiredAddr = materializeRValue(Desired);
@@ -1952,7 +1977,7 @@
 void AtomicInfo::EmitAtomicUpdate(
     llvm::AtomicOrdering AO, const llvm::function_ref<RValue(RValue)> &UpdateOp,
     bool IsVolatile) {
-  if (shouldUseLibcall()) {
+  if (shouldUseLibcall(TargetInfo::AtomicOperationKind::LoadStore)) {
     EmitAtomicUpdateLibcall(AO, UpdateOp, IsVolatile);
   } else {
     EmitAtomicUpdateOp(AO, UpdateOp, IsVolatile);
@@ -1961,7 +1986,7 @@
 
 void AtomicInfo::EmitAtomicUpdate(llvm::AtomicOrdering AO, RValue UpdateRVal,
                                   bool IsVolatile) {
-  if (shouldUseLibcall()) {
+  if (shouldUseLibcall(TargetInfo::AtomicOperationKind::LoadStore)) {
     EmitAtomicUpdateLibcall(AO, UpdateRVal, IsVolatile);
   } else {
     EmitAtomicUpdateOp(AO, UpdateRVal, IsVolatile);
@@ -2006,7 +2031,7 @@
     }
 
     // Check whether we should use a library call.
-    if (atomics.shouldUseLibcall()) {
+    if (atomics.shouldUseLibcall(TargetInfo::AtomicOperationKind::LoadStore)) {
       // Produce a source address.
       Address srcAddr = atomics.materializeRValue(rvalue);
 
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -5023,12 +5023,64 @@
   }
 }
 
+static Optional<TargetInfo::AtomicOperationKind>
+getTargetAtomicOp(BinaryOperatorKind BO) {
+  switch (BO) {
+  case BO_Add:
+  case BO_Sub:
+    return TargetInfo::AtomicOperationKind::AddSub;
+  case BO_And:
+  case BO_Or:
+  case BO_Xor:
+    return TargetInfo::AtomicOperationKind::LogicOp;
+  case BO_LT:
+  case BO_GT:
+    return TargetInfo::AtomicOperationKind::MinMax;
+  case BO_Assign:
+    return TargetInfo::AtomicOperationKind::Xchg;
+  case BO_Mul:
+  case BO_Div:
+  case BO_Rem:
+  case BO_Shl:
+  case BO_Shr:
+  case BO_LAnd:
+  case BO_LOr:
+  case BO_PtrMemD:
+  case BO_PtrMemI:
+  case BO_LE:
+  case BO_GE:
+  case BO_EQ:
+  case BO_NE:
+  case BO_Cmp:
+  case BO_AddAssign:
+  case BO_SubAssign:
+  case BO_AndAssign:
+  case BO_OrAssign:
+  case BO_XorAssign:
+  case BO_MulAssign:
+  case BO_DivAssign:
+  case BO_RemAssign:
+  case BO_ShlAssign:
+  case BO_ShrAssign:
+  case BO_Comma:
+    return None;
+  }
+}
+
 static std::pair<bool, RValue> emitOMPAtomicRMW(CodeGenFunction &CGF, LValue X,
                                                 RValue Update,
                                                 BinaryOperatorKind BO,
                                                 llvm::AtomicOrdering AO,
                                                 bool IsXLHSInRHSPart) {
   ASTContext &Context = CGF.getContext();
+  auto getAtomicSupport = [&](BinaryOperatorKind BO, LValue X) {
+    auto K = getTargetAtomicOp(BO);
+    if (!K.hasValue())
+      return TargetInfo::AtomicSupportKind::Unsupported;
+    return Context.getTargetInfo().getAtomicSupport(
+        K.getValue(), Context.getTypeSize(X.getType()),
+        Context.toBits(X.getAlignment()));
+  };
   // Allow atomicrmw only if 'x' and 'update' are integer values, lvalue for 'x'
   // expression is simple and atomic is allowed for the given type for the
   // target platform.
@@ -5038,8 +5090,7 @@
        (Update.getScalarVal()->getType() !=
         X.getAddress(CGF).getElementType())) ||
       !X.getAddress(CGF).getElementType()->isIntegerTy() ||
-      !Context.getTargetInfo().hasBuiltinAtomic(
-          Context.getTypeSize(X.getType()), Context.toBits(X.getAlignment())))
+      getAtomicSupport(BO, X) != TargetInfo::AtomicSupportKind::LockFree)
     return std::make_pair(false, RValue::get(nullptr));
 
   llvm::AtomicRMWInst::BinOp RMWOp;
Index: clang/lib/Sema/SemaChecking.cpp
===================================================================
--- clang/lib/Sema/SemaChecking.cpp
+++ clang/lib/Sema/SemaChecking.cpp
@@ -101,6 +101,8 @@
 
 using namespace clang;
 using namespace sema;
+using AtomicSupportKind = TargetInfo::AtomicSupportKind;
+using TargetAtomicOpKind = TargetInfo::AtomicOperationKind;
 
 SourceLocation Sema::getLocationOfStringLiteralByte(const StringLiteral *SL,
                                                     unsigned ByteNo) const {
@@ -4794,7 +4796,8 @@
   case AtomicExpr::AO__atomic_add_fetch:
   case AtomicExpr::AO__atomic_sub_fetch:
     IsAddSub = true;
-    LLVM_FALLTHROUGH;
+    Form = Arithmetic;
+    break;
   case AtomicExpr::AO__c11_atomic_fetch_and:
   case AtomicExpr::AO__c11_atomic_fetch_or:
   case AtomicExpr::AO__c11_atomic_fetch_xor:
@@ -4809,6 +4812,8 @@
   case AtomicExpr::AO__atomic_or_fetch:
   case AtomicExpr::AO__atomic_xor_fetch:
   case AtomicExpr::AO__atomic_nand_fetch:
+    Form = Arithmetic;
+    break;
   case AtomicExpr::AO__c11_atomic_fetch_min:
   case AtomicExpr::AO__c11_atomic_fetch_max:
   case AtomicExpr::AO__opencl_atomic_fetch_min:
@@ -4902,9 +4907,9 @@
   // For an arithmetic operation, the implied arithmetic must be well-formed.
   if (Form == Arithmetic) {
     // gcc does not enforce these rules for GNU atomics, but we do so for sanity.
-    if (IsAddSub && !ValType->isIntegerType()
-        && !ValType->isPointerType()) {
-      Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int_or_ptr)
+    if (IsAddSub && !ValType->isIntegerType() && !ValType->isPointerType() &&
+        !ValType->isFloatingType()) {
+      Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int_ptr_or_fp)
           << IsC11 << Ptr->getType() << Ptr->getSourceRange();
       return ExprError();
     }
@@ -5031,7 +5036,9 @@
         // passed by address. For the rest, GNU uses by-address and C11 uses
         // by-value.
         assert(Form != Load);
-        if (Form == Init || (Form == Arithmetic && ValType->isIntegerType()))
+        if (Form == Arithmetic && ValType->isPointerType())
+          Ty = Context.getPointerDiffType();
+        else if (Form == Init || Form == Arithmetic)
           Ty = ValType;
         else if (Form == Copy || Form == Xchg) {
           if (IsPassedByAddress) {
@@ -5040,9 +5047,7 @@
                                  ExprRange.getBegin());
           }
           Ty = ByValType;
-        } else if (Form == Arithmetic)
-          Ty = Context.getPointerDiffType();
-        else {
+        } else {
           Expr *ValArg = APIOrderedArgs[i];
           // The value pointer is always dereferenced, a nullptr is undefined.
           CheckNonNullArgument(*this, ValArg, ExprRange.getBegin());
@@ -5144,16 +5149,13 @@
   AtomicExpr *AE = new (Context)
       AtomicExpr(ExprRange.getBegin(), SubExprs, ResultType, Op, RParenLoc);
 
-  if ((Op == AtomicExpr::AO__c11_atomic_load ||
-       Op == AtomicExpr::AO__c11_atomic_store ||
-       Op == AtomicExpr::AO__opencl_atomic_load ||
-       Op == AtomicExpr::AO__opencl_atomic_store ) &&
-      Context.AtomicUsesUnsupportedLibcall(AE))
-    Diag(AE->getBeginLoc(), diag::err_atomic_load_store_uses_lib)
-        << ((Op == AtomicExpr::AO__c11_atomic_load ||
-             Op == AtomicExpr::AO__opencl_atomic_load)
-                ? 0
-                : 1);
+  auto TargetOp = AE->getTargetAtomicOp();
+  if (Context.getTargetAtomicSupport(TargetOp, AtomTy) ==
+      TargetInfo::AtomicSupportKind::Unsupported) {
+    Diag(AE->getBeginLoc(), diag::err_atomic_op_unsupported)
+        << (unsigned)TargetOp << AtomTy << Ptr->getSourceRange();
+    return ExprError();
+  }
 
   if (ValType->isExtIntType()) {
     Diag(Ptr->getExprLoc(), diag::err_atomic_builtin_ext_int_prohibit);
Index: clang/test/CodeGen/fp-atomic-ops.c
===================================================================
--- /dev/null
+++ clang/test/CodeGen/fp-atomic-ops.c
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=amdgcn-amd-amdhsa \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
+
+// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=aarch64-linux-gnu \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
+
+// RUN: %clang_cc1 %s -emit-llvm -O0 -o - -triple=armv8-apple-ios7.0 \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT %s
+
+// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=hexagon \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
+
+// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=mips64-mti-linux-gnu \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
+
+// RUN: %clang_cc1 %s -emit-llvm -O0 -o - -triple=i686-linux-gnu \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT %s
+
+// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=x86_64-linux-gnu \
+// RUN:   | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
+
+typedef enum memory_order {
+  memory_order_relaxed = __ATOMIC_RELAXED,
+  memory_order_acquire = __ATOMIC_ACQUIRE,
+  memory_order_release = __ATOMIC_RELEASE,
+  memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  memory_order_seq_cst = __ATOMIC_SEQ_CST
+} memory_order;
+
+void test(float *f, float ff, double *d, double dd) {
+  // FLOAT: atomicrmw fadd float* {{.*}} monotonic
+  __atomic_fetch_add(f, ff, memory_order_relaxed);
+
+  // FLOAT: atomicrmw fsub float* {{.*}} monotonic
+  __atomic_fetch_sub(f, ff, memory_order_relaxed);
+
+#ifdef DOUBLE
+  // DOUBLE: atomicrmw fadd double* {{.*}} monotonic
+  __atomic_fetch_add(d, dd, memory_order_relaxed);
+
+  // DOUBLE: atomicrmw fsub double* {{.*}} monotonic
+  __atomic_fetch_sub(d, dd, memory_order_relaxed);
+#endif
+}
Index: clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
+// RUN:   -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
+// RUN:   -fnative-half-arguments-and-returns | FileCheck %s
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+__device__ float ffp1(float *p) {
+  // CHECK-LABEL: @_Z4ffp1Pf
+  // CHECK: atomicrmw fadd float* {{.*}} monotonic
+  return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+}
+
+__device__ double ffp2(double *p) {
+  // CHECK-LABEL: @_Z4ffp2Pd
+  // CHECK: atomicrmw fsub double* {{.*}} monotonic
+  return __atomic_fetch_sub(p, 1.0, memory_order_relaxed);
+}
+
+// long double is the same as double for amdgcn.
+__device__ long double ffp3(long double *p) {
+  // CHECK-LABEL: @_Z4ffp3Pe
+  // CHECK: atomicrmw fsub double* {{.*}} monotonic
+  return __atomic_fetch_sub(p, 1.0L, memory_order_relaxed);
+}
+
+__device__ double ffp4(double *p, float f) {
+  // CHECK-LABEL: @_Z4ffp4Pdf
+  // CHECK: fpext float {{.*}} to double
+  // CHECK: atomicrmw fsub double* {{.*}} monotonic
+  return __atomic_fetch_sub(p, f, memory_order_relaxed);
+}
+
+__device__ double ffp5(double *p, int i) {
+  // CHECK-LABEL: @_Z4ffp5Pdi
+  // CHECK: sitofp i32 {{.*}} to double
+  // CHECK: atomicrmw fsub double* {{.*}} monotonic
+  return __atomic_fetch_sub(p, i, memory_order_relaxed);
+}
Index: clang/test/CodeGenOpenCL/atomic-ops.cl
===================================================================
--- clang/test/CodeGenOpenCL/atomic-ops.cl
+++ clang/test/CodeGenOpenCL/atomic-ops.cl
@@ -1,12 +1,17 @@
-// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -O0 -o - -triple=amdgcn-amd-amdhsa-amdgizcl | opt -instnamer -S | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -O0 -o - -triple=amdgcn-amd-amdhsa \
+// RUN:   | opt -instnamer -S | FileCheck %s
 
 // Also test serialization of atomic operations here, to avoid duplicating the test.
-// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-pch -O0 -o %t -triple=amdgcn-amd-amdhsa-amdgizcl
-// RUN: %clang_cc1 %s -cl-std=CL2.0 -include-pch %t -O0 -triple=amdgcn-amd-amdhsa-amdgizcl -emit-llvm -o - | opt -instnamer -S | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-pch -O0 -o %t -triple=amdgcn-amd-amdhsa
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -include-pch %t -O0 -triple=amdgcn-amd-amdhsa \
+// RUN:   -emit-llvm -o - | opt -instnamer -S | FileCheck %s
 
 #ifndef ALREADY_INCLUDED
 #define ALREADY_INCLUDED
 
+#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
+
 typedef __INTPTR_TYPE__ intptr_t;
 typedef int int8 __attribute__((ext_vector_type(8)));
 
@@ -185,6 +190,18 @@
   return __opencl_atomic_exchange(d, 2, memory_order_seq_cst, memory_scope_work_group);
 }
 
+float ff4(global atomic_float *d, float a) {
+  // CHECK-LABEL: @ff4
+  // CHECK: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
+  return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
+}
+
+float ff5(global atomic_double *d, double a) {
+  // CHECK-LABEL: @ff5
+  // CHECK: atomicrmw fadd double addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
+  return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
+}
+
 // CHECK-LABEL: @atomic_init_foo
 void atomic_init_foo()
 {
Index: clang/test/Sema/atomic-ops.c
===================================================================
--- clang/test/Sema/atomic-ops.c
+++ clang/test/Sema/atomic-ops.c
@@ -99,7 +99,8 @@
 #define _AS2 __attribute__((address_space(2)))
 
 void f(_Atomic(int) *i, const _Atomic(int) *ci,
-       _Atomic(int*) *p, _Atomic(float) *d,
+       _Atomic(int*) *p, _Atomic(float) *f, _Atomic(double) *d,
+       _Atomic(long double) *ld,
        int *I, const int *CI,
        int **P, float *D, struct S *s1, struct S *s2) {
   __c11_atomic_init(I, 5); // expected-error {{pointer to _Atomic}}
@@ -114,7 +115,7 @@
 
   __c11_atomic_load(i, memory_order_seq_cst);
   __c11_atomic_load(p, memory_order_seq_cst);
-  __c11_atomic_load(d, memory_order_seq_cst);
+  __c11_atomic_load(f, memory_order_seq_cst);
   __c11_atomic_load(ci, memory_order_seq_cst);
 
   int load_n_1 = __atomic_load_n(I, memory_order_relaxed);
@@ -137,7 +138,7 @@
 
   __c11_atomic_store(i, 1, memory_order_seq_cst);
   __c11_atomic_store(p, 1, memory_order_seq_cst); // expected-warning {{incompatible integer to pointer conversion}}
-  (int)__c11_atomic_store(d, 1, memory_order_seq_cst); // expected-error {{operand of type 'void'}}
+  (int)__c11_atomic_store(f, 1, memory_order_seq_cst); // expected-error {{operand of type 'void'}}
 
   __atomic_store_n(I, 4, memory_order_release);
   __atomic_store_n(I, 4.0, memory_order_release);
@@ -166,20 +167,22 @@
 
   __c11_atomic_fetch_add(i, 1, memory_order_seq_cst);
   __c11_atomic_fetch_add(p, 1, memory_order_seq_cst);
-  __c11_atomic_fetch_add(d, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer or pointer}}
+  __c11_atomic_fetch_add(f, 1.0f, memory_order_seq_cst);
+  __c11_atomic_fetch_add(d, 1.0, memory_order_seq_cst);
+  __c11_atomic_fetch_add(ld, 1.0, memory_order_seq_cst); // expected-error {{atomic add/sub of '_Atomic(long double)' type requires runtime support that is not available for this target}}
 
-  __atomic_fetch_add(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer or pointer}}
+  __atomic_fetch_add(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer, pointer or supported floating point type}}
   __atomic_fetch_sub(I, 3, memory_order_seq_cst);
   __atomic_fetch_sub(P, 3, memory_order_seq_cst);
-  __atomic_fetch_sub(D, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer or pointer}}
-  __atomic_fetch_sub(s1, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer or pointer}}
+  __atomic_fetch_sub(D, 3, memory_order_seq_cst);
+  __atomic_fetch_sub(s1, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer, pointer or supported floating point type}}
   __atomic_fetch_min(D, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
   __atomic_fetch_max(P, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
   __atomic_fetch_max(p, 3);                       // expected-error {{too few arguments to function call, expected 3, have 2}}
 
   __c11_atomic_fetch_and(i, 1, memory_order_seq_cst);
   __c11_atomic_fetch_and(p, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer}}
-  __c11_atomic_fetch_and(d, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer}}
+  __c11_atomic_fetch_and(f, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer}}
 
   __atomic_fetch_and(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer}}
   __atomic_fetch_or(I, 3, memory_order_seq_cst);
@@ -189,12 +192,12 @@
 
   _Bool cmpexch_1 = __c11_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst);
   _Bool cmpexch_2 = __c11_atomic_compare_exchange_strong(p, P, (int*)1, memory_order_seq_cst, memory_order_seq_cst);
-  _Bool cmpexch_3 = __c11_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
+  _Bool cmpexch_3 = __c11_atomic_compare_exchange_strong(f, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
   (void)__c11_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
 
   _Bool cmpexchw_1 = __c11_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst);
   _Bool cmpexchw_2 = __c11_atomic_compare_exchange_weak(p, P, (int*)1, memory_order_seq_cst, memory_order_seq_cst);
-  _Bool cmpexchw_3 = __c11_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
+  _Bool cmpexchw_3 = __c11_atomic_compare_exchange_weak(f, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
   (void)__c11_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
 
   _Bool cmpexch_4 = __atomic_compare_exchange_n(I, I, 5, 1, memory_order_seq_cst, memory_order_seq_cst);
Index: clang/test/Sema/atomic-requires-library-error.c
===================================================================
--- clang/test/Sema/atomic-requires-library-error.c
+++ clang/test/Sema/atomic-requires-library-error.c
@@ -14,7 +14,7 @@
 
 void structAtomicStore() {
   struct foo f = {0};
-  __c11_atomic_store(&bigAtomic, f, 5); // expected-error {{atomic store requires runtime support that is not available for this target}}
+  __c11_atomic_store(&bigAtomic, f, 5); // expected-error {{atomic C11 load/store of '_Atomic(struct foo)' type requires runtime support that is not available for this target}}
 
   struct bar b = {0};
   __atomic_store(&smallThing, &b, 5);
@@ -23,7 +23,7 @@
 }
 
 void structAtomicLoad() {
-  struct foo f = __c11_atomic_load(&bigAtomic, 5); // expected-error {{atomic load requires runtime support that is not available for this target}}
+  struct foo f = __c11_atomic_load(&bigAtomic, 5); // expected-error {{atomic C11 load/store of '_Atomic(struct foo)' type requires runtime support that is not available for this target}}
   struct bar b;
   __atomic_load(&smallThing, &b, 5);
 
Index: clang/test/SemaCUDA/amdgpu-atomic-ops.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/amdgpu-atomic-ops.cu
@@ -0,0 +1,27 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 %s -verify -fsyntax-only -triple=amdgcn-amd-amdhsa \
+// RUN:   -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
+// RUN:   -fnative-half-arguments-and-returns
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+__device__ _Float16 test_Flot16(_Float16 *p) {
+  return __atomic_fetch_sub(p, 1.0f16, memory_order_relaxed);
+  // expected-error@-1 {{atomic add/sub of '_Float16' type requires runtime support that is not available for this target}}
+}
+
+__device__ __fp16 test_fp16(__fp16 *p) {
+  return __atomic_fetch_sub(p, 1.0f16, memory_order_relaxed);
+  // expected-error@-1 {{atomic add/sub of '__fp16' type requires runtime support that is not available for this target}}
+}
+
+struct BigStruct {
+  int data[128];
+};
+
+void test_big(BigStruct *p1, BigStruct *p2) {
+  __atomic_load(p1, p2, memory_order_relaxed);
+  // expected-error@-1 {{atomic load/store of 'BigStruct' type requires runtime support that is not available for this target}}
+}
Index: clang/test/SemaOpenCL/atomic-ops.cl
===================================================================
--- clang/test/SemaOpenCL/atomic-ops.cl
+++ clang/test/SemaOpenCL/atomic-ops.cl
@@ -1,10 +1,13 @@
-// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only -triple=spir64
-// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only -triple=amdgcn-amdhsa-amd-opencl
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify=expected,spir \
+// RUN:   -fsyntax-only -triple=spir64
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only \
+// RUN:   -triple=amdgcn-amd-amdhsa
 
 // Basic parsing/Sema tests for __opencl_atomic_*
 
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 typedef __INTPTR_TYPE__ intptr_t;
 typedef int int8 __attribute__((ext_vector_type(8)));
@@ -36,7 +39,7 @@
 
 atomic_int gn;
 void f(atomic_int *i, const atomic_int *ci,
-       atomic_intptr_t *p, atomic_float *d,
+       atomic_intptr_t *p, atomic_float *f, atomic_double *d, atomic_half *h, // expected-error {{unknown type name 'atomic_half'}}
        int *I, const int *CI,
        intptr_t *P, float *D, struct S *s1, struct S *s2,
        global atomic_int *i_g, local atomic_int *i_l, private atomic_int *i_p,
@@ -57,37 +60,38 @@
 
   __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(p, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_load(d, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_load(f, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(i_c, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-constant _Atomic type ('__constant atomic_int *' (aka '__constant _Atomic(int) *') invalid)}}
 
   __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_store(p, 1, memory_order_seq_cst, memory_scope_work_group);
-  (int)__opencl_atomic_store(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{operand of type 'void' where arithmetic or pointer type is required}}
+  (int)__opencl_atomic_store(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{operand of type 'void' where arithmetic or pointer type is required}}
 
   int exchange_1 = __opencl_atomic_exchange(i, 1, memory_order_seq_cst, memory_scope_work_group);
   int exchange_2 = __opencl_atomic_exchange(I, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to _Atomic}}
 
   __opencl_atomic_fetch_add(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_add(p, 1, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_add(f, 1.0f, memory_order_seq_cst, memory_scope_work_group); // spir-error {{atomic add/sub of '__generic atomic_float' (aka '__generic _Atomic(float)') type requires runtime support that is not available for this target}}
+  __opencl_atomic_fetch_add(d, 1.0, memory_order_seq_cst, memory_scope_work_group); // spir-error {{atomic add/sub of '__generic atomic_double' (aka '__generic _Atomic(double)') type requires runtime support that is not available for this target}}
   __opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_and(p, 1, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_and(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
 
   __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
-  __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_min(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_max(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
 
   bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
-  bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
+  bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
   (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}}
 
   bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
-  bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
+  bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
   (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}}
 
   // Pointers to different address spaces are allowed.