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.