https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/213810
>From dbf7ad30a38fbbbb54d448e339ca9af2e6968c57 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Mon, 3 Aug 2026 15:12:47 -0500 Subject: [PATCH 1/3] [LLVM] Unify atomic scope handling between LLVM and Clang Summary: This PR attempts to unify the atomic scope handling between clang and LLVM. The difficulty is that clang emits ABI-mandated integers for these while the backends use arbitrary strings. We still want the flexibility of these being arbitrary strings, but we should put this in a single source of truth. The main motivation is so IR passes can determine which scope clang used when creating these, and so future additions all go in one place. This was modeled after the AtomicOrdering.h file. --- clang/include/clang/Basic/SyncScope.h | 17 +-- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 24 ++-- clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 95 ++++++++----- clang/lib/CodeGen/TargetInfo.h | 30 ++++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 31 +--- clang/lib/CodeGen/Targets/NVPTX.cpp | 26 +--- clang/lib/CodeGen/Targets/SPIR.cpp | 26 +--- clang/lib/Frontend/InitPreprocessor.cpp | 25 ++-- llvm/include/llvm/TargetParser/AtomicScope.h | 134 ++++++++++++++++++ .../Target/AMDGPU/AMDGPUBarrierLatency.cpp | 14 +- .../Target/AMDGPU/AMDGPUMachineModuleInfo.cpp | 31 ++-- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 5 +- llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 21 +-- llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 4 +- .../Target/SPIRV/SPIRVInstructionSelector.cpp | 17 +-- llvm/lib/Target/SPIRV/SPIRVUtils.cpp | 16 ++- llvm/lib/Target/SPIRV/SPIRVUtils.h | 4 +- .../TargetParser/AtomicScopeTest.cpp | 37 +++++ llvm/unittests/TargetParser/CMakeLists.txt | 1 + 19 files changed, 374 insertions(+), 184 deletions(-) create mode 100644 llvm/include/llvm/TargetParser/AtomicScope.h create mode 100644 llvm/unittests/TargetParser/AtomicScopeTest.cpp diff --git a/clang/include/clang/Basic/SyncScope.h b/clang/include/clang/Basic/SyncScope.h index 7776c3d83a77d..719c8f764c079 100644 --- a/clang/include/clang/Basic/SyncScope.h +++ b/clang/include/clang/Basic/SyncScope.h @@ -17,6 +17,7 @@ #include "clang/Basic/LangOptions.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/StringRef.h" +#include "llvm/TargetParser/AtomicScope.h" #include <memory> namespace clang { @@ -240,15 +241,15 @@ class AtomicScopeHIPModel : public AtomicScopeModel { class AtomicScopeGenericModel : public AtomicScopeModel { public: /// The enum values match predefined built-in macros __MEMORY_SCOPE_*. - /// These may be present in pch files or bitcode so preserve existing values - /// when adding a new ID. + /// These are ABI-sensitive (they may be present in pch files or bitcode) so + /// existing values must be preserved. enum ID { - System = 0, - Device = 1, - Workgroup = 2, - Wavefront = 3, - Single = 4, - Cluster = 5, + System = static_cast<unsigned>(llvm::AtomicScope::System), + Device = static_cast<unsigned>(llvm::AtomicScope::Device), + Workgroup = static_cast<unsigned>(llvm::AtomicScope::Workgroup), + Wavefront = static_cast<unsigned>(llvm::AtomicScope::Wavefront), + Single = static_cast<unsigned>(llvm::AtomicScope::Single), + Cluster = static_cast<unsigned>(llvm::AtomicScope::Cluster), Count, Last = Count - 1 }; diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 720f956130cc1..a6d6c04589007 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -24,6 +24,7 @@ #include "llvm/IR/MemoryModelRelaxationAnnotations.h" #include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/AtomicOrdering.h" +#include "llvm/TargetParser/AtomicScope.h" using namespace clang; using namespace CodeGen; @@ -380,11 +381,15 @@ static Value *emitFPIntBuiltin(CodeGenFunction &CGF, return CGF.Builder.CreateCall(F, {Src0, Src1}); } -static inline StringRef mapScopeToSPIRV(StringRef AMDGCNScope) { - if (AMDGCNScope == "agent") - return "device"; - if (AMDGCNScope == "wavefront") - return "subgroup"; +// When the target is SPIR-V (spirv64-amd-amdhsa) re-spell the scope for that +// target by parsing as AMDGPU and re-emitting it. +static inline StringRef mapScopeToSPIRV(const llvm::Triple &TargetTriple, + StringRef AMDGCNScope) { + static const llvm::Triple AMDGPU("amdgcn-amd-amdhsa"); + if (auto Parsed = llvm::parseAtomicScopeIRString(AMDGPU, AMDGCNScope)) + if (auto Str = llvm::getAtomicScopeIRString(TargetTriple, Parsed->first, + Parsed->second)) + return *Str; return AMDGCNScope; } @@ -449,7 +454,7 @@ void CodeGenFunction::ProcessOrderScopeAMDGCN(Value *Order, Value *Scope, StringRef scp; if (llvm::getConstantStringInfo(Scope, scp)) { if (getTarget().getTriple().isSPIRV()) - scp = mapScopeToSPIRV(scp); + scp = mapScopeToSPIRV(getTarget().getTriple(), scp); SSID = getLLVMContext().getOrInsertSyncScopeID(scp); return; } @@ -2004,10 +2009,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, // // The global/flat cases need to use agent scope to consistently produce // the native instruction instead of a cmpxchg expansion. - if (getTarget().getTriple().isSPIRV()) - SSID = getLLVMContext().getOrInsertSyncScopeID("device"); - else - SSID = getLLVMContext().getOrInsertSyncScopeID("agent"); + SSID = + getLLVMContext().getOrInsertSyncScopeID(*llvm::getAtomicScopeIRString( + getTarget().getTriple(), llvm::AtomicScope::Device)); AO = AtomicOrdering::Monotonic; // The v2bf16 builtin uses i16 instead of a natural bfloat type. diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp index 48b6866147281..64fdae9d8934d 100644 --- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp @@ -13,6 +13,7 @@ #include "CGBuiltin.h" #include "clang/Basic/TargetBuiltins.h" #include "llvm/IR/IntrinsicsNVPTX.h" +#include "llvm/TargetParser/AtomicScope.h" using namespace clang; using namespace CodeGen; @@ -352,28 +353,28 @@ static Value *MakeLdg(CodeGenFunction &CGF, const CallExpr *E) { return LD; } -// Set `Scope` to: -// - "block" for _cta builtins, and -// - "" for _sys builtins. +// `Scope` is AtomicScope::Workgroup for _cta builtins and AtomicScope::System +// for _sys builtins. static Value *MakeScopedAtomicRMW(CodeGenFunction &CGF, const CallExpr *E, llvm::AtomicRMWInst::BinOp Kind, - StringRef Scope) { + llvm::AtomicScope Scope) { Address Ptr = CGF.EmitPointerWithAlignment(E->getArg(0)); Value *Val = CGF.EmitScalarExpr(E->getArg(1)); - llvm::SyncScope::ID SSID = CGF.getLLVMContext().getOrInsertSyncScopeID(Scope); + llvm::SyncScope::ID SSID = CGF.getLLVMContext().getOrInsertSyncScopeID( + *llvm::getAtomicScopeIRString(CGF.getTarget().getTriple(), Scope)); return CGF.Builder.CreateAtomicRMW(Kind, Ptr, Val, llvm::AtomicOrdering::Monotonic, SSID); } -// Set `Scope` to: -// - "block" for _cta builtins, and -// - "" for _sys builtins. +// `Scope` is AtomicScope::Workgroup for _cta builtins and AtomicScope::System +// for _sys builtins. static Value *MakeScopedAtomicCAS(CodeGenFunction &CGF, const CallExpr *E, - StringRef Scope) { + llvm::AtomicScope Scope) { Address Ptr = CGF.EmitPointerWithAlignment(E->getArg(0)); Value *Cmp = CGF.EmitScalarExpr(E->getArg(1)); Value *New = CGF.EmitScalarExpr(E->getArg(2)); - llvm::SyncScope::ID SSID = CGF.getLLVMContext().getOrInsertSyncScopeID(Scope); + llvm::SyncScope::ID SSID = CGF.getLLVMContext().getOrInsertSyncScopeID( + *llvm::getAtomicScopeIRString(CGF.getTarget().getTriple(), Scope)); Value *Pair = CGF.Builder.CreateAtomicCmpXchg( Ptr, Cmp, New, llvm::AtomicOrdering::Monotonic, llvm::AtomicOrdering::Monotonic, SSID); @@ -597,101 +598,123 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, case NVPTX::BI__nvvm_atom_cta_add_gen_i: case NVPTX::BI__nvvm_atom_cta_add_gen_l: case NVPTX::BI__nvvm_atom_cta_add_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Add, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Add, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_add_gen_i: case NVPTX::BI__nvvm_atom_sys_add_gen_l: case NVPTX::BI__nvvm_atom_sys_add_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Add, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Add, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_add_gen_f: case NVPTX::BI__nvvm_atom_cta_add_gen_d: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::FAdd, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::FAdd, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_add_gen_f: case NVPTX::BI__nvvm_atom_sys_add_gen_d: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::FAdd, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::FAdd, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_xchg_gen_i: case NVPTX::BI__nvvm_atom_cta_xchg_gen_l: case NVPTX::BI__nvvm_atom_cta_xchg_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xchg, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xchg, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_xchg_gen_i: case NVPTX::BI__nvvm_atom_sys_xchg_gen_l: case NVPTX::BI__nvvm_atom_sys_xchg_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xchg, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xchg, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_max_gen_i: case NVPTX::BI__nvvm_atom_cta_max_gen_l: case NVPTX::BI__nvvm_atom_cta_max_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Max, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Max, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_cta_max_gen_ui: case NVPTX::BI__nvvm_atom_cta_max_gen_ul: case NVPTX::BI__nvvm_atom_cta_max_gen_ull: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMax, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMax, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_max_gen_i: case NVPTX::BI__nvvm_atom_sys_max_gen_l: case NVPTX::BI__nvvm_atom_sys_max_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Max, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Max, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_sys_max_gen_ui: case NVPTX::BI__nvvm_atom_sys_max_gen_ul: case NVPTX::BI__nvvm_atom_sys_max_gen_ull: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMax, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMax, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_min_gen_i: case NVPTX::BI__nvvm_atom_cta_min_gen_l: case NVPTX::BI__nvvm_atom_cta_min_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Min, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Min, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_cta_min_gen_ui: case NVPTX::BI__nvvm_atom_cta_min_gen_ul: case NVPTX::BI__nvvm_atom_cta_min_gen_ull: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMin, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMin, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_min_gen_i: case NVPTX::BI__nvvm_atom_sys_min_gen_l: case NVPTX::BI__nvvm_atom_sys_min_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Min, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Min, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_sys_min_gen_ui: case NVPTX::BI__nvvm_atom_sys_min_gen_ul: case NVPTX::BI__nvvm_atom_sys_min_gen_ull: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMin, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UMin, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_inc_gen_ui: return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UIncWrap, - "block"); + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_cta_dec_gen_ui: return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UDecWrap, - "block"); + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_inc_gen_ui: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UIncWrap, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UIncWrap, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_sys_dec_gen_ui: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UDecWrap, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::UDecWrap, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_and_gen_i: case NVPTX::BI__nvvm_atom_cta_and_gen_l: case NVPTX::BI__nvvm_atom_cta_and_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::And, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::And, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_and_gen_i: case NVPTX::BI__nvvm_atom_sys_and_gen_l: case NVPTX::BI__nvvm_atom_sys_and_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::And, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::And, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_or_gen_i: case NVPTX::BI__nvvm_atom_cta_or_gen_l: case NVPTX::BI__nvvm_atom_cta_or_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Or, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Or, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_or_gen_i: case NVPTX::BI__nvvm_atom_sys_or_gen_l: case NVPTX::BI__nvvm_atom_sys_or_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Or, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Or, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_xor_gen_i: case NVPTX::BI__nvvm_atom_cta_xor_gen_l: case NVPTX::BI__nvvm_atom_cta_xor_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xor, "block"); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xor, + llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_xor_gen_i: case NVPTX::BI__nvvm_atom_sys_xor_gen_l: case NVPTX::BI__nvvm_atom_sys_xor_gen_ll: - return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xor, ""); + return MakeScopedAtomicRMW(*this, E, llvm::AtomicRMWInst::Xor, + llvm::AtomicScope::System); case NVPTX::BI__nvvm_atom_cta_cas_gen_us: case NVPTX::BI__nvvm_atom_cta_cas_gen_i: case NVPTX::BI__nvvm_atom_cta_cas_gen_l: case NVPTX::BI__nvvm_atom_cta_cas_gen_ll: - return MakeScopedAtomicCAS(*this, E, "block"); + return MakeScopedAtomicCAS(*this, E, llvm::AtomicScope::Workgroup); case NVPTX::BI__nvvm_atom_sys_cas_gen_us: case NVPTX::BI__nvvm_atom_sys_cas_gen_i: case NVPTX::BI__nvvm_atom_sys_cas_gen_l: case NVPTX::BI__nvvm_atom_sys_cas_gen_ll: - return MakeScopedAtomicCAS(*this, E, ""); + return MakeScopedAtomicCAS(*this, E, llvm::AtomicScope::System); case NVPTX::BI__nvvm_match_all_sync_i32p: case NVPTX::BI__nvvm_match_all_sync_i64p: { Value *Mask = EmitScalarExpr(E->getArg(0)); diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h index 9405cf0124b29..fae32068b1570 100644 --- a/clang/lib/CodeGen/TargetInfo.h +++ b/clang/lib/CodeGen/TargetInfo.h @@ -23,6 +23,7 @@ #include "clang/Basic/TargetInfo.h" #include "llvm/ADT/SmallString.h" #include "llvm/ADT/StringRef.h" +#include "llvm/TargetParser/AtomicScope.h" namespace llvm { class Constant; @@ -35,6 +36,35 @@ namespace clang { class CXXRecordDecl; class Decl; +/// Collapses a clang sync scope onto the target-neutral llvm::AtomicScope. +inline llvm::AtomicScope getAtomicScope(SyncScope S) { + switch (S) { + case SyncScope::HIPSingleThread: + case SyncScope::SingleScope: + return llvm::AtomicScope::Single; + case SyncScope::HIPWavefront: + case SyncScope::OpenCLSubGroup: + case SyncScope::WavefrontScope: + return llvm::AtomicScope::Wavefront; + case SyncScope::HIPWorkgroup: + case SyncScope::OpenCLWorkGroup: + case SyncScope::WorkgroupScope: + return llvm::AtomicScope::Workgroup; + case SyncScope::HIPCluster: + case SyncScope::ClusterScope: + return llvm::AtomicScope::Cluster; + case SyncScope::HIPAgent: + case SyncScope::OpenCLDevice: + case SyncScope::DeviceScope: + return llvm::AtomicScope::Device; + case SyncScope::SystemScope: + case SyncScope::HIPSystem: + case SyncScope::OpenCLAllSVMDevices: + return llvm::AtomicScope::System; + } + llvm_unreachable("Invalid sync scope"); +} + namespace CodeGen { class ABIInfo; class CallArgList; diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 7b37f3f7f9b6e..ea1b1ef683444 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -510,32 +510,11 @@ StringRef AMDGPUTargetCodeGenInfo::getLLVMSyncScopeStr( Scope <= SyncScope::OpenCLSubGroup && Ordering != llvm::AtomicOrdering::SequentiallyConsistent); - switch (Scope) { - case SyncScope::HIPSingleThread: - case SyncScope::SingleScope: - return IsOneAs ? "singlethread-one-as" : "singlethread"; - case SyncScope::HIPWavefront: - case SyncScope::OpenCLSubGroup: - case SyncScope::WavefrontScope: - return IsOneAs ? "wavefront-one-as" : "wavefront"; - case SyncScope::HIPCluster: - case SyncScope::ClusterScope: - assert(!IsOneAs && "OpenCL does not have cluster scope"); - return "cluster"; - case SyncScope::HIPWorkgroup: - case SyncScope::OpenCLWorkGroup: - case SyncScope::WorkgroupScope: - return IsOneAs ? "workgroup-one-as" : "workgroup"; - case SyncScope::HIPAgent: - case SyncScope::OpenCLDevice: - case SyncScope::DeviceScope: - return IsOneAs ? "agent-one-as" : "agent"; - case SyncScope::SystemScope: - case SyncScope::HIPSystem: - case SyncScope::OpenCLAllSVMDevices: - return IsOneAs ? "one-as" : ""; - } - llvm_unreachable("Unknown SyncScope enum"); + llvm::AtomicScope AS = getAtomicScope(Scope); + assert((AS != llvm::AtomicScope::Cluster || !IsOneAs) && + "OpenCL does not have cluster scope"); + return *llvm::getAtomicScopeIRString(getABIInfo().getTarget().getTriple(), AS, + IsOneAs); } void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata( diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 2e307fc01ccda..00fe271c85b0b 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -307,30 +307,8 @@ bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const { StringRef NVPTXTargetCodeGenInfo::getLLVMSyncScopeStr( const LangOptions &LangOpts, SyncScope Scope, llvm::AtomicOrdering Ordering) const { - switch (Scope) { - case SyncScope::HIPSingleThread: - case SyncScope::SingleScope: - return "singlethread"; - case SyncScope::HIPWavefront: - case SyncScope::OpenCLSubGroup: - case SyncScope::WavefrontScope: - case SyncScope::HIPWorkgroup: - case SyncScope::OpenCLWorkGroup: - case SyncScope::WorkgroupScope: - return "block"; - case SyncScope::HIPCluster: - case SyncScope::ClusterScope: - return "cluster"; - case SyncScope::HIPAgent: - case SyncScope::OpenCLDevice: - case SyncScope::DeviceScope: - return "device"; - case SyncScope::SystemScope: - case SyncScope::HIPSystem: - case SyncScope::OpenCLAllSVMDevices: - return ""; - } - llvm_unreachable("Unknown SyncScope enum"); + return *llvm::getAtomicScopeIRString(getABIInfo().getTarget().getTriple(), + getAtomicScope(Scope)); } llvm::Constant * diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index c8f7c03474be9..a260469133304 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -560,30 +560,8 @@ void SPIRVTargetCodeGenInfo::setTargetAttributes( StringRef SPIRVTargetCodeGenInfo::getLLVMSyncScopeStr( const LangOptions &, SyncScope Scope, llvm::AtomicOrdering) const { - switch (Scope) { - case SyncScope::HIPSingleThread: - case SyncScope::SingleScope: - return "singlethread"; - case SyncScope::HIPWavefront: - case SyncScope::OpenCLSubGroup: - case SyncScope::WavefrontScope: - return "subgroup"; - case SyncScope::HIPCluster: - case SyncScope::ClusterScope: - case SyncScope::HIPWorkgroup: - case SyncScope::OpenCLWorkGroup: - case SyncScope::WorkgroupScope: - return "workgroup"; - case SyncScope::HIPAgent: - case SyncScope::OpenCLDevice: - case SyncScope::DeviceScope: - return "device"; - case SyncScope::SystemScope: - case SyncScope::HIPSystem: - case SyncScope::OpenCLAllSVMDevices: - return ""; - } - return ""; + return *llvm::getAtomicScopeIRString(getABIInfo().getTarget().getTriple(), + getAtomicScope(Scope)); } void SPIRVTargetCodeGenInfo::setTargetAtomicMetadata( diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 8b6ff844d0daa..58b637776c3e7 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -894,21 +894,20 @@ static void InitializePredefinedMacros(const TargetInfo &TI, Builder.defineMacro("__ATOMIC_SEQ_CST", "5"); // Define macros for the clang atomic scopes. - Builder.defineMacro("__MEMORY_SCOPE_SYSTEM", "0"); - Builder.defineMacro("__MEMORY_SCOPE_DEVICE", "1"); - Builder.defineMacro("__MEMORY_SCOPE_WRKGRP", "2"); - Builder.defineMacro("__MEMORY_SCOPE_WVFRNT", "3"); - Builder.defineMacro("__MEMORY_SCOPE_SINGLE", "4"); - Builder.defineMacro("__MEMORY_SCOPE_CLUSTR", "5"); + Builder.defineMacro("__MEMORY_SCOPE_SYSTEM", + Twine(AtomicScopeGenericModel::System)); + Builder.defineMacro("__MEMORY_SCOPE_DEVICE", + Twine(AtomicScopeGenericModel::Device)); + Builder.defineMacro("__MEMORY_SCOPE_WRKGRP", + Twine(AtomicScopeGenericModel::Workgroup)); + Builder.defineMacro("__MEMORY_SCOPE_WVFRNT", + Twine(AtomicScopeGenericModel::Wavefront)); + Builder.defineMacro("__MEMORY_SCOPE_SINGLE", + Twine(AtomicScopeGenericModel::Single)); + Builder.defineMacro("__MEMORY_SCOPE_CLUSTR", + Twine(AtomicScopeGenericModel::Cluster)); // Define macros for the OpenCL memory scope. - // The values should match AtomicScopeOpenCLModel::ID enum. - static_assert( - static_cast<unsigned>(AtomicScopeOpenCLModel::WorkGroup) == 1 && - static_cast<unsigned>(AtomicScopeOpenCLModel::Device) == 2 && - static_cast<unsigned>(AtomicScopeOpenCLModel::AllSVMDevices) == 3 && - static_cast<unsigned>(AtomicScopeOpenCLModel::SubGroup) == 4, - "Invalid OpenCL memory scope enum definition"); Builder.defineMacro("__OPENCL_MEMORY_SCOPE_WORK_ITEM", "0"); Builder.defineMacro("__OPENCL_MEMORY_SCOPE_WORK_GROUP", "1"); Builder.defineMacro("__OPENCL_MEMORY_SCOPE_DEVICE", "2"); diff --git a/llvm/include/llvm/TargetParser/AtomicScope.h b/llvm/include/llvm/TargetParser/AtomicScope.h new file mode 100644 index 0000000000000..c338e6b91f27d --- /dev/null +++ b/llvm/include/llvm/TargetParser/AtomicScope.h @@ -0,0 +1,134 @@ +//===-- llvm/TargetParser/AtomicScope.h ---Atomic Scope--------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_TARGETPARSER_ATOMICSCOPE_H +#define LLVM_TARGETPARSER_ATOMICSCOPE_H + +#include "llvm/ADT/StringRef.h" +#include "llvm/ADT/StringSwitch.h" +#include "llvm/TargetParser/Triple.h" +#include <optional> +#include <utility> + +namespace llvm { + +/// Target-neutral memory synchronization scopes. +/// +/// The underlying values are ABI-sensitive and should not be changed. +enum class AtomicScope : unsigned { + System = 0, // __MEMORY_SCOPE_SYSTEM + Device = 1, // __MEMORY_SCOPE_DEVICE + Workgroup = 2, // __MEMORY_SCOPE_WRKGRP + Wavefront = 3, // __MEMORY_SCOPE_WVFRNT + Single = 4, // __MEMORY_SCOPE_SINGLE + Cluster = 5, // __MEMORY_SCOPE_CLUSTR +}; + +/// Returns the LLVM IR syncscope string that \p T uses to spell \p S. +inline std::optional<StringRef> +getAtomicScopeIRString(const Triple &T, AtomicScope S, + bool OneAddressSpace = false) { + if (T.isAMDGPU()) { + switch (S) { + case AtomicScope::Single: + return OneAddressSpace ? "singlethread-one-as" : "singlethread"; + case AtomicScope::Wavefront: + return OneAddressSpace ? "wavefront-one-as" : "wavefront"; + case AtomicScope::Workgroup: + return OneAddressSpace ? "workgroup-one-as" : "workgroup"; + case AtomicScope::Cluster: + return OneAddressSpace ? "cluster-one-as" : "cluster"; + case AtomicScope::Device: + return OneAddressSpace ? "agent-one-as" : "agent"; + case AtomicScope::System: + return OneAddressSpace ? "one-as" : ""; + } + return std::nullopt; + } + if (T.isNVPTX()) { + switch (S) { + case AtomicScope::Single: + return "singlethread"; + // NVPTX has no distinct wavefront/subgroup scope; it folds into block. + case AtomicScope::Wavefront: + case AtomicScope::Workgroup: + return "block"; + case AtomicScope::Cluster: + return "cluster"; + case AtomicScope::Device: + return "device"; + case AtomicScope::System: + return ""; + } + return std::nullopt; + } + if (T.isSPIRV()) { + switch (S) { + case AtomicScope::Single: + return "singlethread"; + case AtomicScope::Wavefront: + return "subgroup"; + // SPIR-V has no cluster scope; it folds into workgroup. + case AtomicScope::Cluster: + case AtomicScope::Workgroup: + return "workgroup"; + case AtomicScope::Device: + return "device"; + case AtomicScope::System: + return ""; + } + return std::nullopt; + } + return std::nullopt; +} + +/// Parses a target syncscope string into its abstract scope, the inverse of +/// getAtomicScopeIRString. Returns the scope and one-as variant if present. +inline std::optional<std::pair<AtomicScope, bool>> +parseAtomicScopeIRString(const Triple &T, StringRef Name) { + using Result = std::optional<std::pair<AtomicScope, bool>>; + auto Make = [](AtomicScope S, bool OneAS) -> std::pair<AtomicScope, bool> { + return {S, OneAS}; + }; + if (T.isAMDGPU()) + return StringSwitch<Result>(Name) + .Case("singlethread", Make(AtomicScope::Single, false)) + .Case("wavefront", Make(AtomicScope::Wavefront, false)) + .Case("workgroup", Make(AtomicScope::Workgroup, false)) + .Case("cluster", Make(AtomicScope::Cluster, false)) + .Case("agent", Make(AtomicScope::Device, false)) + .Case("", Make(AtomicScope::System, false)) + .Case("singlethread-one-as", Make(AtomicScope::Single, true)) + .Case("wavefront-one-as", Make(AtomicScope::Wavefront, true)) + .Case("workgroup-one-as", Make(AtomicScope::Workgroup, true)) + .Case("cluster-one-as", Make(AtomicScope::Cluster, true)) + .Case("agent-one-as", Make(AtomicScope::Device, true)) + .Case("one-as", Make(AtomicScope::System, true)) + .Default(std::nullopt); + if (T.isNVPTX()) + return StringSwitch<Result>(Name) + .Case("singlethread", Make(AtomicScope::Single, false)) + .Case("block", Make(AtomicScope::Workgroup, false)) + .Case("cluster", Make(AtomicScope::Cluster, false)) + .Case("device", Make(AtomicScope::Device, false)) + .Case("", Make(AtomicScope::System, false)) + .Default(std::nullopt); + if (T.isSPIRV()) + return StringSwitch<Result>(Name) + .Case("singlethread", Make(AtomicScope::Single, false)) + .Case("subgroup", Make(AtomicScope::Wavefront, false)) + .Case("workgroup", Make(AtomicScope::Workgroup, false)) + .Case("device", Make(AtomicScope::Device, false)) + .Case("", Make(AtomicScope::System, false)) + .Default(std::nullopt); + return std::nullopt; +} + +} // end namespace llvm + +#endif // LLVM_TARGETPARSER_ATOMICSCOPE_H diff --git a/llvm/lib/Target/AMDGPU/AMDGPUBarrierLatency.cpp b/llvm/lib/Target/AMDGPU/AMDGPUBarrierLatency.cpp index 5be42ec5e1589..a11bf2644b15b 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUBarrierLatency.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUBarrierLatency.cpp @@ -26,6 +26,7 @@ #include "SIInstrInfo.h" #include "llvm/CodeGen/ScheduleDAGInstrs.h" #include "llvm/Support/CommandLine.h" +#include "llvm/TargetParser/AtomicScope.h" using namespace llvm; @@ -44,17 +45,22 @@ class BarrierLatency : public ScheduleDAGMutation { public: BarrierLatency(MachineFunction *MF) { LLVMContext &Context = MF->getFunction().getContext(); + const Triple &TT = MF->getSubtarget<GCNSubtarget>().getTargetTriple(); + auto ScopeID = [&](AtomicScope Scope, bool OneAS) { + return Context.getOrInsertSyncScopeID( + *getAtomicScopeIRString(TT, Scope, OneAS)); + }; IgnoredScopes.insert(SyncScope::SingleThread); - IgnoredScopes.insert(Context.getOrInsertSyncScopeID("wavefront")); - IgnoredScopes.insert(Context.getOrInsertSyncScopeID("wavefront-one-as")); - IgnoredScopes.insert(Context.getOrInsertSyncScopeID("singlethread-one-as")); + IgnoredScopes.insert(ScopeID(AtomicScope::Wavefront, /*OneAS=*/false)); + IgnoredScopes.insert(ScopeID(AtomicScope::Wavefront, /*OneAS=*/true)); + IgnoredScopes.insert(ScopeID(AtomicScope::Single, /*OneAS=*/true)); const GCNSubtarget &ST = MF->getSubtarget<GCNSubtarget>(); bool TgSplit = ST.hasTgSplitSupport() && AMDGPU::isTgSplitEnabled(MF->getFunction()); if (!ST.requiresWaitOnWorkgroupReleaseFence(TgSplit)) { // Prior to GFX10 workgroup scope does not normally require waitcnts - IgnoredScopes.insert(Context.getOrInsertSyncScopeID("workgroup")); + IgnoredScopes.insert(ScopeID(AtomicScope::Workgroup, /*OneAS=*/false)); } } void apply(ScheduleDAGInstrs *DAG) override; diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMachineModuleInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMachineModuleInfo.cpp index d09b7cffe9f29..146fbc1e12268 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMachineModuleInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUMachineModuleInfo.cpp @@ -14,25 +14,32 @@ #include "AMDGPUMachineModuleInfo.h" #include "llvm/IR/Module.h" +#include "llvm/Target/TargetMachine.h" +#include "llvm/TargetParser/AtomicScope.h" using namespace llvm; AMDGPUMachineModuleInfo::AMDGPUMachineModuleInfo(const MachineModuleInfo &MMI) : MachineModuleInfoELF(MMI) { LLVMContext &CTX = MMI.getModule()->getContext(); - AgentSSID = CTX.getOrInsertSyncScopeID("agent"); - WorkgroupSSID = CTX.getOrInsertSyncScopeID("workgroup"); - WavefrontSSID = CTX.getOrInsertSyncScopeID("wavefront"); - ClusterSSID = CTX.getOrInsertSyncScopeID("cluster"); - SystemOneAddressSpaceSSID = - CTX.getOrInsertSyncScopeID("one-as"); - AgentOneAddressSpaceSSID = - CTX.getOrInsertSyncScopeID("agent-one-as"); + const Triple &TT = MMI.getTarget().getTargetTriple(); + + auto InsertScope = [&](AtomicScope Scope, bool OneAS) { + return CTX.getOrInsertSyncScopeID( + *getAtomicScopeIRString(TT, Scope, OneAS)); + }; + AgentSSID = InsertScope(AtomicScope::Device, /*OneAS=*/false); + WorkgroupSSID = InsertScope(AtomicScope::Workgroup, /*OneAS=*/false); + WavefrontSSID = InsertScope(AtomicScope::Wavefront, /*OneAS=*/false); + ClusterSSID = InsertScope(AtomicScope::Cluster, /*OneAS=*/false); + SystemOneAddressSpaceSSID = InsertScope(AtomicScope::System, /*OneAS=*/true); + AgentOneAddressSpaceSSID = InsertScope(AtomicScope::Device, /*OneAS=*/true); WorkgroupOneAddressSpaceSSID = - CTX.getOrInsertSyncScopeID("workgroup-one-as"); + InsertScope(AtomicScope::Workgroup, /*OneAS=*/true); WavefrontOneAddressSpaceSSID = - CTX.getOrInsertSyncScopeID("wavefront-one-as"); + InsertScope(AtomicScope::Wavefront, /*OneAS=*/true); SingleThreadOneAddressSpaceSSID = - CTX.getOrInsertSyncScopeID("singlethread-one-as"); - ClusterOneAddressSpaceSSID = CTX.getOrInsertSyncScopeID("cluster-one-as"); + InsertScope(AtomicScope::Single, /*OneAS=*/true); + ClusterOneAddressSpaceSSID = + InsertScope(AtomicScope::Cluster, /*OneAS=*/true); } diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index 9159e425dfffd..f867b15b05df5 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -48,6 +48,7 @@ #include "llvm/Support/CommandLine.h" #include "llvm/Support/KnownBits.h" #include "llvm/Support/ModRef.h" +#include "llvm/TargetParser/AtomicScope.h" #include "llvm/Transforms/Utils/LowerAtomic.h" #include <optional> @@ -20580,7 +20581,9 @@ SITargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *RMW) const { auto SSID = RMW->getSyncScopeID(); bool HasSystemScope = SSID == SyncScope::System || - SSID == RMW->getContext().getOrInsertSyncScopeID("one-as"); + SSID == RMW->getContext().getOrInsertSyncScopeID(*getAtomicScopeIRString( + getTargetMachine().getTargetTriple(), AtomicScope::System, + /*OneAddressSpace=*/true)); auto Op = RMW->getOperation(); switch (Op) { diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp index a0e3b1cc8e47f..af44ee053a734 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp @@ -35,6 +35,7 @@ #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/FormatVariadic.h" #include "llvm/Support/MathExtras.h" +#include "llvm/TargetParser/AtomicScope.h" #include <optional> using namespace llvm; @@ -57,7 +58,7 @@ namespace { struct NVPTXScopes { NVPTXScopes() = default; - NVPTXScopes(LLVMContext &C); + NVPTXScopes(LLVMContext &C, const Triple &T); NVPTX::Scope operator[](SyncScope::ID ID) const; bool empty() const; @@ -173,7 +174,8 @@ NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm, bool NVPTXDAGToDAGISel::runOnMachineFunction(MachineFunction &MF) { Subtarget = &MF.getSubtarget<NVPTXSubtarget>(); - Scopes = NVPTXScopes(MF.getFunction().getContext()); + Scopes = NVPTXScopes(MF.getFunction().getContext(), + MF.getTarget().getTargetTriple()); return SelectionDAGISel::runOnMachineFunction(MF); } @@ -1873,12 +1875,15 @@ bool NVPTXDAGToDAGISel::tryFence(SDNode *N) { return true; } -NVPTXScopes::NVPTXScopes(LLVMContext &C) : Context(&C) { - Scopes[C.getOrInsertSyncScopeID("singlethread")] = NVPTX::Scope::Thread; - Scopes[C.getOrInsertSyncScopeID("")] = NVPTX::Scope::System; - Scopes[C.getOrInsertSyncScopeID("block")] = NVPTX::Scope::Block; - Scopes[C.getOrInsertSyncScopeID("cluster")] = NVPTX::Scope::Cluster; - Scopes[C.getOrInsertSyncScopeID("device")] = NVPTX::Scope::Device; +NVPTXScopes::NVPTXScopes(LLVMContext &C, const Triple &T) : Context(&C) { + auto ScopeID = [&](AtomicScope Scope) { + return C.getOrInsertSyncScopeID(*getAtomicScopeIRString(T, Scope)); + }; + Scopes[ScopeID(AtomicScope::Single)] = NVPTX::Scope::Thread; + Scopes[ScopeID(AtomicScope::System)] = NVPTX::Scope::System; + Scopes[ScopeID(AtomicScope::Workgroup)] = NVPTX::Scope::Block; + Scopes[ScopeID(AtomicScope::Cluster)] = NVPTX::Scope::Cluster; + Scopes[ScopeID(AtomicScope::Device)] = NVPTX::Scope::Device; } NVPTX::Scope NVPTXScopes::operator[](SyncScope::ID ID) const { diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp index 513b3149dd28f..8c72ba3d8b316 100644 --- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp @@ -2591,8 +2591,8 @@ SPIRVEmitIntrinsicsImpl::visitAtomicCmpXchgInst(AtomicCmpXchgInst &I) { IRBuilder<> B(I.getParent()); B.SetInsertPoint(&I); SmallVector<Value *> Args(I.operands()); - Args.push_back(B.getInt32( - static_cast<uint32_t>(getMemScope(I.getContext(), I.getSyncScopeID())))); + Args.push_back(B.getInt32(static_cast<uint32_t>( + getMemScope(TM.getTargetTriple(), I.getContext(), I.getSyncScopeID())))); // Per SPIR-V spec atomic ops must combine the ordering bits with the // storage-class bit. const SPIRVSubtarget &ST = TM.getSubtarget<SPIRVSubtarget>(*I.getFunction()); diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp index 2ac7398479683..70cdc857cff5c 100644 --- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp @@ -2091,8 +2091,8 @@ bool SPIRVInstructionSelector::selectAtomicLoad(Register ResVReg, const MachineMemOperand &MemOp = **I.memoperands_begin(); assert(MemOp.isAtomic()); - uint32_t Scope = - static_cast<uint32_t>(getMemScope(Context, MemOp.getSyncScopeID())); + uint32_t Scope = static_cast<uint32_t>( + getMemScope(STI.getTargetTriple(), Context, MemOp.getSyncScopeID())); Register ScopeReg = buildI32Constant(Scope, I); AtomicOrdering AO = MemOp.getSuccessOrdering(); @@ -2240,8 +2240,8 @@ bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &I) const { const MachineMemOperand &MemOp = **I.memoperands_begin(); assert(MemOp.isAtomic()); - uint32_t Scope = - static_cast<uint32_t>(getMemScope(Context, MemOp.getSyncScopeID())); + uint32_t Scope = static_cast<uint32_t>( + getMemScope(STI.getTargetTriple(), Context, MemOp.getSyncScopeID())); Register ScopeReg = buildI32Constant(Scope, I); AtomicOrdering AO = MemOp.getSuccessOrdering(); @@ -2515,8 +2515,9 @@ bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg, unsigned NegateOpcode) const { assert(I.hasOneMemOperand()); const MachineMemOperand *MemOp = *I.memoperands_begin(); - uint32_t Scope = static_cast<uint32_t>(getMemScope( - GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID())); + uint32_t Scope = static_cast<uint32_t>( + getMemScope(STI.getTargetTriple(), GR.CurMF->getFunction().getContext(), + MemOp->getSyncScopeID())); Register ScopeReg = buildI32Constant(Scope, I); Register Ptr = I.getOperand(1).getReg(); @@ -2642,8 +2643,8 @@ bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const { uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO)); Register MemSemReg = buildI32Constant(MemSem, I); SyncScope::ID Ord = SyncScope::ID(I.getOperand(1).getImm()); - uint32_t Scope = static_cast<uint32_t>( - getMemScope(GR.CurMF->getFunction().getContext(), Ord)); + uint32_t Scope = static_cast<uint32_t>(getMemScope( + STI.getTargetTriple(), GR.CurMF->getFunction().getContext(), Ord)); Register ScopeReg = buildI32Constant(Scope, I); MachineBasicBlock &BB = *I.getParent(); BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpMemoryBarrier)) diff --git a/llvm/lib/Target/SPIRV/SPIRVUtils.cpp b/llvm/lib/Target/SPIRV/SPIRVUtils.cpp index 236da61192f1a..f5ef1b3ed92ce 100644 --- a/llvm/lib/Target/SPIRV/SPIRVUtils.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVUtils.cpp @@ -27,6 +27,7 @@ #include "llvm/IR/IntrinsicInst.h" #include "llvm/IR/IntrinsicsSPIRV.h" #include "llvm/Support/MathExtras.h" +#include "llvm/TargetParser/AtomicScope.h" #include <queue> #include <vector> @@ -448,18 +449,19 @@ SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord) { llvm_unreachable(nullptr); } -SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id) { +SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, + SyncScope::ID Id) { // Named by // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_scope_id. // We don't need aliases for Invocation and CrossDevice, as we already have // them covered by "singlethread" and "" strings respectively (see // implementation of LLVMContext::LLVMContext()). - static const llvm::SyncScope::ID SubGroup = - Ctx.getOrInsertSyncScopeID("subgroup"); - static const llvm::SyncScope::ID WorkGroup = - Ctx.getOrInsertSyncScopeID("workgroup"); - static const llvm::SyncScope::ID Device = - Ctx.getOrInsertSyncScopeID("device"); + auto ScopeID = [&](AtomicScope Scope) { + return Ctx.getOrInsertSyncScopeID(*getAtomicScopeIRString(TT, Scope)); + }; + static const llvm::SyncScope::ID SubGroup = ScopeID(AtomicScope::Wavefront); + static const llvm::SyncScope::ID WorkGroup = ScopeID(AtomicScope::Workgroup); + static const llvm::SyncScope::ID Device = ScopeID(AtomicScope::Device); if (Id == llvm::SyncScope::SingleThread) return SPIRV::Scope::Invocation; diff --git a/llvm/lib/Target/SPIRV/SPIRVUtils.h b/llvm/lib/Target/SPIRV/SPIRVUtils.h index eac06d1a7036e..b95f09eba95f1 100644 --- a/llvm/lib/Target/SPIRV/SPIRVUtils.h +++ b/llvm/lib/Target/SPIRV/SPIRVUtils.h @@ -37,6 +37,7 @@ class MachineIRBuilder; class MachineRegisterInfo; class Register; class StringRef; +class Triple; class SPIRVInstrInfo; class SPIRVSubtarget; class SPIRVGlobalRegistry; @@ -285,7 +286,8 @@ getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC); SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord); -SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id); +SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, + SyncScope::ID Id); // Find def instruction for the given ConstReg, walking through // spv_track_constant and ASSIGN_TYPE instructions. Updates ConstReg by def diff --git a/llvm/unittests/TargetParser/AtomicScopeTest.cpp b/llvm/unittests/TargetParser/AtomicScopeTest.cpp new file mode 100644 index 0000000000000..62df1a7a1a3d7 --- /dev/null +++ b/llvm/unittests/TargetParser/AtomicScopeTest.cpp @@ -0,0 +1,37 @@ +//===- AtomicScopeTest.cpp ------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "llvm/TargetParser/AtomicScope.h" +#include "gtest/gtest.h" + +using namespace llvm; + +namespace { + +TEST(AtomicScopeTest, AbiValues) { + EXPECT_EQ(static_cast<unsigned>(AtomicScope::System), 0u); + EXPECT_EQ(static_cast<unsigned>(AtomicScope::Device), 1u); + EXPECT_EQ(static_cast<unsigned>(AtomicScope::Workgroup), 2u); + EXPECT_EQ(static_cast<unsigned>(AtomicScope::Wavefront), 3u); + EXPECT_EQ(static_cast<unsigned>(AtomicScope::Single), 4u); + EXPECT_EQ(static_cast<unsigned>(AtomicScope::Cluster), 5u); +} + +TEST(AtomicScopeTest, UnknownStringIsNone) { + Triple NVPTX("nvptx64-nvidia-cuda"); + EXPECT_FALSE(parseAtomicScopeIRString(NVPTX, "bogus").has_value()); + EXPECT_FALSE(parseAtomicScopeIRString(NVPTX, "agent").has_value()); +} + +TEST(AtomicScopeTest, NonGPUTargetIsNone) { + Triple X86("x86_64-unknown-linux-gnu"); + EXPECT_FALSE(getAtomicScopeIRString(X86, AtomicScope::System).has_value()); + EXPECT_FALSE(parseAtomicScopeIRString(X86, "").has_value()); +} + +} // namespace diff --git a/llvm/unittests/TargetParser/CMakeLists.txt b/llvm/unittests/TargetParser/CMakeLists.txt index cb63202ea8028..9ef532603517b 100644 --- a/llvm/unittests/TargetParser/CMakeLists.txt +++ b/llvm/unittests/TargetParser/CMakeLists.txt @@ -4,6 +4,7 @@ set(LLVM_LINK_COMPONENTS ) add_llvm_unittest(TargetParserTests + AtomicScopeTest.cpp CSKYTargetParserTest.cpp Host.cpp NVPTXTargetParserTest.cpp >From 3e71bf56247b34db5a583bcf755ae1214aee4ae7 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Tue, 4 Aug 2026 16:22:22 -0500 Subject: [PATCH 2/3] comments --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 13 +++++---- llvm/include/llvm/TargetParser/AtomicScope.h | 22 ++++++++------- .../TargetParser/AtomicScopeTest.cpp | 28 +++++++++++++++++++ 3 files changed, 48 insertions(+), 15 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index a6d6c04589007..d10bfe178ac3b 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -386,10 +386,12 @@ static Value *emitFPIntBuiltin(CodeGenFunction &CGF, static inline StringRef mapScopeToSPIRV(const llvm::Triple &TargetTriple, StringRef AMDGCNScope) { static const llvm::Triple AMDGPU("amdgcn-amd-amdhsa"); - if (auto Parsed = llvm::parseAtomicScopeIRString(AMDGPU, AMDGCNScope)) - if (auto Str = llvm::getAtomicScopeIRString(TargetTriple, Parsed->first, - Parsed->second)) + if (auto Parsed = llvm::parseAtomicScopeIRString(AMDGPU, AMDGCNScope)) { + auto [Scope, IsSingleAddressSpace] = *Parsed; + if (auto Str = llvm::getAtomicScopeIRString(TargetTriple, Scope, + IsSingleAddressSpace)) return *Str; + } return AMDGCNScope; } @@ -451,10 +453,11 @@ void CodeGenFunction::ProcessOrderScopeAMDGCN(Value *Order, Value *Scope, AO = mapCABIAtomicOrdering(ord); // Some of the atomic builtins take the scope as a string name. + const llvm::Triple &TargetTriple = getTarget().getTriple(); StringRef scp; if (llvm::getConstantStringInfo(Scope, scp)) { - if (getTarget().getTriple().isSPIRV()) - scp = mapScopeToSPIRV(getTarget().getTriple(), scp); + if (TargetTriple.isSPIRV()) + scp = mapScopeToSPIRV(TargetTriple, scp); SSID = getLLVMContext().getOrInsertSyncScopeID(scp); return; } diff --git a/llvm/include/llvm/TargetParser/AtomicScope.h b/llvm/include/llvm/TargetParser/AtomicScope.h index c338e6b91f27d..9ab4c0e9dd5d1 100644 --- a/llvm/include/llvm/TargetParser/AtomicScope.h +++ b/llvm/include/llvm/TargetParser/AtomicScope.h @@ -32,21 +32,21 @@ enum class AtomicScope : unsigned { /// Returns the LLVM IR syncscope string that \p T uses to spell \p S. inline std::optional<StringRef> getAtomicScopeIRString(const Triple &T, AtomicScope S, - bool OneAddressSpace = false) { + bool IsSingleAddressSpace = false) { if (T.isAMDGPU()) { switch (S) { case AtomicScope::Single: - return OneAddressSpace ? "singlethread-one-as" : "singlethread"; + return IsSingleAddressSpace ? "singlethread-one-as" : "singlethread"; case AtomicScope::Wavefront: - return OneAddressSpace ? "wavefront-one-as" : "wavefront"; + return IsSingleAddressSpace ? "wavefront-one-as" : "wavefront"; case AtomicScope::Workgroup: - return OneAddressSpace ? "workgroup-one-as" : "workgroup"; + return IsSingleAddressSpace ? "workgroup-one-as" : "workgroup"; case AtomicScope::Cluster: - return OneAddressSpace ? "cluster-one-as" : "cluster"; + return IsSingleAddressSpace ? "cluster-one-as" : "cluster"; case AtomicScope::Device: - return OneAddressSpace ? "agent-one-as" : "agent"; + return IsSingleAddressSpace ? "agent-one-as" : "agent"; case AtomicScope::System: - return OneAddressSpace ? "one-as" : ""; + return IsSingleAddressSpace ? "one-as" : ""; } return std::nullopt; } @@ -88,12 +88,14 @@ getAtomicScopeIRString(const Triple &T, AtomicScope S, } /// Parses a target syncscope string into its abstract scope, the inverse of -/// getAtomicScopeIRString. Returns the scope and one-as variant if present. +/// getAtomicScopeIRString. Returns the scope and whether it is the single +/// address space variant. inline std::optional<std::pair<AtomicScope, bool>> parseAtomicScopeIRString(const Triple &T, StringRef Name) { using Result = std::optional<std::pair<AtomicScope, bool>>; - auto Make = [](AtomicScope S, bool OneAS) -> std::pair<AtomicScope, bool> { - return {S, OneAS}; + auto Make = [](AtomicScope S, + bool IsSingleAddressSpace) -> std::pair<AtomicScope, bool> { + return {S, IsSingleAddressSpace}; }; if (T.isAMDGPU()) return StringSwitch<Result>(Name) diff --git a/llvm/unittests/TargetParser/AtomicScopeTest.cpp b/llvm/unittests/TargetParser/AtomicScopeTest.cpp index 62df1a7a1a3d7..1160e411d0209 100644 --- a/llvm/unittests/TargetParser/AtomicScopeTest.cpp +++ b/llvm/unittests/TargetParser/AtomicScopeTest.cpp @@ -34,4 +34,32 @@ TEST(AtomicScopeTest, NonGPUTargetIsNone) { EXPECT_FALSE(parseAtomicScopeIRString(X86, "").has_value()); } +TEST(AtomicScopeTest, RoundTrip) { + static constexpr AtomicScope Scopes[] = { + AtomicScope::System, AtomicScope::Device, AtomicScope::Workgroup, + AtomicScope::Wavefront, AtomicScope::Single, AtomicScope::Cluster}; + const Triple Targets[] = {Triple("amdgcn-amd-amdhsa"), + Triple("nvptx64-nvidia-cuda"), + Triple("spirv64-unknown-unknown")}; + + for (const Triple &T : Targets) { + for (AtomicScope S : Scopes) { + for (bool IsSingleAddressSpace : {false, true}) { + auto Str = getAtomicScopeIRString(T, S, IsSingleAddressSpace); + if (!Str) + continue; + auto Parsed = parseAtomicScopeIRString(T, *Str); + ASSERT_TRUE(Parsed.has_value()) + << "target=" << T.str() << " string='" << Str->str() << "'"; + + auto ReEmitted = + getAtomicScopeIRString(T, Parsed->first, Parsed->second); + ASSERT_TRUE(ReEmitted.has_value()); + EXPECT_EQ(*ReEmitted, *Str) + << "target=" << T.str() << " string='" << Str->str() << "'"; + } + } + } +} + } // namespace >From 56e0295bbfbbb6af82988c393d64d002477b6c0d Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Wed, 5 Aug 2026 08:50:50 -0500 Subject: [PATCH 3/3] new test --- clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp b/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp index 547524580e566..6bb3c773f38db 100644 --- a/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp +++ b/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp @@ -153,6 +153,22 @@ void test_mixed() { __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "global"); __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "local", "global", "local", "local"); } + +// GCN-LABEL: define dso_local void @_Z12test_clusterv( +// GCN-SAME: ) #[[ATTR0]] { +// GCN-NEXT: entry: +// GCN-NEXT: fence syncscope("cluster") seq_cst +// GCN-NEXT: ret void +// +// AMDGCNSPIRV-LABEL: define spir_func void @_Z12test_clusterv( +// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] { +// AMDGCNSPIRV-NEXT: entry: +// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst +// AMDGCNSPIRV-NEXT: ret void +// +void test_cluster() { + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "cluster"); +} //. // GCN: [[META2]] = !{!"amdgpu-synchronize-as", !"local"} // GCN: [[META3]] = !{!"amdgpu-synchronize-as", !"global"} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
