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 <[email protected]> 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]] = !{} +//. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
