Let clang atomic builtins fetch add/sub support floating point types

[yaxunl]

Summary

Recently atomicrmw started to support fadd/fsub:

https://reviews.llvm.org/D53965

However clang atomic builtins fetch add/sub still does not support emitting atomicrmw fadd/fsub.

This patch adds that.

[jfb]

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

Also, @ldionne / @EricWF / @mclow.lists do you need this in libc++ for floating-point atomic support?

[yaxunl]

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

For x86_64, amdgcn, aarch64, armv7, mips64, it is translated to cmpxchg by AtomicExpandPass and backends did codegen successfully.

For hexagon, riscv32, it is translated to call of __atomic_fetch_add_4 for fadd float. This is concerning. Probably we need to add __atomic_fetch_{add|sub}_{f16|f32|f64} ?

[arsenm]

clang/lib/CodeGen/CGAtomic.cpp

612–614

Should this really be based on the type, or should the builtin name be different for FP?

[__simt__]

In D71726#1792852, @yaxunl wrote:

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

For x86_64, amdgcn, aarch64, armv7, mips64, it is translated to cmpxchg by AtomicExpandPass and backends did codegen successfully.

For hexagon, riscv32, it is translated to call of __atomic_fetch_add_4 for fadd float. This is concerning. Probably we need to add __atomic_fetch_{add|sub}_{f16|f32|f64} ?

For systems that have load-link/store-conditional architectures, like ARM / PPC / base RISC-V without extension, I would imagine that using a cmpxchg loop is much worse than simply doing the floating-point add/sub in the middle of the atomic mini-transaction. I'm sure that we want back-ends to be capable of implementing this better than what this pass is doing, even when they don't have "native" fp atomics.

You listed amdgcn... what does this do on nvptx?

[arsenm]

In D71726#1801346, @__simt__ wrote:

In D71726#1792852, @yaxunl wrote:

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

For x86_64, amdgcn, aarch64, armv7, mips64, it is translated to cmpxchg by AtomicExpandPass and backends did codegen successfully.

For hexagon, riscv32, it is translated to call of __atomic_fetch_add_4 for fadd float. This is concerning. Probably we need to add __atomic_fetch_{add|sub}_{f16|f32|f64} ?

For systems that have load-link/store-conditional architectures, like ARM / PPC / base RISC-V without extension, I would imagine that using a cmpxchg loop is much worse than simply doing the floating-point add/sub in the middle of the atomic mini-transaction. I'm sure that we want back-ends to be capable of implementing this better than what this pass is doing, even when they don't have "native" fp atomics.

You listed amdgcn... what does this do on nvptx?

Targets can implement shouldExpandAtomicRMWInIR for the desired behavior, which NVPTX currently does not implement. Looking at AtomicExpandPass, it looks like either cmpxchg or LLSC expansions should work for the FP atomics already

[yaxunl]

rebase

[yaxunl]

In D71726#2039319, @arsenm wrote:

In D71726#1801346, @__simt__ wrote:

In D71726#1792852, @yaxunl wrote:

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

For x86_64, amdgcn, aarch64, armv7, mips64, it is translated to cmpxchg by AtomicExpandPass and backends did codegen successfully.

For hexagon, riscv32, it is translated to call of __atomic_fetch_add_4 for fadd float. This is concerning. Probably we need to add __atomic_fetch_{add|sub}_{f16|f32|f64} ?

For systems that have load-link/store-conditional architectures, like ARM / PPC / base RISC-V without extension, I would imagine that using a cmpxchg loop is much worse than simply doing the floating-point add/sub in the middle of the atomic mini-transaction. I'm sure that we want back-ends to be capable of implementing this better than what this pass is doing, even when they don't have "native" fp atomics.

You listed amdgcn... what does this do on nvptx?

Targets can implement shouldExpandAtomicRMWInIR for the desired behavior, which NVPTX currently does not implement. Looking at AtomicExpandPass, it looks like either cmpxchg or LLSC expansions should work for the FP atomics already

nvptx is similar to hexagon and riscv32, where fp atomics is translated to call of __atomic_fetch_add_4.

Since currently only amdgcn supports fp atomics, I am going to add a TargetInfo hook about whether fp atomics is supported and only emit fp atomics for targets supporting it.

clang/lib/CodeGen/CGAtomic.cpp

612–614

I think the original name is better. They are exactly what they are intended to be. They were not able to handle fp types therefore they used to emit diagnostics when fp types were passed to them. However now they are able to handle fp types.

[ldionne]

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

Also, @ldionne / @EricWF / @mclow.lists do you need this in libc++ for floating-point atomic support?

Yes, I guess we do in order to implement fetch_add & friends on floating point types (https://wg21.link/P0020R6).

The builtins would need to work on float, double and long double. The code seems to suggest it does, however the tests only check for float. Does this support __atomic_fetch_{add,sub} on double and long double?

[ldionne]

clang/test/CodeGen/atomic-ops.c

296 ↗

(On Diff #265325)

Sorry if that's a dumb question, but I'm a bit confused: p is a float*, but then we add a double 1.0 to it. Is that intended, or should that be double *p instead (or 1.0f)?

[yaxunl]

In D71726#2047566, @ldionne wrote:

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

Also, @ldionne / @EricWF / @mclow.lists do you need this in libc++ for floating-point atomic support?

Yes, I guess we do in order to implement fetch_add & friends on floating point types (https://wg21.link/P0020R6).

The builtins would need to work on float, double and long double. The code seems to suggest it does, however the tests only check for float. Does this support __atomic_fetch_{add,sub} on double and long double?

It depends on target. For x86_64, __atomic_fetch_{add,sub} on double and long double are translated to __atomic_fetch_sub_8 and __atomic_fetch_sub_16.
For amdgcn, __atomic_fetch_{add,sub} on double is translated to fp atomic insts. long double is the same as double on amdgcn.

clang/test/CodeGen/atomic-ops.c

296 ↗

(On Diff #265325)

In this case, the value type is converted to the pointee type of the pointer operand.

[jfb]

In D71726#2047566, @ldionne wrote:

In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

Also, @ldionne / @EricWF / @mclow.lists do you need this in libc++ for floating-point atomic support?

Yes, I guess we do in order to implement fetch_add & friends on floating point types (https://wg21.link/P0020R6).

The builtins would need to work on float, double and long double. The code seems to suggest it does, however the tests only check for float. Does this support __atomic_fetch_{add,sub} on double and long double?

libc++ could implement atomic<float> using a cmpxchg loop with bit_cast and the FP instruction in most cases, and only use these builtins if available.

[yaxunl]

Added TargetInfo::isFPAtomicFetchAddSubSupported to guard fp atomic.

[tra]

clang/include/clang/Basic/TargetInfo.h

1477

I think it should be predicated on specific type.
E.g. NVPTX supports atomic ops on fp32 ~everywhere, but fp64 atomic add/sub is only supported on newer GPUs.
And then there's fp16...

[ldionne]

clang/test/CodeGen/atomic-ops.c

296 ↗

(On Diff #265325)

Ok, thanks for the clarification. Yeah, it was a dumb question after all. I still think it should be made clearer by using 1.0f.

[yaxunl]

clang/include/clang/Basic/TargetInfo.h

1477

will do and add tests for fp16

clang/test/CodeGen/atomic-ops.c

296 ↗

(On Diff #265325)

this test has been removed. the new tests do not have this issue.

[yaxunl]

check supported fp atomics by bits.

[ldionne]

clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu

27

Nitpick, but this should be 1.0L to be consistent.

[tra]

clang/include/clang/Basic/TargetInfo.h

1477

The number of bits alone may not be sufficient to differentiate the FP variants.
E.g. 16-bit floats currently have 2 variants: IEEE FP16 and BFloat16 (supported by intel and newer NVIDIA GPUs).
CUDA-11 has introduced TF32 FP format, so we're likely to have more than one 32-bit FP type, too.
I think PPC has an odd long double variant represented as pair of 64-bit doubles.

[yaxunl]

clang/include/clang/Basic/TargetInfo.h

1477

will use llvm::fltSemantics for checking, which should cover different fp types.

clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu

27

done

[yaxunl]

use llvm::fltSemantics for checking

[jyknight]

Why not have clang always emit atomicrmw for floats, and let AtomicExpandPass handle legalizing that into integer atomics if necessary, rather than adding a target hook in clang?

[yaxunl]

In D71726#2165424, @jyknight wrote:

Why not have clang always emit atomicrmw for floats, and let AtomicExpandPass handle legalizing that into integer atomics if necessary, rather than adding a target hook in clang?

Not all targets can legalize fp atomics by AtomicExpandPass. Some targets need library support.

[jfb]

In D71726#2165445, @yaxunl wrote:

In D71726#2165424, @jyknight wrote:

Why not have clang always emit atomicrmw for floats, and let AtomicExpandPass handle legalizing that into integer atomics if necessary, rather than adding a target hook in clang?

Not all targets can legalize fp atomics by AtomicExpandPass. Some targets need library support.

What are they missing? It can be expanded to a cmpxchg loop with bitcast to an integer type of the same size.

[jyknight]

In D71726#2165445, @yaxunl wrote:

In D71726#2165424, @jyknight wrote:

Why not have clang always emit atomicrmw for floats, and let AtomicExpandPass handle legalizing that into integer atomics if necessary, rather than adding a target hook in clang?

Not all targets can legalize fp atomics by AtomicExpandPass. Some targets need library support.

That isn't true, because you can do so generically with a cmpxchg loop, assuming that size of atomic is supported by the target. This might not be the most efficient lowering choice, but it's always possible as a fallback. (And if the size is too large, then AtomicExpandPass will lower the cmpxchg to the libatomic call.)

If a target wants to tell AtomicExpandPass that fp add/sub are supported, and then lower the resulting ATOMIC_LOAD_FSUB sdag node into a libcall of its choice, that's also ok (as long as the libcall is lock-free).

[yaxunl]

In D71726#2165494, @jyknight wrote:

In D71726#2165445, @yaxunl wrote:

In D71726#2165424, @jyknight wrote:

Why not have clang always emit atomicrmw for floats, and let AtomicExpandPass handle legalizing that into integer atomics if necessary, rather than adding a target hook in clang?

Not all targets can legalize fp atomics by AtomicExpandPass. Some targets need library support.

That isn't true, because you can do so generically with a cmpxchg loop, assuming that size of atomic is supported by the target. This might not be the most efficient lowering choice, but it's always possible as a fallback. (And if the size is too large, then AtomicExpandPass will lower the cmpxchg to the libatomic call.)

If a target wants to tell AtomicExpandPass that fp add/sub are supported, and then lower the resulting ATOMIC_LOAD_FSUB sdag node into a libcall of its choice, that's also ok (as long as the libcall is lock-free).

how about other fp types e.g. bf16, half, long double? Do we need to diagnose them or not?

[yaxunl]

Make IEEE single and double type as supported for fp atomics in all targets by default. This is based on the assumption that AtomicExpandPass or its ongoing work is sufficient to support fp atomics for all targets. This is to facilitate middle end and backend end development to support fp atomics.

If a target would like to treat single and double fp atomics as unsupported, it can override the default behavior in its own TargetInfo.

[yaxunl]

ping. I think I have addressed all the issues in FE. I think issues in AtomicExpandPass should be addressed by separate patches. Can we land this? Thanks.

[tra]

LGTM, modulo couple of nits.

@jyknight are you OK with this?

In D71726#2179428, @yaxunl wrote:

Make IEEE single and double type as supported for fp atomics in all targets by default. This is based on the assumption that AtomicExpandPass or its ongoing work is sufficient to support fp atomics for all targets. This is to facilitate middle end and backend end development to support fp atomics.

If a target would like to treat single and double fp atomics as unsupported, it can override the default behavior in its own TargetInfo.

Do we have sufficient test coverage on all platforms to make sure we're not generating something that LLVM can't handle everywhere?
If not, perhaps we should default to unsupported and only enable it for known working targets.

clang/lib/CodeGen/CGAtomic.cpp

913–915

ShouldCastToIntPtrTy = !MemTy->isFloatingType();

clang/test/Sema/atomic-ops.c

102–104

Rename arguments? d -> f, d2 -> d, d3 -> ld ?

[yaxunl]

revised by Artem's comments.

[yaxunl]

added tests for targets supporting fp atomics.

[yaxunl]

In D71726#2182667, @tra wrote:

LGTM, modulo couple of nits.

@jyknight are you OK with this?

In D71726#2179428, @yaxunl wrote:

Make IEEE single and double type as supported for fp atomics in all targets by default. This is based on the assumption that AtomicExpandPass or its ongoing work is sufficient to support fp atomics for all targets. This is to facilitate middle end and backend end development to support fp atomics.

If a target would like to treat single and double fp atomics as unsupported, it can override the default behavior in its own TargetInfo.

Do we have sufficient test coverage on all platforms to make sure we're not generating something that LLVM can't handle everywhere?
If not, perhaps we should default to unsupported and only enable it for known working targets.

I updated TargetInfo for fp atomic support for common targets. Basically by default fp atomic support is now off. It is enabled only for targets which do not generate lib calls for fp atomics. This is because the availability of lib call depends on platform, so it is up to the Target owners to determine whether the support is available if lib call is needed. For those targets which are able to generate llvm fp atomic instructions, fp atomic support is enabled in clang, and tests are added to cover them.

clang/lib/CodeGen/CGAtomic.cpp

913–915

done

clang/test/Sema/atomic-ops.c

102–104

done

[yaxunl]

ping

[jyknight]

In D71726#2182667, @tra wrote:

If a target would like to treat single and double fp atomics as unsupported, it can override the default behavior in its own TargetInfo.

I really don't think this should be a target option at all. Every target can support the atomic fadd/fsub IR instruction (via lowering to a cmpxchg loop if nothing else). If it doesn't work, that's a bug in LLVM. We shouldn't be adding target hooks in Clang to workaround LLVM bugs, rather, we should fix them.

There is one nit -- atomicrmw doesn't (yet) support specifying alignment. There's work now to fix that, but until that's submitted, only naturally-aligned atomicrmw instructions can be created. So, for now, supporting only a naturally-aligned floating-point add would be a reasonable temporary measure.

Do we have sufficient test coverage on all platforms to make sure we're not generating something that LLVM can't handle everywhere?

Probably not.

If not, perhaps we should default to unsupported and only enable it for known working targets.

No, I don't think that's a good way to go. We should fix LLVM if it's broken.

[jyknight]

clang/lib/CodeGen/CGAtomic.cpp

959

convertToAtomicIntPointer does more than just cast to an int pointer, are you sure the rest is not necessary for fp types?

clang/lib/Sema/SemaChecking.cpp

5034–5036

This is confusing, and took me a bit to understand what you're doing. I'd suggest reordering the clauses, putting the pointer case first, e.g.:

if (Form == Arithmetic && ValType->isPointerType())
  Ty = Context.getPointerDiffType();
else if (Form == Init || Form == Arithmetic)
  Ty = ValType;
else if (Form == Copy || Form == Xchg) .....
else ......
...

[jyknight]

Oh, one more note, C11 has -- and clang already supports -- _Atomic long double x; x += 4; via lowering to a cmpxchg loop. Now that we have an LLVM IR representation for atomicrmw fadd/fsub, clang should be lowering the _Atomic += to that, too. (Doesn't need to be in this patch, but it should be done.)

[yaxunl]

In D71726#2207148, @jyknight wrote:

In D71726#2182667, @tra wrote:

If a target would like to treat single and double fp atomics as unsupported, it can override the default behavior in its own TargetInfo.

I really don't think this should be a target option at all. Every target can support the atomic fadd/fsub IR instruction (via lowering to a cmpxchg loop if nothing else). If it doesn't work, that's a bug in LLVM. We shouldn't be adding target hooks in Clang to workaround LLVM bugs, rather, we should fix them.

There is one nit -- atomicrmw doesn't (yet) support specifying alignment. There's work now to fix that, but until that's submitted, only naturally-aligned atomicrmw instructions can be created. So, for now, supporting only a naturally-aligned floating-point add would be a reasonable temporary measure.

clang does not always emit atomic instructions for atomic builtins. Clang may emit lib calls for atomic builtins. Basically clang checks target info about max atomic inline width and if the desired atomic operation exceeds the supported atomic inline width, clang will emit lib calls for atomic builtins. The rationale is that the lib calls may be faster than the IR generated by the LLVM pass. This behavior has long existed and it also applies to fp atomics. I don't think emitting lib calls for atomic builtins is a bug. However, this does introduce the issue about whether the library functions for atomics are available for a specific target. As I said, only the target owners have the answer and therefore I introduced the target hook.

Do we have sufficient test coverage on all platforms to make sure we're not generating something that LLVM can't handle everywhere?

Probably not.

In clang, we only test IR generation, as is done for other atomic builtins. fp atomics do not have less coverage compared with other atomic builtins. Actually for other atomic builtins we do not even test them on different targets. The ISA generation of fp atomics should be done in llvm tests and should not be blocking clang change.

clang/lib/CodeGen/CGAtomic.cpp

959

it is not needed for fp types. If the value type does not match the pointer type, clang automatically inserts proper llvm instructions to convert the value type to a value type that matches the pointer type. Two codegen tests are added (atomic_fetch_add(double*, float) and atomic_fetch_add(double*, int)) to test such situations.

clang/lib/Sema/SemaChecking.cpp

5034–5036

done

[yaxunl]

Revised by James' comments.

[yaxunl]

ping

[rjmccall]

In D71726#2207700, @yaxunl wrote:

clang does not always emit atomic instructions for atomic builtins. Clang may emit lib calls for atomic builtins. Basically clang checks target info about max atomic inline width and if the desired atomic operation exceeds the supported atomic inline width, clang will emit lib calls for atomic builtins. The rationale is that the lib calls may be faster than the IR generated by the LLVM pass. This behavior has long existed and it also applies to fp atomics. I don't think emitting lib calls for atomic builtins is a bug. However, this does introduce the issue about whether the library functions for atomics are available for a specific target. As I said, only the target owners have the answer and therefore I introduced the target hook.

If we want the frontend to emit an error when the target doesn't support library-based atomics, that seems fine, but there's no reason to only do so for floating-point types. That is, we should have a TargetInfo method that asks whether atomics at a given size and alignment are supported at all, similar to what we have for "builtin" (lock-free) atomics, and we should check it for all the atomic types and operations.

Actually, maybe we should take the existing hook and have it return one of { LockFree, Library, Unsupported }.

[jyknight]

In D71726#2207700, @yaxunl wrote:

clang does not always emit atomic instructions for atomic builtins. Clang may emit lib calls for atomic builtins. Basically clang checks target info about max atomic inline width and if the desired atomic operation exceeds the supported atomic inline width, clang will emit lib calls for atomic builtins. The rationale is that the lib calls may be faster than the IR generated by the LLVM pass. This behavior has long existed and it also applies to fp atomics. I don't think emitting lib calls for atomic builtins is a bug. However, this does introduce the issue about whether the library functions for atomics are available for a specific target. As I said, only the target owners have the answer and therefore I introduced the target hook.

The LLVM AtomicExpandPass is _also_ introducing libcalls (or cmpxchg loops), as is appropriate for a given target. We currently, redundantly, support the same thing in two places. It's a long-term goal of mine to simplify the atomics code in clang, by deferring more of it to LLVM, but some prerequisites (e.g. supporting misaligned atomicrmw) are not yet in place. The intent is that it is always valid to emit the LLVM atomic IR, and it will be transformed into whatever is best on a given target. As such, there's no reason to restrict these clang intrinsics.

Yes, there are no generically available libcalls for atomic float math -- but that's okay -- let LLVM handle transform into a cmpxchg loop when required.

[rjmccall]

Yes, there are no generically available libcalls for atomic float math -- but that's okay -- let LLVM handle transform into a cmpxchg loop when required.

I suspect Yaxun's target cannot provide libcalls at all, which is why he wants to diagnose up-front. But I agree that we should be thinking about this uniformly, and that his target should be diagnosing *all* unsupported atomics.

[yaxunl]

revised by John's comments. Added target hook and diagnostics for generic atomic operations.

[yaxunl]

In D71726#2351069, @rjmccall wrote:

Yes, there are no generically available libcalls for atomic float math -- but that's okay -- let LLVM handle transform into a cmpxchg loop when required.

I suspect Yaxun's target cannot provide libcalls at all, which is why he wants to diagnose up-front. But I agree that we should be thinking about this uniformly, and that his target should be diagnosing *all* unsupported atomics.

amdgpu target currently does not support atomic libcalls. I added a target hook for atomic operation support and diagnostics for generic atomic operations by John's comments.

Clang has existing diagnostics for unsupported atomic load/store for some platforms, and functions about atomic support scattered in target info, AST context, and codegen. This change refactors these codes and unify them as a target hook.

[yaxunl]

ping

[yaxunl]

ping

[yaxunl]

ping

[yaxunl]

@rjmccall I have addressed the comments about diagnostics. Could you please review it? Thanks.

[rjmccall]

clang/include/clang/Basic/TargetInfo.h

1478

This shouldn't be here; if you have places that don't always represent an atomic operation, queries for the kind should return an Optional<AtomicOperationKind> from the classification.

1479

atomic_init is not actually an atomic operation, so there's never an inherent reason it can't be supported.

In general, I am torn about this list, because it's simultaneously rather fine-grained while not seeming nearly fine-grained enough to be truly general. What's actually going on on your target? You have ISA support for doing some specific operations atomically, but not a general atomic compare-and-swap operation? Which means that you then cannot support support other operations?

It is unfortunate that our layering prevents TargetInfo from simply being passed the appropriate expression.

1497

I think this reflects our current strategies for emitting atomics, but it's a somewhat misleading enum in general because this isn't an exhaustive list of the options — there are certainly possible inline expansions that aren't lock-free. (For example, you could have an inline spin-lock embedded in the atomic object.) The goal of this enum is so that TargetInfo only has to have one hook for checking atomic operations? I would be happier if you included an inline-but-not-lock-free alternative in this enum, even if it's never currently used, so that clients can do the right test.

1501

Why is this needed as a separate hook?

clang/lib/AST/ASTContext.cpp

11046

Should this be a method on AtomicExpr? It seems like an intrinsic, target-independent property of the expression.

clang/lib/Basic/TargetInfo.cpp

870

Darwin targets should all be subclasses of DarwinTargetInfo in OSTargets.h, so you should be able to just override this there instead of having it in the base case.

clang/lib/Basic/Targets/AArch64.h

143

Why can't targets reliably expand this to an atomic compare-and-exchange if they support that for the target width?

[yaxunl]

clang/include/clang/Basic/TargetInfo.h

1478

Removed.

1479

The target hook getAtomicSupport needs an argument for atomic operation. Since not all targets support fp add/sub, we need an enum for add/sub. Since certain release of iOS/macOS does not support C11 load/store, we need an enum for C11 load/store. We could define the enums as {AddSub, C11LoadStore, Other}. However, this would cause a difficulty for emitting diagnostic message for unsupported atomic operations since we map this enum to a string for the atomic operation and use it in the diagnostic message. 'Other' would be mapped to 'other atomic operation' which is not clear what it is.

1497

Added InlineWithLock

1501

Most target shares getAtomicSupport except FP atomic support, so define a virtual function for FP atomic support and let getAtomicSupport call it.

clang/lib/AST/ASTContext.cpp

11046

Yes. moved to AtomicExpr

clang/lib/Basic/TargetInfo.cpp

870

done

clang/lib/Basic/Targets/AArch64.h

143

There are some bugs in either the middle end or backend causing this not working. For example, half type atomic fadd on amdgcn is not lowered to cmpxchg and the backend has isel failure, bf16 type atomic fadd on arm is not lowered to cmpxchg and the backend has isel failure. The support for each fp type needs to be done case by case. So far there is no target support atomic fadd/sub with half and bf16 type.

[yaxunl]

revised by John's comments

[jyknight]

I still have the same fundamental objection as before to the parts of this patch for prohibiting FP add/sub on some targets.

If a particular LLVM target cannot handle transforming an FP add/sub (or any other RMW operations!) into the correct cmpxchg or LL/SC loop, that's a bug in the backend which should be fixed. I don't see why we ought to add a bunch of functionality in the frontend to workaround this?

(Some of the other changes, e.g. to diagnose lack of support for large atomics is useful, though.)

[rjmccall]

clang/include/clang/Basic/TargetInfo.h

1479

It's not obviously true that not all targets support FP add/sub, though. Any target that provides compare-and-swap at the width of an FP type can do an atomic FP add/sub at that width; it might be less efficient than it would be with specific ISA support, but that's true for a lot of atomic operations. Surely it's better to just fix whatever bugs LLVM has with lowering atomic FP add/sub than to add more abstraction to Clang to handle a special case that shouldn't exist.

I don't know what issues Darwin has with C11 load/store; that might be a more compelling reason to have this abstraction, although again it seems strange that we're outlawing a specific operation when in principle we can just emit it less efficiently.

clang/lib/Basic/Targets/AArch64.h

143

Are we legalizing atomicrmw to cmpxchg loops in the backend instead of as LLVM IR pass? That seems like an architectural mistake. Regardless, this bug should just be fixed.

[yaxunl]

This patch focuses on clang work for enabling fp atomics. There is a middle end pass for lowering fp atomics to cmpxchg, however not all targets enable it or enable it properly. From clang point of view, those targets are not ready to say they support fp atomics, therefore it diagnose those situations and let clang fail gracefully instead of crashing with isel failure or missing symbols in linker.

I have limited resources to work on middle end and backend for all targets. If a backend really cares about fp atomics, they should fix the atomic lowering pass then enable fp atomics support in clang Target info.

This patch implements fp atomic support in clang. It does not make things worse in regarding the bugs in middle ends and backends. I think it is not beneficial to blocking this clang change due to middle end and backend issues.

[yaxunl]

If the concern is that diagnose fp atomics as unsupported hinders middle end and backend work for fixing fp atomic issues, how about adding a -fenable-fp-atomics to clang which can override target info about fp atomics support.

[jyknight]

My concern is that this is treating a backend _bug_ as if it were just an optional feature. But it's not the case that it might be reasonable to either implement or not implement this in a backend -- it should be implemented, and those that don't are buggy.

I'd be happier with just having an ISEL failure when you try to use fp atomics on broken targets, rather than adding all this code and configuration to Clang in order to avoid that. (And, of course, the target maintainers should also fix them)

[tra]

In D71726#2536966, @jyknight wrote:

My concern is that this is treating a backend _bug_ as if it were just an optional feature. But it's not the case that it might be reasonable to either implement or not implement this in a backend -- it should be implemented, and those that don't are buggy.

I'd be happier with just having an ISEL failure when you try to use fp atomics on broken targets, rather than adding all this code and configuration to Clang in order to avoid that. (And, of course, the target maintainers should also fix them)

+1. I agree with James.

Removing code is often harder than adding it. When you're adding things, you're the only user. Once things are in, they will start growing dependencies that will need to be dealt with if you ever want to remove the code.

Clean solution that works for AMDGPU only for now is better than a potentially permanent workaround.

[yaxunl]

In D71726#2537054, @tra wrote:

In D71726#2536966, @jyknight wrote:

My concern is that this is treating a backend _bug_ as if it were just an optional feature. But it's not the case that it might be reasonable to either implement or not implement this in a backend -- it should be implemented, and those that don't are buggy.

I'd be happier with just having an ISEL failure when you try to use fp atomics on broken targets, rather than adding all this code and configuration to Clang in order to avoid that. (And, of course, the target maintainers should also fix them)

+1. I agree with James.

Removing code is often harder than adding it. When you're adding things, you're the only user. Once things are in, they will start growing dependencies that will need to be dealt with if you ever want to remove the code.

Clean solution that works for AMDGPU only for now is better than a potentially permanent workaround.

For amdgpu target, we do need diagnose unsupported atomics (not limited to fp atomics) since we do not support libcall due to ISA level linking not supported. This is something we cannot fix in a short time and we would rather diagnose it than confusing the users with missing symbols in lld.

For other targets, I can make changes to assume fp atomics are supported if width is within max inline atomic width of the target. Basically this will let fp atomics emitted for these targets and assuming middle end or backend will handle them properly.

[rjmccall]

In D71726#2537101, @yaxunl wrote:

In D71726#2537054, @tra wrote:

In D71726#2536966, @jyknight wrote:

My concern is that this is treating a backend _bug_ as if it were just an optional feature. But it's not the case that it might be reasonable to either implement or not implement this in a backend -- it should be implemented, and those that don't are buggy.

I'd be happier with just having an ISEL failure when you try to use fp atomics on broken targets, rather than adding all this code and configuration to Clang in order to avoid that. (And, of course, the target maintainers should also fix them)

+1. I agree with James.

Removing code is often harder than adding it. When you're adding things, you're the only user. Once things are in, they will start growing dependencies that will need to be dealt with if you ever want to remove the code.

Clean solution that works for AMDGPU only for now is better than a potentially permanent workaround.

For amdgpu target, we do need diagnose unsupported atomics (not limited to fp atomics) since we do not support libcall due to ISA level linking not supported. This is something we cannot fix in a short time and we would rather diagnose it than confusing the users with missing symbols in lld.

Diagnosing that you don't support atomics your target can't reasonably support is completely fine. (You could actually actually inline a locking approach if you really wanted to, though; Microsoft's std::atomic does that in the general case, although admittedly that's library code.) I would like to understand whether that's really type-specific or just size-specific, though, and I don't think we've gotten a plain answer about that. Is it true that amdgpu simply does not have a generic cmpxchg?

For other targets, I can make changes to assume fp atomics are supported if width is within max inline atomic width of the target. Basically this will let fp atomics emitted for these targets and assuming middle end or backend will handle them properly.

I think that's reasonable.

[jyknight]

In D71726#2537101, @yaxunl wrote:

For amdgpu target, we do need diagnose unsupported atomics (not limited to fp atomics) since we do not support libcall due to ISA level linking not supported. This is something we cannot fix in a short time and we would rather diagnose it than confusing the users with missing symbols in lld.

If this is limited simply to not supporting oversized or misaligned atomics, I'd find that a lot less objectionable. At that point you just need a single boolean variable/accessor for whether the target can support atomic library calls. I note that we already have warning messages: warn_atomic_op_misaligned and warn_atomic_op_oversized. Maybe those can just be promoted to errors on AMDGPU.

[yaxunl]

In D71726#2537378, @jyknight wrote:

In D71726#2537101, @yaxunl wrote:

For amdgpu target, we do need diagnose unsupported atomics (not limited to fp atomics) since we do not support libcall due to ISA level linking not supported. This is something we cannot fix in a short time and we would rather diagnose it than confusing the users with missing symbols in lld.

If this is limited simply to not supporting oversized or misaligned atomics, I'd find that a lot less objectionable. At that point you just need a single boolean variable/accessor for whether the target can support atomic library calls. I note that we already have warning messages: warn_atomic_op_misaligned and warn_atomic_op_oversized. Maybe those can just be promoted to errors on AMDGPU.

Good points. Will do.

[yaxunl]

Revised by James, Artem, and John's comments.

[yaxunl]

ping

[yaxunl]

@jyknight @rjmccall ping. diagnostic issue addressed.

[yaxunl]

@rjmccall @jyknight Ping. Any further concerns? Thanks.

[yaxunl]

Re-use existing warning instead of introducing new diagnostics.

Ping. Can some one help review this patch? I believe all comments addressed. Thanks.

[tra]

@jyknight - James, do you have further concerns about the patch?

clang/lib/Driver/ToolChains/Clang.cpp

6454 ↗

(On Diff #332658)

If we rely on promoting the warnings to errors for correctness, I think we may need a more robust mechanism to enforce that than trying to guess the state based on provided options.
E.g. can these diagnostics be enabled/disabled with a wider scope option like -W[no-]extra or -W[no-]all?

Maybe we should add a cc1-only option --enforce-atomic-alignment and use that to determine if misalignment should be an error at the point where we issue the diagnostics?

6457 ↗

(On Diff #332658)

This should be else if, or, maybe use llvm::StringSwitch()instead:

DiagAtomicLibCall = llvm::StringSwitch<bool>(A->getValue())
   .Case("no-error=atomic-alignment", false)
   .Case("error=atomic-alignment", true)
   .Default(DiagAtomicLibCall)

[yaxunl]

separate diagnosing unaligned atomc for amdgpu to another review.

[yaxunl]

In D71726#2645269, @tra wrote:

@jyknight - James, do you have further concerns about the patch?

I separated the change about diagnosing unaligned atomics for amdgpu to https://reviews.llvm.org/D99201 since these two changes are orthogonal.

[rjmccall]

clang/lib/Sema/SemaChecking.cpp

4907

Does LLVM support atomics on all floating-point types?

[yaxunl]

clang/lib/Sema/SemaChecking.cpp

4907

LLVM IR parser requires atomicrmw value operand must have size of power of 2, therefore LLVM does not support atomicrmw on x86_fp80 which has size of 80 bytes. LLVM supports atomicrmw on all other floating-point types (bfloat, half, float, double, fp128, ppc_fp128).

[rjmccall]

clang/lib/Sema/SemaChecking.cpp

4907

Okay. So this needs to check the underlying FP semantics and disallow atomics on unsupported types.

[yaxunl]

clang/lib/Sema/SemaChecking.cpp

4907

will do

[yaxunl]

Revised by John's comments. Do not allow atomic fetch add with x86_fp80.

[rjmccall]

Alright, mostly looks good.

clang/lib/Sema/SemaChecking.cpp

4907

Could you extract this whole condition into a function and make it a bit more readable?

[yaxunl]

revised by John's comments

[rjmccall]

Thanks, LGTM

clang/include/clang/AST/ASTContext.h

 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/Linkage.h"
 #include "clang/Basic/OperatorKinds.h"
 #include "clang/Basic/PartialDiagnostic.h"
 #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"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/DenseSet.h"
 #include "llvm/ADT/FoldingSet.h"
 #include "llvm/ADT/IntrusiveRefCntPtr.h"
 #include "llvm/ADT/MapVector.h"
public:
   QualType getIntTypeForBitwidth(unsigned DestWidth,
                                  unsigned Signed) const;
 
   /// getRealTypeForBitwidth -
   /// sets floating point QualTy according to specified bitwidth.
   /// Returns empty type if there is no appropriate target types.
   QualType getRealTypeForBitwidth(unsigned DestWidth, bool ExplicitIEEE) const;
 
-  bool AtomicUsesUnsupportedLibcall(const AtomicExpr *E) const;
+  TargetInfo::AtomicOperationKind getTargetAtomicOp(const AtomicExpr *E) const;
+
+  TargetInfo::AtomicSupportKind
+  getTargetAtomicSupport(TargetInfo::AtomicOperationKind Op,
+                         QualType AtomicTy) const;
 
   const LangOptions& getLangOpts() const { return LangOpts; }
 
   // If this condition is false, typo correction must be performed eagerly
   // rather than delayed in many places, as it makes use of dependent types.
   // the condition is false for clang's C-only codepath, as it doesn't support
   // dependent types yet.
   bool isDependenceAllowed() const {

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_atomic_op_needs_non_const_atomic : Error<
   "address argument to atomic operation must be a pointer to non-%select{const|constant}0 _Atomic "
   "type (%1 invalid)">;
 def err_atomic_op_needs_non_const_pointer : Error<
   "address argument to atomic operation must be a pointer to non-const "
   "type (%0 invalid)">;
 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)">;
 def err_atomic_op_needs_atomic_int : Error<
   "address argument to atomic operation must be a pointer to "
   "%select{|atomic }0integer (%1 invalid)">;
 def warn_atomic_op_has_invalid_memory_order : Warning<
   "memory order argument to atomic operation is invalid">,
   InGroup<DiagGroup<"atomic-memory-ordering">>;
 def err_atomic_op_has_invalid_synch_scope : Error<
   "synchronization scope argument to atomic operation is invalid">;
 def warn_atomic_implicit_seq_cst : Warning<
   "implicit use of sequentially-consistent atomic may incur stronger memory barriers than necessary">,
   InGroup<DiagGroup<"atomic-implicit-seq-cst">>, DefaultIgnore;
 
 def err_overflow_builtin_must_be_int : Error<
   "operand argument to overflow builtin must be an integer (%0 invalid)">;
 def err_overflow_builtin_must_be_ptr_int : Error<
   "result argument to overflow builtin must be a pointer "
   "to a non-const integer (%0 invalid)">;
 def err_overflow_builtin_ext_int_max_size : Error<
   "__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{operation|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)">;
 def err_nontemporal_builtin_must_be_pointer_intfltptr_or_vector : Error<
   "address argument to nontemporal builtin must be a pointer to integer, float, "
   "pointer, or a vector of such types (%0 invalid)">;
 
 def err_deleted_function_use : Error<"attempt to use a deleted function">;

clang/include/clang/Basic/TargetInfo.h

public:
   /// ever be supported for the given target
   unsigned getMaxAtomicPromoteWidth() const { return MaxAtomicPromoteWidth; }
   /// Return the maximum width lock-free atomic operation which can be
   /// inlined given the supported features of the given target.
   unsigned getMaxAtomicInlineWidth() const { return MaxAtomicInlineWidth; }
   /// 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; }
   /// Return default simd alignment for the given target. Generally, this
   /// value is type-specific, but this alignment can be used for most of the
   /// types for the given target.
   unsigned getSimdDefaultAlign() const { return SimdDefaultAlign; }
 
virtual bool validateTarget(DiagnosticsEngine &Diags) const {
     return true;
   }
 
   virtual void setAuxTarget(const TargetInfo *Aux) {}
 
   /// Whether target allows debuginfo types for decl only variables.
   virtual bool allowDebugInfoForExternalVar() const { return false; }
 
+  /// Abstraction of source level atomic operations.
+  enum class AtomicOperationKind {
+    Unsupported,
+    Init,
+    C11LoadStore,
+    LoadStore,
+    AddSub,
+    MinMax,
+    LogicOp,
+    Xchg,
+    CmpXchg,
+  };
+
+  /// What is emitted in LLVM IR by clang for the atomic operation:
+  /// LockFree - LLVM atomic instructions
+  /// Library - call of library functions
+  /// Unsupported - diagnostics
+  enum class AtomicSupportKind {
+    LockFree,
+    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);
   virtual uint64_t getPointerWidthV(unsigned AddrSpace) const {
     return PointerWidth;
   }
   virtual uint64_t getPointerAlignV(unsigned AddrSpace) const {
     return PointerAlign;

clang/lib/AST/ASTContext.cpp

           TemplateParamObjectDecls.FindNodeOrInsertPos(ID, InsertPos))
     return Existing;
 
   TemplateParamObjectDecl *New = TemplateParamObjectDecl::Create(*this, T, V);
   TemplateParamObjectDecls.InsertNode(New, InsertPos);
   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::AtomicOperationKind
+ASTContext::getTargetAtomicOp(const AtomicExpr *E) const {
+  switch (E->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;
+  }
+}
+
+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
 ASTContext::ObjCMethodsAreEqual(const ObjCMethodDecl *MethodDecl,
                                 const ObjCMethodDecl *MethodImpl) {
   // No point trying to match an unavailable/deprecated mothod.
   if (MethodDecl->hasAttr<UnavailableAttr>()
       || MethodDecl->hasAttr<DeprecatedAttr>())

clang/lib/Basic/TargetInfo.cpp

void TargetInfo::CheckFixedPointBits() const {
   assert(getLongAccumIBits() >= getUnsignedLongAccumIBits());
 }
 
 void TargetInfo::copyAuxTarget(const TargetInfo *Aux) {
   auto *Target = static_cast<TransferrableTargetInfo*>(this);
   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 (Op == AtomicOperationKind::Unsupported)
+    return AtomicSupportKind::Unsupported;
+
+  if (&FS != &llvm::APFloat::Bogus() && Op == AtomicOperationKind::AddSub)
+    return getFPAtomicAddSubSupport(FS);
+
+  const llvm::Triple &T = getTriple();
+  if (Op == AtomicOperationKind::C11LoadStore && T.isOSDarwin() &&
+      ((T.isiOS() && T.isOSVersionLT(7)) ||
+       (T.isMacOSX() && T.isOSVersionLT(10, 9))) &&
+      (AtomicWidthInBits != AlignmentInBits ||
+       AtomicWidthInBits > getMaxAtomicInlineWidth())) {
+    return TargetInfo::AtomicSupportKind::Unsupported;
+  }
+  return AtomicWidthInBits <= AlignmentInBits &&
+                 AtomicWidthInBits <= getMaxAtomicInlineWidth() &&
+                 (AtomicWidthInBits <= getCharWidth() ||
+                  llvm::isPowerOf2_64(AtomicWidthInBits / getCharWidth()))
+             ? AtomicSupportKind::LockFree
+             : AtomicSupportKind::Library;
+}

clang/lib/Basic/Targets/AArch64.h

public:
   }
 
   int getEHDataRegisterNumber(unsigned RegNo) const override;
 
   const char *getBFloat16Mangling() const override { return "u6__bf16"; };
   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 {
 public:
   AArch64leTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts);
 
   void getTargetDefines(const LangOptions &Opts,
                             MacroBuilder &Builder) const override;

clang/lib/Basic/Targets/AMDGPU.h

if (!isAMDGCN(getTriple()))
       return llvm::None;
     // When -target-cpu is not set, we assume generic code that it is valid
     // for all GPU and use an empty string as target ID to represent that.
     if (GPUKind == llvm::AMDGPU::GK_NONE)
       return std::string("");
     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
 } // namespace clang
 
 #endif // LLVM_CLANG_LIB_BASIC_TARGETS_AMDGPU_H

clang/lib/Basic/Targets/ARM.h

public:
 
   int getEHDataRegisterNumber(unsigned RegNo) const override;
 
   bool hasSjLjLowering() const override;
 
   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 {
 public:
   ARMleTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts);
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 };

clang/lib/Basic/Targets/Hexagon.h

public:
   }
 
   bool isTinyCore() const {
     // We can write more stricter checks later.
     return CPU.find('t') != std::string::npos;
   }
 
   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
 #endif // LLVM_CLANG_LIB_BASIC_TARGETS_HEXAGON_H

clang/lib/Basic/Targets/Mips.h

public:
   bool hasInt128Type() const override {
     return (ABI == "n32" || ABI == "n64") || getTargetOpts().ForceEnableInt128;
   }
 
   unsigned getUnwindWordWidth() const override;
 
   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
 
 #endif // LLVM_CLANG_LIB_BASIC_TARGETS_MIPS_H

clang/lib/Basic/Targets/X86.h

uint64_t getPointerWidthV(unsigned AddrSpace) const override {
     if (AddrSpace == ptr64)
       return 64;
     return PointerWidth;
   }
 
   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
 class LLVM_LIBRARY_VISIBILITY X86_32TargetInfo : public X86TargetInfo {
 public:
   X86_32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
       : X86TargetInfo(Triple, Opts) {
     DoubleAlign = LongLongAlign = 32;

clang/lib/CodeGen/CGAtomic.cpp

class AtomicInfo {
     CodeGenFunction &CGF;
     QualType AtomicTy;
     QualType ValueTy;
     uint64_t AtomicSizeInBits;
     uint64_t ValueSizeInBits;
     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()) {
         AtomicTy = lvalue.getType();
         if (auto *ATy = AtomicTy->getAs<AtomicType>())
           ValueTy = ATy->getValueType();
         else
           ValueTy = AtomicTy;
AtomicInfo(CodeGenFunction &CGF, LValue &lvalue)
         AtomicTy = ValueTy = CGF.getContext().getExtVectorType(
             lvalue.getType(), cast<llvm::FixedVectorType>(
                                   lvalue.getExtVectorAddress().getElementType())
                                   ->getNumElements());
         AtomicSizeInBits = C.getTypeSize(AtomicTy);
         AtomicAlign = ValueAlign = lvalue.getAlignment();
         LVal = lvalue;
       }
-      UseLibcall = !C.getTargetInfo().hasBuiltinAtomic(
-          AtomicSizeInBits, C.toBits(lvalue.getAlignment()));
     }
 
     QualType getAtomicType() const { return AtomicTy; }
     QualType getValueType() const { return ValueTy; }
     CharUnits getAtomicAlignment() const { return AtomicAlign; }
     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())
         return LVal.getPointer(CGF);
       else if (LVal.isBitField())
         return LVal.getBitFieldPointer();
       else if (LVal.isVectorElt())
         return LVal.getVectorPointer();
static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
   case AtomicExpr::AO__c11_atomic_exchange:
   case AtomicExpr::AO__opencl_atomic_exchange:
   case AtomicExpr::AO__atomic_exchange_n:
   case AtomicExpr::AO__atomic_exchange:
     Op = llvm::AtomicRMWInst::Xchg;
     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:
     PostOpMinMax = true;
     LLVM_FALLTHROUGH;
   case AtomicExpr::AO__c11_atomic_fetch_min:
   case AtomicExpr::AO__opencl_atomic_fetch_min:
   case AtomicExpr::AO__atomic_fetch_min:
RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E) {
 
   auto TInfo = getContext().getTypeInfoInChars(AtomicTy);
   uint64_t Size = TInfo.Width.getQuantity();
   unsigned MaxInlineWidthInBits = getTarget().getMaxAtomicInlineWidth();
 
   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);
 
   DiagnosticsEngine &Diags = CGM.getDiags();
 
   if (Misaligned) {
     Diags.Report(E->getBeginLoc(), diag::warn_atomic_op_misaligned)
         << (int)TInfo.Width.getQuantity()
if (MemTy->isPointerType()) {
       CharUnits PointeeIncAmt =
           getContext().getTypeSizeInChars(MemTy->getPointeeType());
       Val1Scalar = Builder.CreateMul(Val1Scalar, CGM.getSize(PointeeIncAmt));
       auto Temp = CreateMemTemp(Val1Ty, ".atomictmp");
       Val1 = Temp;
       EmitStoreOfScalar(Val1Scalar, MakeAddrLValue(Temp, Val1Ty));
       break;
     }
-      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:
   case AtomicExpr::AO__opencl_atomic_exchange:
   case AtomicExpr::AO__atomic_store_n:
   case AtomicExpr::AO__atomic_exchange_n:
   case AtomicExpr::AO__c11_atomic_fetch_and:
   case AtomicExpr::AO__c11_atomic_fetch_or:
RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E) {
   QualType RValTy = E->getType().getUnqualifiedType();
 
   // The inlined atomics only function on iN types, where N is a power of 2. We
   // need to make sure (via temporaries if necessary) that all incoming values
   // are compatible.
   LValue AtomicVal = MakeAddrLValue(Ptr, AtomicTy);
   AtomicInfo Atomics(*this, AtomicVal);
 
+  if (ShouldCastToIntPtrTy) {
     Ptr = Atomics.emitCastToAtomicIntPointer(Ptr);
-  if (Val1.isValid()) Val1 = Atomics.convertToAtomicIntPointer(Val1);
-  if (Val2.isValid()) Val2 = Atomics.convertToAtomicIntPointer(Val2);
-  if (Dest.isValid())
+    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())
+  } 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) {
     bool UseOptimizedLibcall = false;
     switch (E->getOp()) {
     case AtomicExpr::AO__c11_atomic_init:
     case AtomicExpr::AO__opencl_atomic_init:
       llvm_unreachable("Already handled above with EmitAtomicInit!");

 /// An LValue is a candidate for having its loads and stores be made atomic if
 /// we are operating under /volatile:ms *and* the LValue itself is volatile and
 /// performing such an operation can be performed without a libcall.
 bool CodeGenFunction::LValueIsSuitableForInlineAtomic(LValue LV) {
   if (!CGM.getCodeGenOpts().MSVolatile) return false;
   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()))
     return false;
   return IsVolatile && AtomicIsInline;
 }
 
 RValue CodeGenFunction::EmitAtomicLoad(LValue LV, SourceLocation SL,
                                        AggValueSlot Slot) {
   llvm::AtomicOrdering AO;
   bool IsVolatile = LV.isVolatileQualified();
   if (LV.getType()->isAtomicType()) {
     AO = llvm::AtomicOrdering::SequentiallyConsistent;
   } else {
     AO = llvm::AtomicOrdering::Acquire;
     IsVolatile = true;
   }
   return EmitAtomicLoad(LV, SL, AO, IsVolatile, Slot);
 }
 
 RValue AtomicInfo::EmitAtomicLoad(AggValueSlot ResultSlot, SourceLocation Loc,
                                   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);
       TempAddr = ResultSlot.getAddress();
     } else
       TempAddr = CreateTempAlloca();
 
     EmitAtomicLoadLibcall(TempAddr.getPointer(), AO, IsVolatile);
std::pair<RValue, llvm::Value *> AtomicInfo::EmitAtomicCompareExchange(
     RValue Expected, RValue Desired, llvm::AtomicOrdering Success,
     llvm::AtomicOrdering Failure, bool IsWeak) {
   if (isStrongerThan(Failure, Success))
     // Don't assert on undefined behavior "failure argument shall be no stronger
     // than the success argument".
     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);
     auto *Res = EmitAtomicCompareExchangeLibcall(ExpectedAddr.getPointer(),
                                                  DesiredAddr.getPointer(),
                                                  Success, Failure);
     return std::make_pair(
         convertAtomicTempToRValue(ExpectedAddr, AggValueSlot::ignored(),
void AtomicInfo::EmitAtomicUpdateOp(llvm::AtomicOrdering AO, RValue UpdateRVal,
   PHI->addIncoming(Res.first, CGF.Builder.GetInsertBlock());
   CGF.Builder.CreateCondBr(Res.second, ExitBB, ContBB);
   CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
 }
 
 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);
   }
 }
 
 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);
   }
 }
 
 void CodeGenFunction::EmitAtomicStore(RValue rvalue, LValue lvalue,
                                       bool isInit) {
void CodeGenFunction::EmitAtomicStore(RValue rvalue, LValue dest,
   // If this is an initialization, just put the value there normally.
   if (LVal.isSimple()) {
     if (isInit) {
       atomics.emitCopyIntoMemory(rvalue);
       return;
     }
 
     // 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);
 
       // void __atomic_store(size_t size, void *mem, void *val, int order)
       CallArgList args;
       args.add(RValue::get(atomics.getAtomicSizeValue()),
                getContext().getSizeType());
       args.add(RValue::get(EmitCastToVoidPtr(atomics.getAtomicPointer())),

clang/lib/CodeGen/CGStmtOpenMP.cpp

static void emitOMPAtomicWriteExpr(CodeGenFunction &CGF,
   case llvm::AtomicOrdering::Monotonic:
     break;
   case llvm::AtomicOrdering::NotAtomic:
   case llvm::AtomicOrdering::Unordered:
     llvm_unreachable("Unexpected ordering.");
   }
 }
 
+static 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 TargetInfo::AtomicOperationKind::Unsupported;
+  }
+}
+
 static std::pair<bool, RValue> emitOMPAtomicRMW(CodeGenFunction &CGF, LValue X,
                                                 RValue Update,
                                                 BinaryOperatorKind BO,
                                                 llvm::AtomicOrdering AO,
                                                 bool IsXLHSInRHSPart) {
   ASTContext &Context = CGF.getContext();
   // 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.
   if (BO == BO_Comma || !Update.isScalar() ||
       !Update.getScalarVal()->getType()->isIntegerTy() || !X.isSimple() ||
       (!isa<llvm::ConstantInt>(Update.getScalarVal()) &&
        (Update.getScalarVal()->getType() !=
         X.getAddress(CGF).getElementType())) ||
       !X.getAddress(CGF).getElementType()->isIntegerTy() ||
-      !Context.getTargetInfo().hasBuiltinAtomic(
-          Context.getTypeSize(X.getType()), Context.toBits(X.getAlignment())))
+      Context.getTargetInfo().getAtomicSupport(
+          getTargetAtomicOp(BO), Context.getTypeSize(X.getType()),
+          Context.toBits(X.getAlignment())) !=
+          TargetInfo::AtomicSupportKind::LockFree)
     return std::make_pair(false, RValue::get(nullptr));
 
   llvm::AtomicRMWInst::BinOp RMWOp;
   switch (BO) {
   case BO_Add:
     RMWOp = llvm::AtomicRMWInst::Add;
     break;
   case BO_Sub:

clang/lib/Sema/SemaChecking.cpp

 #include <functional>
 #include <limits>
 #include <string>
 #include <tuple>
 #include <utility>
 
 using namespace clang;
 using namespace sema;
+using AtomicSupportKind = TargetInfo::AtomicSupportKind;
+using TargetAtomicOpKind = TargetInfo::AtomicOperationKind;
 
 SourceLocation Sema::getLocationOfStringLiteralByte(const StringLiteral *SL,
                                                     unsigned ByteNo) const {
   return SL->getLocationOfByte(ByteNo, getSourceManager(), LangOpts,
                                Context.getTargetInfo());
 }
 
 /// Checks that a call expression's argument count is the desired number.
ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
   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:
     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:
   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:
+    Form = Arithmetic;
+    break;
   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:
if (ValType.isConstQualified()) {
           << Ptr->getType() << Ptr->getSourceRange();
       return ExprError();
     }
   }
 
   // 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();
     }
     if (!IsAddSub && !ValType->isIntegerType()) {
       Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int)
           << IsC11 << Ptr->getType() << Ptr->getSourceRange();
       return ExprError();
     }
if (i < NumVals[Form] + 1) {
         // Nothing else to do: we already know all we want about this pointer.
         continue;
       case 1:
         // The second argument is the non-atomic operand. For arithmetic, this
         // is always passed by value, and for a compare_exchange it is always
         // 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) {
             // The value pointer is always dereferenced, a nullptr is undefined.
             CheckNonNullArgument(*this, APIOrderedArgs[i],
                                  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());
           LangAS AS = LangAS::Default;
           // Keep address space of non-atomic pointer type.
           if (const PointerType *PtrTy =
                   ValArg->getType()->getAs<PointerType>()) {
             AS = PtrTy->getPointeeType().getAddressSpace();
if (Optional<llvm::APSInt> Result =
             << Scope->getSourceRange();
     }
     SubExprs.push_back(Scope);
   }
 
   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 = Context.getTargetAtomicOp(AE);
+  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);
     return ExprError();
   }
 
   return AE;
 }

clang/test/CodeGen/fp-atomic-ops.c

+// 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
+}

clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu

+// 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);
+}

clang/test/CodeGenOpenCL/atomic-ops.cl

-// 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)));
 
 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,

 }
 
 float ff3(atomic_float *d) {
   // CHECK-LABEL: @ff3
   // CHECK: atomicrmw xchg i32* {{.*}} syncscope("workgroup") seq_cst
   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()
 {
   // CHECK-NOT: atomic
   // CHECK: store
   __opencl_atomic_init(&j, 42);
 
   // CHECK-NOT: atomic

clang/test/Sema/atomic-ops.c

 _Static_assert(__atomic_always_lock_free(4, &i64), "");
 _Static_assert(!__atomic_always_lock_free(8, &i32), "");
 _Static_assert(__atomic_always_lock_free(8, &i64), "");
 
 #define _AS1 __attribute__((address_space(1)))
 #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}}
   __c11_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const _Atomic(int) *' invalid)}}
 
   __c11_atomic_load(0); // expected-error {{too few arguments to function}}
   __c11_atomic_load(0,0,0); // expected-error {{too many arguments to function}}
   __c11_atomic_store(0,0,0); // expected-error {{address argument to atomic builtin must be a pointer}}
   __c11_atomic_store((int*)0,0,0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic}}
   __c11_atomic_store(i, 0, memory_order_relaxed);
   __c11_atomic_store(ci, 0, memory_order_relaxed); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const _Atomic(int) *' invalid)}}
 
   __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);
   int *load_n_2 = __atomic_load_n(P, memory_order_relaxed);
   float load_n_3 = __atomic_load_n(D, memory_order_relaxed); // expected-error {{must be a pointer to integer or pointer}}
   __atomic_load_n(s1, memory_order_relaxed); // expected-error {{must be a pointer to integer or pointer}}
   load_n_1 = __atomic_load_n(CI, memory_order_relaxed);
 
   __atomic_load(i, I, memory_order_relaxed); // expected-error {{must be a pointer to a trivially-copyable type}}
   __atomic_load(CI, I, memory_order_relaxed);
 
   __atomic_load(I, i, memory_order_relaxed); // expected-warning {{passing '_Atomic(int) *' to parameter of type 'int *'}}
   __atomic_load(I, *P, memory_order_relaxed);
   __atomic_load(I, *P, memory_order_relaxed, 42); // expected-error {{too many arguments}}
   (int)__atomic_load(I, I, memory_order_seq_cst); // expected-error {{operand of type 'void'}}
   __atomic_load(s1, s2, memory_order_acquire);
   __atomic_load(CI, I, memory_order_relaxed);
   __atomic_load(I, CI, memory_order_relaxed); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
   __atomic_load(CI, CI, memory_order_relaxed); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
 
   __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);
   __atomic_store_n(CI, 4, memory_order_release); // expected-error {{address argument to atomic operation must be a pointer to non-const type ('const int *' invalid)}}
   __atomic_store_n(I, P, memory_order_release); // expected-warning {{parameter of type 'int'}}
   __atomic_store_n(i, 1, memory_order_release); // expected-error {{must be a pointer to integer or pointer}}
   __atomic_store_n(s1, *s2, memory_order_release); // expected-error {{must be a pointer to integer or pointer}}
   __atomic_store_n(I, I, memory_order_release); // expected-warning {{incompatible pointer to integer conversion passing 'int *' to parameter of type 'int'; dereference with *}}
void f(_Atomic(int) *i, const _Atomic(int) *ci,
   __atomic_exchange(s1, I, P, memory_order_seq_cst); // expected-warning 2{{parameter of type 'struct S *'}}
   (int)__atomic_exchange(s1, s2, s2, memory_order_seq_cst); // expected-error {{operand of type 'void'}}
   __atomic_exchange(I, I, I, memory_order_seq_cst);
   __atomic_exchange(CI, I, I, memory_order_seq_cst); // expected-error {{address argument to atomic operation must be a pointer to non-const type ('const int *' invalid)}}
   __atomic_exchange(I, I, CI, memory_order_seq_cst); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
 
   __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);
   __atomic_fetch_xor(P, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
   __atomic_fetch_or(D, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
   __atomic_fetch_and(s1, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
 
   _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);
   _Bool cmpexch_5 = __atomic_compare_exchange_n(I, P, 5, 0, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{; dereference with *}}
   _Bool cmpexch_6 = __atomic_compare_exchange_n(I, I, P, 0, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{passing 'int **' to parameter of type 'int'}}
   (void)__atomic_compare_exchange_n(CI, I, 5, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-error {{address argument to atomic operation must be a pointer to non-const type ('const int *' invalid)}}
   (void)__atomic_compare_exchange_n(I, CI, 5, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
 

clang/test/Sema/atomic-requires-library-error.c

 // RUN: %clang_cc1 %s -triple=i686-apple-darwin9 -verify
 // rdar://13973577
 
 struct foo {
   int big[128];
 };
 struct bar {
   char c[3];
 };
 
 struct bar smallThing;
 struct foo bigThing;
 _Atomic(struct foo) bigAtomic;
 
 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);
 
   __atomic_store(&bigThing, &f, 5);
 }
 
 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);
 
   __atomic_load(&bigThing, &f, 5);
 }

clang/test/SemaCUDA/amdgpu-atomic-ops.cu

+// 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}}
+}

clang/test/SemaOpenCL/atomic-ops.cl

-// 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)));
 
 typedef enum memory_order {
   memory_order_relaxed = __ATOMIC_RELAXED,
   memory_order_acquire = __ATOMIC_ACQUIRE,
   memory_order_release = __ATOMIC_RELEASE,

 
 char i8;
 short i16;
 int i32;
 int8 i64;
 
 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,
        constant atomic_int *i_c) {
   __opencl_atomic_init(I, 5); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
   __opencl_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
 
   __opencl_atomic_load(0); // expected-error {{too few arguments to function call, expected 3, have 1}}
   __opencl_atomic_load(0, 0, 0, 0); // expected-error {{too many arguments to function call, expected 3, have 4}}
   __opencl_atomic_store(0,0,0,0); // expected-error {{address argument to atomic builtin must be a pointer}}
   __opencl_atomic_store((int *)0, 0, 0, 0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
   __opencl_atomic_store(i, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(ci, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
   __opencl_atomic_store(i_g, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(i_l, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(i_p, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(i_c, 0, memory_order_relaxed, 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_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.
   bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
   __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
   __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
   __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group);