[clang] [Clang][AMDGPU] Add ``amdgcn_av("none")`` attribute for atomic expressions (PR #199622)
Sameer Sahasrabuddhe via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 15 01:21:35 PDT 2026
https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/199622
>From 4788542685f82933c2df59e25727d4b6ab8b274e 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 amdgcn_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 | 9 +-
clang/lib/CodeGen/CGBuiltin.cpp | 68 +++++----
clang/lib/CodeGen/CGStmt.cpp | 5 +
clang/lib/CodeGen/CodeGenFunction.h | 4 +
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 4 +
clang/lib/CodeGen/Targets/AMDGPU.cpp | 6 +
clang/lib/Sema/SemaStmtAttr.cpp | 59 ++++++++
.../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, 437 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 1a5bd2301dfc8..de6996f6f9823 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 AMDGCNAV : StmtAttr, TargetSpecificAttr<TargetAMDGPU> {
+ let Spellings = [Clang<"amdgcn_av">];
+ let Args = [StringArgument<"Mode">];
+ let Documentation = [AMDGCNAVDocs];
+}
+
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..e410dad399c2b 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 AMDGCNAVDocs : 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 a
+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::amdgcn_av("none")]] __atomic_thread_fence(__ATOMIC_SEQ_CST);
+ [[clang::amdgcn_av("none")]] __atomic_fetch_add(ptr, 1, __ATOMIC_ACQ_REL);
+
+ // Also works with _Atomic type qualifier operations.
+ _Atomic int *p;
+ [[clang::amdgcn_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 79583534b9bbd..e213211bd132a 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -2006,3 +2006,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<"amdgcn-av">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 3ff7e30d9f1f7..c081b34305651 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -14342,6 +14342,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_amdgcn_av_requires_expr : Warning<
+ "%0 attribute can only be applied to an expression">,
+ InGroup<AvailableVisibleAMDGPU>;
+def warn_amdgcn_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..29671169d645d 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)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index a41ac703dafed..901cd1aa607a1 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -5671,24 +5671,28 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
Value *Order = EmitScalarExpr(E->getArg(0));
if (isa<llvm::ConstantInt>(Order)) {
int ord = cast<llvm::ConstantInt>(Order)->getZExtValue();
+ llvm::FenceInst *Fence = nullptr;
switch (ord) {
case 0: // memory_order_relaxed
default: // invalid order
break;
case 1: // memory_order_consume
case 2: // memory_order_acquire
- Builder.CreateFence(llvm::AtomicOrdering::Acquire, SSID);
+ Fence = Builder.CreateFence(llvm::AtomicOrdering::Acquire, SSID);
break;
case 3: // memory_order_release
- Builder.CreateFence(llvm::AtomicOrdering::Release, SSID);
+ Fence = Builder.CreateFence(llvm::AtomicOrdering::Release, SSID);
break;
case 4: // memory_order_acq_rel
- Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease, SSID);
+ Fence = Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease, SSID);
break;
case 5: // memory_order_seq_cst
- Builder.CreateFence(llvm::AtomicOrdering::SequentiallyConsistent, SSID);
+ Fence = Builder.CreateFence(
+ llvm::AtomicOrdering::SequentiallyConsistent, SSID);
break;
}
+ if (Fence)
+ getTargetHooks().setTargetAtomicMetadata(*this, *Fence);
return RValue::get(nullptr);
}
@@ -5703,23 +5707,29 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
llvm::SwitchInst *SI = Builder.CreateSwitch(Order, ContBB);
Builder.SetInsertPoint(AcquireBB);
- Builder.CreateFence(llvm::AtomicOrdering::Acquire, SSID);
+ getTargetHooks().setTargetAtomicMetadata(
+ *this, *Builder.CreateFence(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);
+ getTargetHooks().setTargetAtomicMetadata(
+ *this, *Builder.CreateFence(llvm::AtomicOrdering::Release, SSID));
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(3), ReleaseBB);
Builder.SetInsertPoint(AcqRelBB);
- Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease, SSID);
+ getTargetHooks().setTargetAtomicMetadata(
+ *this,
+ *Builder.CreateFence(llvm::AtomicOrdering::AcquireRelease, SSID));
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(4), AcqRelBB);
Builder.SetInsertPoint(SeqCstBB);
- Builder.CreateFence(llvm::AtomicOrdering::SequentiallyConsistent, SSID);
+ getTargetHooks().setTargetAtomicMetadata(
+ *this, *Builder.CreateFence(
+ llvm::AtomicOrdering::SequentiallyConsistent, SSID));
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(5), SeqCstBB);
@@ -5737,40 +5747,43 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
SyncScope SS = ScopeModel->isValid(Scp->getZExtValue())
? ScopeModel->map(Scp->getZExtValue())
: ScopeModel->map(ScopeModel->getFallBackValue());
+ llvm::FenceInst *Fence = nullptr;
switch (Ord->getZExtValue()) {
case 0: // memory_order_relaxed
default: // invalid order
break;
case 1: // memory_order_consume
case 2: // memory_order_acquire
- Builder.CreateFence(
+ Fence = Builder.CreateFence(
llvm::AtomicOrdering::Acquire,
getTargetHooks().getLLVMSyncScopeID(getLangOpts(), SS,
llvm::AtomicOrdering::Acquire,
getLLVMContext()));
break;
case 3: // memory_order_release
- Builder.CreateFence(
+ Fence = Builder.CreateFence(
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()));
+ Fence = Builder.CreateFence(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()));
+ Fence = Builder.CreateFence(
+ llvm::AtomicOrdering::SequentiallyConsistent,
+ getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), SS, llvm::AtomicOrdering::SequentiallyConsistent,
+ getLLVMContext()));
break;
}
+ if (Fence)
+ getTargetHooks().setTargetAtomicMetadata(*this, *Fence);
return RValue::get(nullptr);
}
@@ -5829,9 +5842,10 @@ 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()));
+ llvm::FenceInst *Fence = Builder.CreateFence(
+ Ordering, getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), SS, Ordering, getLLVMContext()));
+ getTargetHooks().setTargetAtomicMetadata(*this, *Fence);
Builder.CreateBr(ContBB);
} else {
llvm::DenseMap<unsigned, llvm::BasicBlock *> BBs;
@@ -5845,9 +5859,11 @@ 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()));
+ llvm::FenceInst *Fence = Builder.CreateFence(
+ Ordering, getTargetHooks().getLLVMSyncScopeID(
+ getLangOpts(), ScopeModel->map(Scp), Ordering,
+ getLLVMContext()));
+ getTargetHooks().setTargetAtomicMetadata(*this, *Fence);
Builder.CreateBr(ContBB);
}
}
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 27e74d966eca1..ff1f2cbf87407 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 amdgcnAVMode;
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::AMDGCNAV:
+ amdgcnAVMode = cast<AMDGCNAVAttr>(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_amdgcnav(AMDGCNAVMode, amdgcnAVMode);
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 e7c24f1f36f1e..d9d6075df0aab 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 amdgcn_av attribute on the current statement,
+ /// or empty if the attribute is not present.
+ StringRef AMDGCNAVMode;
+
/// HLSL Branch attribute.
HLSLControlFlowHintAttr::Spelling HLSLControlFlowAttr =
HLSLControlFlowHintAttr::SpellingNotCalculated;
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 66e4e688e33ee..ba316f2f9b05b 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -1905,6 +1905,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 +2018,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
Builder.CreateAtomicRMW(BinOp, Ptr, Val, AO, SSID);
if (Volatile)
RMW->setVolatile(true);
+ if (!AMDGCNAVMode.empty()) {
+ llvm::MMRAMetadata::appendTags(*RMW, {{"amdgcn-av", AMDGCNAVMode}});
+ }
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..77c25366bb64e 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,11 @@ void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata(
AtomicInst.setMetadata(llvm::LLVMContext::MD_noalias_addrspace, ASRange);
}
+ if (!CGF.AMDGCNAVMode.empty()) {
+ llvm::MMRAMetadata::appendTags(AtomicInst,
+ {{"amdgcn-av", CGF.AMDGCNAVMode}});
+ }
+
if (!RMW)
return;
diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp
index 3e9dc9c5e7e08..fb858ff4926ca 100644
--- a/clang/lib/Sema/SemaStmtAttr.cpp
+++ b/clang/lib/Sema/SemaStmtAttr.cpp
@@ -363,6 +363,63 @@ 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 *handleAMDGCNAVAttr(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_amdgcn_av_requires_atomic) << A;
+ return nullptr;
+ }
+ } else {
+ S.Diag(A.getLoc(), diag::warn_amdgcn_av_requires_expr) << A;
+ return nullptr;
+ }
+
+ return ::new (S.Context) AMDGCNAVAttr(S.Context, A, Mode);
+}
+
static Attr *handleLikely(Sema &S, Stmt *St, const ParsedAttr &A,
SourceRange Range) {
@@ -742,6 +799,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_AMDGCNAV:
+ return handleAMDGCNAVAttr(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..531b89e4c74d2
--- /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::amdgcn_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::amdgcn_av("none")]] *p = val; // expected-warning {{'clang::amdgcn_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::amdgcn_av("none")]] foo(); // expected-warning {{'clang::amdgcn_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::amdgcn_av("none")]] for (;;) { // expected-warning {{'clang::amdgcn_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::amdgcn_av("none")]] if (__atomic_load_n(p, __ATOMIC_ACQUIRE)) { // expected-warning {{'clang::amdgcn_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..5c291aed37b66
--- /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::amdgcn_av("none")]] on AMDGPU-specific atomic builtins
+// is rejected with a warning.
+
+void test_atomic_inc32_av(volatile unsigned *p, unsigned val) {
+ [[clang::amdgcn_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::amdgcn_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::amdgcn_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::amdgcn_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..dc68addeeacd7
--- /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::amdgcn_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::amdgcn_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::amdgcn_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::amdgcn_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..b9e2a45eba60c
--- /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::amdgcn_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::amdgcn_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::amdgcn_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::amdgcn_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::amdgcn_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::amdgcn_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