[clang] [Clang][AMDGPU] Add ``amdgcn_av("none")`` attribute for atomic expressions (PR #199622)
Sameer Sahasrabuddhe via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 30 03:23:43 PDT 2026
https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/199622
>From 21c30bc99dac05a432dda1502b0b9f3e3ef1d6e3 Mon Sep 17 00:00:00 2001
From: Sameer Sahasrabuddhe <sameer.sahasrabuddhe at amd.com>
Date: Tue, 7 Jul 2026 19:57:07 +0530
Subject: [PATCH] [Clang][AMDGPU] Add amdgpu_av("none") attribute for atomic
expressions
Add a statement attribute that suppresses MakeAvailable/MakeVisible
cache operations on AMDGPU atomic instructions while preserving memory
ordering (waits).
The attribute takes a string argument specifying the mode. Currently "none" is
the only supported mode. The resulting atomic or fence instruction carries !mmra
!{!"amdgcn-av", !"none"} metadata.
Assisted-By: Claude Opus 4.6
---
clang/include/clang/Basic/Attr.td | 6 +
clang/include/clang/Basic/AttrDocs.td | 26 ++++
clang/include/clang/Basic/DiagnosticGroups.td | 2 +
.../clang/Basic/DiagnosticSemaKinds.td | 7 +
clang/lib/CodeGen/CGAtomic.cpp | 16 ++-
clang/lib/CodeGen/CGBuiltin.cpp | 53 +++----
clang/lib/CodeGen/CGStmt.cpp | 5 +
clang/lib/CodeGen/CodeGenFunction.h | 13 ++
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 10 ++
clang/lib/CodeGen/Targets/AMDGPU.cpp | 3 +
clang/lib/Sema/SemaStmtAttr.cpp | 60 ++++++++
.../CodeGen/AMDGPU/amdgcn-av-non-atomic.cpp | 43 ++++++
.../AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp | 23 ++++
.../AMDGPU/amdgcn-av-none-attr-c-atomic.c | 72 ++++++++++
.../CodeGen/AMDGPU/amdgcn-av-none-attr.cpp | 130 ++++++++++++++++++
15 files changed, 442 insertions(+), 27 deletions(-)
create mode 100644 clang/test/CodeGen/AMDGPU/amdgcn-av-non-atomic.cpp
create mode 100644 clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp
create mode 100644 clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr-c-atomic.c
create mode 100644 clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr.cpp
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 39c672322d515..4afd5bda853e7 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -2532,6 +2532,12 @@ def AMDGPUMaxNumWorkGroups : InheritableAttr {
let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
}
+def AMDGPUAvailableVisible : StmtAttr, TargetSpecificAttr<TargetAMDGPU> {
+ let Spellings = [Clang<"amdgpu_av">];
+ let Args = [StringArgument<"Mode">];
+ let Documentation = [AMDGPUAvailableVisibleDocs];
+}
+
def BPFPreserveAccessIndex : InheritableAttr,
TargetSpecificAttr<TargetBPF> {
let Spellings = [Clang<"preserve_access_index">];
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 05e4cb0870652..19c8ee53b951c 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -3604,6 +3604,32 @@ An error will be given if:
}];
}
+def AMDGPUAvailableVisibleDocs : Documentation {
+ let Category = DocCatAMDGPUAttributes;
+ let Content = [{
+
+This attribute controls availability and visibility as described in the `AMDGPU
+Memory Model <https://llvm.org/docs/AMDGPUMemoryModel.html>`__. When placed on
+an atomic expression or fence, the resulting atomic or fence instruction carries
+the corresponding *AV Metadata*.
+
+The attribute takes a string literal as an argument, which currently has only
+one supported value:
+
+- ``"none"``: Disable MakeAvailable and MakeVisible semantics on release and
+ acquire operations respectively.
+
+.. code-block:: c++
+
+ [[clang::amdgpu_av("none")]] __atomic_thread_fence(__ATOMIC_SEQ_CST);
+ [[clang::amdgpu_av("none")]] __atomic_fetch_add(ptr, 1, __ATOMIC_ACQ_REL);
+
+ // Also works with _Atomic type qualifier operations.
+ _Atomic int *p;
+ [[clang::amdgpu_av("none")]] *p += 1;
+ }];
+}
+
def DocCatCallingConvs : DocumentationCategory<"Calling Conventions"> {
let Content = [{
Clang supports several different calling conventions, depending on the target
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index b7072634cccf3..698f15c57aa9e 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -2007,3 +2007,5 @@ def ExperimentalOption : DiagGroup<"experimental-option">;
// Warnings about unguarded usages of AMDGPU target specific constructs
def UnguardedBuiltinUsageAMDGPU : DiagGroup<"amdgpu-unguarded-builtin-usage">;
+
+def AvailableVisibleAMDGPU : DiagGroup<"amdgpu-av">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index d9d0d485f16ac..3e98d264a0a09 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -14324,6 +14324,13 @@ def note_amdgcn_unguarded_builtin_silence
def err_amdgcn_global_or_flat_pointer_required : Error<"builtin requires a global or generic pointer">;
+def warn_amdgpu_av_requires_expr : Warning<
+ "%0 attribute can only be applied to an expression">,
+ InGroup<AvailableVisibleAMDGPU>;
+def warn_amdgpu_av_requires_atomic : Warning<
+ "%0 attribute only applies to atomic expressions or Clang atomic builtins">,
+ InGroup<AvailableVisibleAMDGPU>;
+
def err_amdgcn_dmask_has_too_many_bits_set
: Error<"dmask argument cannot have more bits set than there are elements "
"in return type">;
diff --git a/clang/lib/CodeGen/CGAtomic.cpp b/clang/lib/CodeGen/CGAtomic.cpp
index 61e13d4925020..cd7b278eab15d 100644
--- a/clang/lib/CodeGen/CGAtomic.cpp
+++ b/clang/lib/CodeGen/CGAtomic.cpp
@@ -631,6 +631,7 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
llvm::LoadInst *Load = CGF.Builder.CreateLoad(Ptr);
Load->setAtomic(Order, Scope);
Load->setVolatile(E->isVolatile());
+ CGF.getTargetHooks().setTargetAtomicMetadata(CGF, *Load, E);
CGF.maybeAttachRangeForLoad(Load, E->getValueType(), E->getExprLoc());
auto *I = CGF.Builder.CreateStore(Load, Dest);
CGF.addInstToCurrentSourceAtom(I, Load);
@@ -648,6 +649,7 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
llvm::StoreInst *Store = CGF.Builder.CreateStore(LoadVal1, Ptr);
Store->setAtomic(Order, Scope);
Store->setVolatile(E->isVolatile());
+ CGF.getTargetHooks().setTargetAtomicMetadata(CGF, *Store, E);
CGF.addInstToCurrentSourceAtom(Store, LoadVal1);
return;
}
@@ -822,6 +824,7 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
CGF.Builder.CreateStore(CGF.Builder.getInt8(0), Ptr);
Store->setAtomic(Order, Scope);
Store->setVolatile(E->isVolatile());
+ CGF.getTargetHooks().setTargetAtomicMetadata(CGF, *Store, E);
CGF.addInstToCurrentSourceAtom(Store, nullptr);
return;
}
@@ -1629,6 +1632,7 @@ llvm::Value *AtomicInfo::EmitAtomicLoadOp(llvm::AtomicOrdering AO,
Addr = castToAtomicIntPointer(Addr);
llvm::LoadInst *Load = CGF.Builder.CreateLoad(Addr, "atomic-load");
Load->setAtomic(AO);
+ CGF.getTargetHooks().setTargetAtomicMetadata(CGF, *Load);
// Other decoration.
if (IsVolatile)
@@ -1796,6 +1800,7 @@ std::pair<llvm::Value *, llvm::Value *> AtomicInfo::EmitAtomicCompareExchangeOp(
// Other decoration.
Inst->setVolatile(LVal.isVolatileQualified());
Inst->setWeak(IsWeak);
+ CGF.getTargetHooks().setTargetAtomicMetadata(CGF, *Inst);
// Okay, turn that back into the original value type.
auto *PreviousVal = CGF.Builder.CreateExtractValue(Inst, /*Idxs=*/0);
@@ -2151,8 +2156,10 @@ void CodeGenFunction::EmitAtomicStore(RValue rvalue, LValue dest,
else if (AO == llvm::AtomicOrdering::AcquireRelease)
AO = llvm::AtomicOrdering::Release;
// Initializations don't need to be atomic.
- if (!isInit)
+ if (!isInit) {
store->setAtomic(AO);
+ getTargetHooks().setTargetAtomicMetadata(*this, *store);
+ }
// Other decoration.
if (IsVolatile)
@@ -2196,6 +2203,13 @@ CodeGenFunction::emitAtomicRMWInst(llvm::AtomicRMWInst::BinOp Op, Address Addr,
return RMW;
}
+llvm::FenceInst *CodeGenFunction::emitAtomicFence(llvm::AtomicOrdering Order,
+ llvm::SyncScope::ID SSID) {
+ llvm::FenceInst *Fence = Builder.CreateFence(Order, SSID);
+ getTargetHooks().setTargetAtomicMetadata(*this, *Fence);
+ return Fence;
+}
+
void CodeGenFunction::EmitAtomicUpdate(
LValue LVal, llvm::AtomicOrdering AO,
const llvm::function_ref<RValue(RValue)> &UpdateOp, bool IsVolatile) {
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index f38a891272474..2dd038da20864 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -5692,16 +5692,16 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
break;
case 1: // memory_order_consume
case 2: // memory_order_acquire
- Builder.CreateFence(llvm::AtomicOrdering::Acquire, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::Acquire, SSID);
break;
case 3: // memory_order_release
- Builder.CreateFence(llvm::AtomicOrdering::Release, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::Release, SSID);
break;
case 4: // memory_order_acq_rel
- Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::AcquireRelease, SSID);
break;
case 5: // memory_order_seq_cst
- Builder.CreateFence(llvm::AtomicOrdering::SequentiallyConsistent, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::SequentiallyConsistent, SSID);
break;
}
return RValue::get(nullptr);
@@ -5718,23 +5718,23 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
llvm::SwitchInst *SI = Builder.CreateSwitch(Order, ContBB);
Builder.SetInsertPoint(AcquireBB);
- Builder.CreateFence(llvm::AtomicOrdering::Acquire, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::Acquire, SSID);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(1), AcquireBB);
SI->addCase(Builder.getInt32(2), AcquireBB);
Builder.SetInsertPoint(ReleaseBB);
- Builder.CreateFence(llvm::AtomicOrdering::Release, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::Release, SSID);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(3), ReleaseBB);
Builder.SetInsertPoint(AcqRelBB);
- Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::AcquireRelease, SSID);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(4), AcqRelBB);
Builder.SetInsertPoint(SeqCstBB);
- Builder.CreateFence(llvm::AtomicOrdering::SequentiallyConsistent, SSID);
+ emitAtomicFence(llvm::AtomicOrdering::SequentiallyConsistent, SSID);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(5), SeqCstBB);
@@ -5758,32 +5758,32 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
break;
case 1: // memory_order_consume
case 2: // memory_order_acquire
- Builder.CreateFence(
+ emitAtomicFence(
llvm::AtomicOrdering::Acquire,
getTargetHooks().getLLVMSyncScopeID(getLangOpts(), SS,
llvm::AtomicOrdering::Acquire,
getLLVMContext()));
break;
case 3: // memory_order_release
- Builder.CreateFence(
+ emitAtomicFence(
llvm::AtomicOrdering::Release,
getTargetHooks().getLLVMSyncScopeID(getLangOpts(), SS,
llvm::AtomicOrdering::Release,
getLLVMContext()));
break;
case 4: // memory_order_acq_rel
- Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease,
- getTargetHooks().getLLVMSyncScopeID(
- getLangOpts(), SS,
- llvm::AtomicOrdering::AcquireRelease,
- getLLVMContext()));
+ emitAtomicFence(llvm::AtomicOrdering::AcquireRelease,
+ getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), SS,
+ llvm::AtomicOrdering::AcquireRelease,
+ getLLVMContext()));
break;
case 5: // memory_order_seq_cst
- Builder.CreateFence(llvm::AtomicOrdering::SequentiallyConsistent,
- getTargetHooks().getLLVMSyncScopeID(
- getLangOpts(), SS,
- llvm::AtomicOrdering::SequentiallyConsistent,
- getLLVMContext()));
+ emitAtomicFence(
+ llvm::AtomicOrdering::SequentiallyConsistent,
+ getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), SS, llvm::AtomicOrdering::SequentiallyConsistent,
+ getLLVMContext()));
break;
}
return RValue::get(nullptr);
@@ -5844,9 +5844,9 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
SyncScope SS = ScopeModel->isValid(Scp->getZExtValue())
? ScopeModel->map(Scp->getZExtValue())
: ScopeModel->map(ScopeModel->getFallBackValue());
- Builder.CreateFence(Ordering,
- getTargetHooks().getLLVMSyncScopeID(
- getLangOpts(), SS, Ordering, getLLVMContext()));
+ emitAtomicFence(Ordering,
+ getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), SS, Ordering, getLLVMContext()));
Builder.CreateBr(ContBB);
} else {
llvm::DenseMap<unsigned, llvm::BasicBlock *> BBs;
@@ -5860,9 +5860,10 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
SI->addCase(Builder.getInt32(Scp), B);
Builder.SetInsertPoint(B);
- Builder.CreateFence(Ordering, getTargetHooks().getLLVMSyncScopeID(
- getLangOpts(), ScopeModel->map(Scp),
- Ordering, getLLVMContext()));
+ emitAtomicFence(Ordering,
+ getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), ScopeModel->map(Scp), Ordering,
+ getLLVMContext()));
Builder.CreateBr(ContBB);
}
}
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 27e74d966eca1..aeaf2e763fcf4 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -783,6 +783,7 @@ void CodeGenFunction::EmitAttributedStmt(const AttributedStmt &S) {
bool noinline = false;
bool alwaysinline = false;
bool noconvergent = false;
+ StringRef amdgpuAVMode;
HLSLControlFlowHintAttr::Spelling flattenOrBranch =
HLSLControlFlowHintAttr::SpellingNotCalculated;
const CallExpr *musttail = nullptr;
@@ -820,6 +821,9 @@ void CodeGenFunction::EmitAttributedStmt(const AttributedStmt &S) {
case attr::Atomic:
AA = cast<AtomicAttr>(A);
break;
+ case attr::AMDGPUAvailableVisible:
+ amdgpuAVMode = cast<AMDGPUAvailableVisibleAttr>(A)->getMode();
+ break;
case attr::HLSLControlFlowHint: {
flattenOrBranch = cast<HLSLControlFlowHintAttr>(A)->getSemanticSpelling();
} break;
@@ -829,6 +833,7 @@ void CodeGenFunction::EmitAttributedStmt(const AttributedStmt &S) {
SaveAndRestore save_noinline(InNoInlineAttributedStmt, noinline);
SaveAndRestore save_alwaysinline(InAlwaysInlineAttributedStmt, alwaysinline);
SaveAndRestore save_noconvergent(InNoConvergentAttributedStmt, noconvergent);
+ SaveAndRestore save_amdgpuav(AMDGPUAvailableVisibleMode, amdgpuAVMode);
SaveAndRestore save_musttail(MustTailCall, musttail);
SaveAndRestore save_flattenOrBranch(HLSLControlFlowAttr, flattenOrBranch);
CGAtomicOptionsRAII AORAII(CGM, AA);
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 3fc9052adb87b..783f97dc354eb 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -624,6 +624,10 @@ class CodeGenFunction : public CodeGenTypeCache {
/// True if the current statement has noconvergent attribute.
bool InNoConvergentAttributedStmt = false;
+ /// The mode string from the amdgpu_av attribute on the current statement,
+ /// or empty if the attribute is not present.
+ StringRef AMDGPUAvailableVisibleMode;
+
/// HLSL Branch attribute.
HLSLControlFlowHintAttr::Spelling HLSLControlFlowAttr =
HLSLControlFlowHintAttr::SpellingNotCalculated;
@@ -4364,6 +4368,12 @@ class CodeGenFunction : public CodeGenTypeCache {
llvm::SyncScope::ID SSID = llvm::SyncScope::System,
const AtomicExpr *AE = nullptr);
+ /// Emit a fence instruction, applying relevant target-specific metadata when
+ /// applicable.
+ llvm::FenceInst *
+ emitAtomicFence(llvm::AtomicOrdering Order,
+ llvm::SyncScope::ID SSID = llvm::SyncScope::System);
+
void EmitAtomicUpdate(LValue LVal, llvm::AtomicOrdering AO,
const llvm::function_ref<RValue(RValue)> &UpdateOp,
bool IsVolatile);
@@ -4984,6 +4994,9 @@ class CodeGenFunction : public CodeGenTypeCache {
void AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,
const CallExpr *E);
+ /// Attach the AMDGPU availability/visibility MMRA to \p Inst when the
+ /// amdgpu_av attribute is active on the current statement.
+ void AddAMDGPUAvailableVisibleMMRA(llvm::Instruction *Inst);
void ProcessOrderScopeAMDGCN(llvm::Value *Order, llvm::Value *Scope,
llvm::AtomicOrdering &AO,
llvm::SyncScope::ID &SSID);
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 72d6f536165c5..cc50ce95e55e6 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -473,6 +473,14 @@ void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,
MMRAMetadata::appendTags(*Inst, MMRAs);
}
+void CodeGenFunction::AddAMDGPUAvailableVisibleMMRA(llvm::Instruction *Inst) {
+ if (AMDGPUAvailableVisibleMode.empty())
+ return;
+
+ constexpr const char *Tag = "amdgcn-av";
+ MMRAMetadata::appendTags(*Inst, {{Tag, AMDGPUAvailableVisibleMode}});
+}
+
static Value *GetAMDGPUPredicate(CodeGenFunction &CGF, Twine Name) {
Constant *SpecId = ConstantInt::getAllOnesValue(CGF.Int32Ty);
@@ -1905,6 +1913,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
FenceInst *Fence = Builder.CreateFence(AO, SSID);
if (E->getNumArgs() > 2)
AddAMDGPUFenceAddressSpaceMMRA(Fence, E);
+ getTargetHooks().setTargetAtomicMetadata(*this, *Fence);
return Fence;
}
case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
@@ -2017,6 +2026,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
Builder.CreateAtomicRMW(BinOp, Ptr, Val, AO, SSID);
if (Volatile)
RMW->setVolatile(true);
+ AddAMDGPUAvailableVisibleMMRA(RMW);
unsigned AddrSpace = Ptr.getType()->getAddressSpace();
if (AddrSpace != llvm::AMDGPUAS::LOCAL_ADDRESS) {
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 7b37f3f7f9b6e..b51996c6d88bf 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -10,6 +10,7 @@
#include "TargetInfo.h"
#include "clang/AST/DeclCXX.h"
#include "llvm/ADT/StringExtras.h"
+#include "llvm/IR/MemoryModelRelaxationAnnotations.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
using namespace clang;
@@ -560,6 +561,8 @@ void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata(
AtomicInst.setMetadata(llvm::LLVMContext::MD_noalias_addrspace, ASRange);
}
+ CGF.AddAMDGPUAvailableVisibleMMRA(&AtomicInst);
+
if (!RMW)
return;
diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp
index 3e9dc9c5e7e08..bad5a3e0fb919 100644
--- a/clang/lib/Sema/SemaStmtAttr.cpp
+++ b/clang/lib/Sema/SemaStmtAttr.cpp
@@ -363,6 +363,64 @@ static Attr *handleMustTailAttr(Sema &S, Stmt *St, const ParsedAttr &A,
return ::new (S.Context) MustTailAttr(S.Context, A);
}
+/// Return true if E is an atomic expression or a fence.
+static bool isAtomicExprOrFence(const Expr *E) {
+ E = E->IgnoreParenCasts();
+
+ if (isa<AtomicExpr>(E))
+ return true;
+
+ // _Atomic type qualifier operations: assignments and compound assignments
+ // to atomic lvalues, and loads from atomic lvalues.
+ if (const auto *BO = dyn_cast<BinaryOperator>(E)) {
+ if (BO->getLHS()->getType()->isAtomicType())
+ return true;
+ } else if (E->getType()->isAtomicType()) {
+ return true;
+ }
+
+ // Target-independent fence builtins.
+ if (const auto *CE = dyn_cast<CallExpr>(E)) {
+ switch (CE->getBuiltinCallee()) {
+ case Builtin::BI__c11_atomic_thread_fence:
+ case Builtin::BI__c11_atomic_signal_fence:
+ case Builtin::BI__atomic_thread_fence:
+ case Builtin::BI__atomic_signal_fence:
+ case Builtin::BI__scoped_atomic_thread_fence:
+ return true;
+ default:
+ break;
+ }
+ }
+
+ return false;
+}
+
+static Attr *handleAMDGPUAvailableVisibleAttr(Sema &S, Stmt *St,
+ const ParsedAttr &A,
+ SourceRange Range) {
+ StringRef Mode;
+ if (!S.checkStringLiteralArgumentAttr(A, 0, Mode))
+ return nullptr;
+
+ if (Mode != "none") {
+ S.Diag(A.getLoc(), diag::warn_attribute_type_not_supported) << A << Mode;
+ return nullptr;
+ }
+
+ if (const auto *E = dyn_cast<Expr>(St)) {
+ if (!isAtomicExprOrFence(E)) {
+ S.Diag(A.getLoc(), diag::warn_amdgpu_av_requires_atomic) << A;
+ return nullptr;
+ }
+ } else {
+ S.Diag(A.getLoc(), diag::warn_amdgpu_av_requires_expr) << A;
+ return nullptr;
+ }
+
+ return ::new (S.Context) AMDGPUAvailableVisibleAttr(S.Context, A, Mode);
+}
+
static Attr *handleLikely(Sema &S, Stmt *St, const ParsedAttr &A,
SourceRange Range) {
@@ -742,6 +800,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
return handleNoInlineAttr(S, St, A, Range);
case ParsedAttr::AT_MustTail:
return handleMustTailAttr(S, St, A, Range);
+ case ParsedAttr::AT_AMDGPUAvailableVisible:
+ return handleAMDGPUAvailableVisibleAttr(S, St, A, Range);
case ParsedAttr::AT_Likely:
return handleLikely(S, St, A, Range);
case ParsedAttr::AT_Unlikely:
diff --git a/clang/test/CodeGen/AMDGPU/amdgcn-av-non-atomic.cpp b/clang/test/CodeGen/AMDGPU/amdgcn-av-non-atomic.cpp
new file mode 100644
index 0000000000000..46fbdec7b7a41
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/amdgcn-av-non-atomic.cpp
@@ -0,0 +1,43 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 %s -emit-llvm -O0 -verify -o - \
+// RUN: -triple=amdgcn-amd-amdhsa | FileCheck %s
+
+// Test that [[clang::amdgpu_av("none")]] on non-atomic statements emits a
+// warning and does NOT produce !mmra metadata.
+
+// CHECK-LABEL: define {{.*}} @_Z16test_plain_storePii(
+// CHECK-NOT: !mmra
+// CHECK: ret void
+void test_plain_store(int *p, int val) {
+ [[clang::amdgpu_av("none")]] *p = val; // expected-warning {{'clang::amdgpu_av' attribute only applies to atomic expressions or Clang atomic builtins}}
+}
+
+// CHECK-LABEL: define {{.*}} @_Z15test_plain_callv(
+// CHECK-NOT: !mmra
+// CHECK: ret void
+extern void foo();
+void test_plain_call() {
+ [[clang::amdgpu_av("none")]] foo(); // expected-warning {{'clang::amdgpu_av' attribute only applies to atomic expressions or Clang atomic builtins}}
+}
+
+// CHECK-LABEL: define {{.*}} @_Z18test_for_with_atomPi(
+// CHECK-NOT: !mmra
+// CHECK: ret void
+void test_for_with_atom(int *p) {
+ // The attribute on a for loop should warn even if the body contains atomics.
+ [[clang::amdgpu_av("none")]] for (;;) { // expected-warning {{'clang::amdgpu_av' attribute can only be applied to an expression}}
+ __atomic_fetch_add(p, 1, __ATOMIC_SEQ_CST);
+ break;
+ }
+}
+
+// The attribute on an if statement should warn even if the condition is atomic.
+// CHECK-LABEL: define {{.*}} @_Z20test_if_atomic_condnPi(
+// CHECK-NOT: !mmra
+// CHECK: ret void
+void test_if_atomic_condn(int *p) {
+ [[clang::amdgpu_av("none")]] if (__atomic_load_n(p, __ATOMIC_ACQUIRE)) { // expected-warning {{'clang::amdgpu_av' attribute can only be applied to an expression}}
+ }
+}
+
+// CHECK-NOT: amdgcn-av
diff --git a/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp
new file mode 100644
index 0000000000000..ada72942dcd4c
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp
@@ -0,0 +1,23 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 %s -verify -emit-llvm -O0 -o - \
+// RUN: -triple=amdgcn-amd-amdhsa \
+// RUN: -target-cpu gfx942
+
+// Test that [[clang::amdgpu_av("none")]] on AMDGPU-specific atomic builtins
+// is rejected with a warning.
+
+void test_atomic_inc32_av(volatile unsigned *p, unsigned val) {
+ [[clang::amdgpu_av("none")]] __builtin_amdgcn_atomic_inc32(p, val, __ATOMIC_SEQ_CST, "agent"); // expected-warning {{only applies to atomic expressions or Clang atomic builtins}}
+}
+
+void test_atomic_dec64_av(volatile unsigned long long *p, unsigned long long val) {
+ [[clang::amdgpu_av("none")]] __builtin_amdgcn_atomic_dec64(p, val, __ATOMIC_ACQUIRE, "agent"); // expected-warning {{only applies to atomic expressions or Clang atomic builtins}}
+}
+
+void test_ds_faddf_av(float __attribute__((address_space(3))) *p, float val) {
+ [[clang::amdgpu_av("none")]] __builtin_amdgcn_ds_faddf(p, val, __ATOMIC_ACQUIRE, __ATOMIC_ACQUIRE, false); // expected-warning {{only applies to atomic expressions or Clang atomic builtins}}
+}
+
+void test_global_atomic_fadd_av(float __attribute__((address_space(1))) *p, float val) {
+ [[clang::amdgpu_av("none")]] __builtin_amdgcn_global_atomic_fadd_f32(p, val); // expected-warning {{only applies to atomic expressions or Clang atomic builtins}}
+}
diff --git a/clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr-c-atomic.c b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr-c-atomic.c
new file mode 100644
index 0000000000000..9ae0e88a27a6a
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr-c-atomic.c
@@ -0,0 +1,72 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
+// RUN: -triple=amdgcn-amd-amdhsa -std=c23 | FileCheck %s
+
+// Test that [[clang::amdgpu_av("none")]] works with _Atomic type
+// qualifier operations, which go through CGExprScalar and CGAtomic
+// codegen paths.
+
+// CHECK-LABEL: define dso_local void @test_atomic_rmw_add(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw add ptr [[TMP0]], i32 1 seq_cst, align 4, !mmra [[META2:![0-9]+]], !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]]
+// CHECK-NEXT: [[TMP2:%.*]] = add i32 [[TMP1]], 1
+// CHECK-NEXT: ret void
+//
+void test_atomic_rmw_add(_Atomic int *p) {
+ [[clang::amdgpu_av("none")]] *p += 1;
+}
+
+// CHECK-LABEL: define dso_local void @test_atomic_store(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store atomic i32 42, ptr [[TMP0]] seq_cst, align 4, !mmra [[META2]]
+// CHECK-NEXT: ret void
+//
+void test_atomic_store(_Atomic int *p) {
+ [[clang::amdgpu_av("none")]] *p = 42;
+}
+
+// CHECK-LABEL: define dso_local void @test_atomic_load(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[ATOMIC_LOAD:%.*]] = load atomic i32, ptr [[TMP0]] seq_cst, align 4, !mmra [[META2]]
+// CHECK-NEXT: ret void
+//
+void test_atomic_load(_Atomic int *p) {
+ [[clang::amdgpu_av("none")]] (void)*p;
+}
+
+// --- Negative control: no attribute ---
+
+// CHECK-LABEL: define dso_local void @test_no_attr(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw add ptr [[TMP0]], i32 1 seq_cst, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]]
+// CHECK-NEXT: [[TMP2:%.*]] = add i32 [[TMP1]], 1
+// CHECK-NEXT: ret void
+//
+void test_no_attr(_Atomic int *p) {
+ *p += 1;
+}
+//.
+// CHECK: [[META2]] = !{!"amdgcn-av", !"none"}
+// CHECK: [[META3]] = !{}
+//.
diff --git a/clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr.cpp b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr.cpp
new file mode 100644
index 0000000000000..38071f466c3a8
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-attr.cpp
@@ -0,0 +1,130 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
+// RUN: -triple=amdgcn-amd-amdhsa | FileCheck %s
+
+// Test that [[clang::amdgpu_av("none")]] on atomic builtins emits
+// !mmra !{!"amdgcn-av", !"none"} metadata on the resulting IR instructions.
+
+// CHECK-LABEL: define dso_local void @_Z21test_atomic_fetch_addPi(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[DOTATOMICTMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[ATOMIC_TEMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] seq_cst, align 4, !mmra [[META2:![0-9]+]], !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]]
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
+// CHECK-NEXT: ret void
+//
+void test_atomic_fetch_add(int *p) {
+ [[clang::amdgpu_av("none")]] __atomic_fetch_add(p, 1, __ATOMIC_SEQ_CST);
+}
+
+// CHECK-LABEL: define dso_local void @_Z17test_atomic_storePi(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[DOTATOMICTMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 42, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: store atomic i32 [[TMP1]], ptr [[TMP0]] release, align 4, !mmra [[META2]]
+// CHECK-NEXT: ret void
+//
+void test_atomic_store(int *p) {
+ [[clang::amdgpu_av("none")]] __atomic_store_n(p, 42, __ATOMIC_RELEASE);
+}
+
+// CHECK-LABEL: define dso_local void @_Z16test_atomic_loadPi(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[ATOMIC_TEMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load atomic i32, ptr [[TMP0]] acquire, align 4, !mmra [[META2]]
+// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
+// CHECK-NEXT: ret void
+//
+void test_atomic_load(int *p) {
+ [[clang::amdgpu_av("none")]] __atomic_load_n(p, __ATOMIC_ACQUIRE);
+}
+
+// CHECK-LABEL: define dso_local void @_Z19test_atomic_cmpxchgPi(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[EXPECTED:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[DOTATOMICTMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[CMPXCHG_BOOL:%.*]] = alloca i8, align 1, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: [[EXPECTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[EXPECTED]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 0, ptr [[EXPECTED_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[EXPECTED_ASCAST]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] seq_cst acquire, align 4, !mmra [[META2]]
+// CHECK-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0
+// CHECK-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1
+// CHECK-NEXT: br i1 [[TMP5]], label [[CMPXCHG_CONTINUE:%.*]], label [[CMPXCHG_STORE_EXPECTED:%.*]]
+// CHECK: cmpxchg.store_expected:
+// CHECK-NEXT: store i32 [[TMP4]], ptr [[EXPECTED_ASCAST]], align 4
+// CHECK-NEXT: br label [[CMPXCHG_CONTINUE]]
+// CHECK: cmpxchg.continue:
+// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[TMP5]] to i8
+// CHECK-NEXT: store i8 [[STOREDV]], ptr addrspace(5) [[CMPXCHG_BOOL]], align 1
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr addrspace(5) [[CMPXCHG_BOOL]], align 1
+// CHECK-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP6]], 0
+// CHECK-NEXT: ret void
+//
+void test_atomic_cmpxchg(int *p) {
+ int expected = 0;
+ [[clang::amdgpu_av("none")]] __atomic_compare_exchange_n(p, &expected, 1, false, __ATOMIC_SEQ_CST, __ATOMIC_ACQUIRE);
+}
+
+// CHECK-LABEL: define dso_local void @_Z24test_atomic_thread_fencev(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: fence seq_cst, !mmra [[META2]]
+// CHECK-NEXT: ret void
+//
+void test_atomic_thread_fence() {
+ [[clang::amdgpu_av("none")]] __atomic_thread_fence(__ATOMIC_SEQ_CST);
+}
+
+// Verify that without the attribute, no !mmra metadata is emitted.
+// CHECK-LABEL: define dso_local void @_Z12test_no_attrPi(
+// CHECK-SAME: ptr noundef [[P:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[DOTATOMICTMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[ATOMIC_TEMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] seq_cst, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]]
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
+// CHECK-NEXT: ret void
+//
+void test_no_attr(int *p) {
+ __atomic_fetch_add(p, 1, __ATOMIC_SEQ_CST);
+}
+//.
+// CHECK: [[META2]] = !{!"amdgcn-av", !"none"}
+// CHECK: [[META3]] = !{}
+//.
More information about the cfe-commits
mailing list