Author: Joseph Huber Date: 2026-08-05T11:53:24-05:00 New Revision: c338734884f52e807436e55f51899bf848a14773
URL: https://github.com/llvm/llvm-project/commit/c338734884f52e807436e55f51899bf848a14773 DIFF: https://github.com/llvm/llvm-project/commit/c338734884f52e807436e55f51899bf848a14773.diff LOG: [LLVM] Unify atomic scope handling between LLVM and Clang (#213810) 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. Added: llvm/include/llvm/TargetParser/AtomicScope.h llvm/unittests/TargetParser/AtomicScopeTest.cpp Modified: clang/include/clang/Basic/SyncScope.h clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp clang/lib/CodeGen/TargetInfo.h clang/lib/CodeGen/Targets/AMDGPU.cpp clang/lib/CodeGen/Targets/NVPTX.cpp clang/lib/CodeGen/Targets/SPIR.cpp clang/lib/Frontend/InitPreprocessor.cpp clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp llvm/lib/Target/AMDGPU/AMDGPUBarrierLatency.cpp llvm/lib/Target/AMDGPU/AMDGPUMachineModuleInfo.cpp llvm/lib/Target/AMDGPU/SIISelLowering.cpp llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp llvm/lib/Target/SPIRV/SPIRVUtils.cpp llvm/lib/Target/SPIRV/SPIRVUtils.h llvm/unittests/TargetParser/CMakeLists.txt Removed: ################################################################################ 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 b40059a1d73c0..e4be3869700cd 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,17 @@ 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)) { + auto [Scope, IsSingleAddressSpace] = *Parsed; + if (auto Str = llvm::getAtomicScopeIRString(TargetTriple, Scope, + IsSingleAddressSpace)) + return *Str; + } return AMDGCNScope; } @@ -446,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(scp); + if (TargetTriple.isSPIRV()) + scp = mapScopeToSPIRV(TargetTriple, scp); SSID = getLLVMContext().getOrInsertSyncScopeID(scp); return; } @@ -2013,10 +2021,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 b51996c6d88bf..9a8fcc546915e 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -511,32 +511,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 65b065dc818f5..349628c3439bf 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/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"} diff --git a/llvm/include/llvm/TargetParser/AtomicScope.h b/llvm/include/llvm/TargetParser/AtomicScope.h new file mode 100644 index 0000000000000..9ab4c0e9dd5d1 --- /dev/null +++ b/llvm/include/llvm/TargetParser/AtomicScope.h @@ -0,0 +1,136 @@ +//===-- 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 IsSingleAddressSpace = false) { + if (T.isAMDGPU()) { + switch (S) { + case AtomicScope::Single: + return IsSingleAddressSpace ? "singlethread-one-as" : "singlethread"; + case AtomicScope::Wavefront: + return IsSingleAddressSpace ? "wavefront-one-as" : "wavefront"; + case AtomicScope::Workgroup: + return IsSingleAddressSpace ? "workgroup-one-as" : "workgroup"; + case AtomicScope::Cluster: + return IsSingleAddressSpace ? "cluster-one-as" : "cluster"; + case AtomicScope::Device: + return IsSingleAddressSpace ? "agent-one-as" : "agent"; + case AtomicScope::System: + return IsSingleAddressSpace ? "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 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 IsSingleAddressSpace) -> std::pair<AtomicScope, bool> { + return {S, IsSingleAddressSpace}; + }; + 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 7b694e9e8719d..72774c7ea97b0 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> @@ -20582,7 +20583,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 5e416374dd829..ce993299096ad 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..b91ca685812c5 --- /dev/null +++ b/llvm/unittests/TargetParser/AtomicScopeTest.cpp @@ -0,0 +1,65 @@ +//===- 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()); +} + +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 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
