Author: Arseniy Obolenskiy Date: 2026-10-14T10:57:36+02:00 New Revision: 34aba97cbff67fadf0e2f68f2aabf11a1d78d966
URL: https://github.com/llvm/llvm-project/commit/34aba97cbff67fadf0e2f68f2aabf11a1d78d966 DIFF: https://github.com/llvm/llvm-project/commit/34aba97cbff67fadf0e2f68f2aabf11a1d78d966.diff LOG: [CIR][AMDGPU][SPIR-V] Complete __builtin_amdgcn_fence lowering implementation (#230829) Emit MMRA (memory model relaxation annotations) for address space arguments, reject unsupported scopes, translate AMDGPU scope names to SPIR-V ones on SPIR-V targets, matching classic codegen Added: clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index f28bfc67d9cc5..cf2a424c34148 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -12,11 +12,14 @@ #include "CIRGenFunction.h" +#include "mlir/Dialect/LLVMIR/LLVMAttrs.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/Value.h" #include "clang/Basic/TargetBuiltins.h" #include "llvm/IR/IntrinsicsAMDGPU.h" #include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/ErrorHandling.h" +#include "llvm/TargetParser/AtomicScope.h" using namespace clang; using namespace clang::CIRGen; @@ -68,6 +71,24 @@ static cir::SyncScopeKind decodeAMDGPUSyncScope(const Expr *arg) { .Default(cir::SyncScopeKind::System); } +static cir::SyncScopeKind convertAtomicScopeToCIR(llvm::AtomicScope scope) { + switch (scope) { + case llvm::AtomicScope::Single: + return cir::SyncScopeKind::SingleThread; + case llvm::AtomicScope::Wavefront: + return cir::SyncScopeKind::Wavefront; + case llvm::AtomicScope::Workgroup: + return cir::SyncScopeKind::Workgroup; + case llvm::AtomicScope::Cluster: + return cir::SyncScopeKind::Cluster; + case llvm::AtomicScope::Device: + return cir::SyncScopeKind::Device; + case llvm::AtomicScope::System: + return cir::SyncScopeKind::System; + } + llvm_unreachable("unknown atomic scope"); +} + /// Emit one of the AMDGPU raw hardware atomic builtins as a cir.atomic.fetch. static mlir::Value emitAMDGPUAtomicRMW(CIRGenFunction &cgf, const CallExpr *expr, @@ -1244,13 +1265,49 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_fence: { - CIRGenBuilderTy &b = getBuilder(); - cir::MemOrder mo = decodeAtomicOrder(expr->getArg(0), getContext()); - cir::SyncScopeKind syncScope = decodeAMDGPUSyncScope(expr->getArg(1)); - cir::SyncScopeKindAttr syncScopeAttr = - cir::SyncScopeKindAttr::get(b.getContext(), syncScope); - cir::AtomicFenceOp::create(b, getLoc(expr->getExprLoc()), mo, - syncScopeAttr); + // The scope uses AMDGPU spelling on every target, like mapScopeToSPIRV. + const llvm::Triple &triple = getTarget().getTriple(); + std::optional<std::pair<llvm::AtomicScope, bool>> parsed; + if (std::optional<std::string> name = + expr->getArg(1)->tryEvaluateString(getContext())) { + parsed = llvm::parseAtomicScopeIRString(llvm::Triple("amdgcn-amd-amdhsa"), + *name); + if (!parsed && triple.isSPIRV()) + parsed = llvm::parseAtomicScopeIRString(triple, *name); + } + // SPIR-V drops one-as, AMDGPU needs a CIR sync scope for it. + if (!parsed || (parsed->second && !triple.isSPIRV())) { + cgm.errorNYI(expr->getSourceRange(), "__builtin_amdgcn_fence scope"); + return mlir::Value{}; + } + // Sorted and unique, like MMRAMetadata::appendTags. + llvm::SmallVector<std::string> addrSpaces; + addrSpaces.reserve(expr->getNumArgs() - 2); + for (unsigned i = 2; i < expr->getNumArgs(); ++i) { + if (std::optional<std::string> as = + expr->getArg(i)->tryEvaluateString(getContext())) + addrSpaces.push_back(std::move(*as)); + else + cgm.error(expr->getExprLoc(), + "expected an address space name as a string literal"); + } + llvm::sort(addrSpaces); + addrSpaces.erase(llvm::unique(addrSpaces), addrSpaces.end()); + llvm::SmallVector<mlir::Attribute> mmras; + mmras.reserve(addrSpaces.size()); + for (const std::string &as : addrSpaces) + mmras.push_back(mlir::LLVM::MMRATagAttr::get( + &getMLIRContext(), "amdgpu-synchronize-as", as)); + + auto fence = cir::AtomicFenceOp::create( + builder, getLoc(expr->getExprLoc()), + decodeAtomicOrder(expr->getArg(0), getContext()), + cir::SyncScopeKindAttr::get(&getMLIRContext(), + convertAtomicScopeToCIR(parsed->first))); + if (!mmras.empty()) + fence->setAttr(mlir::LLVM::LLVMDialect::getMmraAttrName(), + mmras.size() == 1 ? mmras.front() + : builder.getArrayAttr(mmras)); return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_atomic_inc32: diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 0ef83e901511b..7b29331b762ad 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -57,6 +57,7 @@ #include "llvm/Support/TimeProfiler.h" #include "llvm/Support/VirtualFileSystem.h" #include "llvm/Support/raw_ostream.h" +#include "llvm/TargetParser/AtomicScope.h" using namespace cir; using namespace llvm; @@ -1363,42 +1364,45 @@ getLLVMMemOrder(std::optional<cir::MemOrder> memorder) { llvm_unreachable("unknown memory order"); } -static llvm::StringRef getLLVMSyncScope(cir::SyncScopeKind syncScope, - mlir::ModuleOp moduleOp) { +static llvm::AtomicScope getAtomicScope(cir::SyncScopeKind syncScope) { switch (syncScope) { case cir::SyncScopeKind::SingleThread: case cir::SyncScopeKind::HIPSingleThread: - return "singlethread"; - case cir::SyncScopeKind::System: - case cir::SyncScopeKind::HIPSystem: - case cir::SyncScopeKind::OpenCLAllSVMDevices: - return ""; - case cir::SyncScopeKind::Device: - case cir::SyncScopeKind::HIPAgent: - case cir::SyncScopeKind::OpenCLDevice: - return "agent"; - case cir::SyncScopeKind::Workgroup: - case cir::SyncScopeKind::HIPWorkgroup: - case cir::SyncScopeKind::OpenCLWorkGroup: { - // NVPTX uses "block" for workgroup sync scope; AMDGPU uses "workgroup". - auto tripleAttr = moduleOp->getAttrOfType<mlir::StringAttr>( - cir::CIRDialect::getTripleAttrName()); - if (tripleAttr && llvm::Triple(tripleAttr.getValue()).isNVPTX()) - return "block"; - return "workgroup"; - } + return llvm::AtomicScope::Single; case cir::SyncScopeKind::Wavefront: case cir::SyncScopeKind::HIPWavefront: - return "wavefront"; + case cir::SyncScopeKind::OpenCLSubGroup: + return llvm::AtomicScope::Wavefront; + case cir::SyncScopeKind::Workgroup: + case cir::SyncScopeKind::HIPWorkgroup: + case cir::SyncScopeKind::OpenCLWorkGroup: + return llvm::AtomicScope::Workgroup; case cir::SyncScopeKind::Cluster: case cir::SyncScopeKind::HIPCluster: - return "cluster"; - case cir::SyncScopeKind::OpenCLSubGroup: - return "sub_group"; + return llvm::AtomicScope::Cluster; + case cir::SyncScopeKind::Device: + case cir::SyncScopeKind::HIPAgent: + case cir::SyncScopeKind::OpenCLDevice: + return llvm::AtomicScope::Device; + case cir::SyncScopeKind::System: + case cir::SyncScopeKind::HIPSystem: + case cir::SyncScopeKind::OpenCLAllSVMDevices: + return llvm::AtomicScope::System; } llvm_unreachable("unknown sync scope"); } +static llvm::StringRef getLLVMSyncScope(cir::SyncScopeKind syncScope, + mlir::ModuleOp moduleOp) { + llvm::AtomicScope scope = getAtomicScope(syncScope); + if (auto triple = moduleOp->getAttrOfType<mlir::StringAttr>( + cir::CIRDialect::getTripleAttrName())) + if (std::optional<llvm::StringRef> name = llvm::getAtomicScopeIRString( + llvm::Triple(triple.getValue()), scope)) + return *name; + return scope == llvm::AtomicScope::Single ? "singlethread" : ""; +} + static std::optional<llvm::StringRef> getLLVMSyncScope(std::optional<cir::SyncScopeKind> syncScope, mlir::ModuleOp moduleOp) { @@ -1495,6 +1499,9 @@ mlir::LogicalResult CIRToLLVMAtomicFenceOpLowering::matchAndRewrite( auto fence = mlir::LLVM::FenceOp::create(rewriter, op.getLoc(), llvmOrder); fence.setSyncscope(getLLVMSyncScope(op.getSyncscope(), op->getParentOfType<mlir::ModuleOp>())); + if (mlir::Attribute mmra = + op->getAttr(mlir::LLVM::LLVMDialect::getMmraAttrName())) + fence->setAttr(mlir::LLVM::LLVMDialect::getMmraAttrName(), mmra); rewriter.replaceOp(op, fence); diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip new file mode 100644 index 0000000000000..c5e515c864d40 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip @@ -0,0 +1,63 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-cir %s -o %t.amd.cir +// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.amd.cir +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-llvm %s -o %t.amd-cir.ll +// RUN: FileCheck --check-prefixes=LLVM,AMD %s --input-file=%t.amd-cir.ll +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o %t.amd.ll +// RUN: FileCheck --check-prefixes=LLVM,AMD %s --input-file=%t.amd.ll +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-cir %s -o %t.spv.cir +// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.spv.cir +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-llvm %s -o %t.spv-cir.ll +// RUN: FileCheck --check-prefixes=LLVM,SPV %s --input-file=%t.spv-cir.ll +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o %t.spv.ll +// RUN: FileCheck --check-prefixes=LLVM,SPV %s --input-file=%t.spv.ll + +#define __device__ __attribute__((device)) + +// CIR-DAG: #[[MMRA_G:.*]] = #llvm.mmra_tag<"amdgpu-synchronize-as":"global"> +// CIR-DAG: #[[MMRA_L:.*]] = #llvm.mmra_tag<"amdgpu-synchronize-as":"local"> + +// CIR-LABEL: @_Z6scopesv +// LLVM-LABEL: @_Z6scopesv +__device__ void scopes() { + // CIR: cir.atomic.fence syncscope(device) seq_cst + // AMD: fence syncscope("agent") seq_cst + // SPV: fence syncscope("device") seq_cst + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent"); + // CIR: cir.atomic.fence syncscope(workgroup) seq_cst + // LLVM: fence syncscope("workgroup") seq_cst + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup"); + // CIR: cir.atomic.fence syncscope(wavefront) acq_rel + // AMD: fence syncscope("wavefront") acq_rel + // SPV: fence syncscope("subgroup") acq_rel + __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "wavefront"); + // CIR: cir.atomic.fence syncscope(single_thread) acquire + // LLVM: fence syncscope("singlethread") acquire + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "singlethread"); + // CIR: cir.atomic.fence syncscope(cluster) release + // AMD: fence syncscope("cluster") release + // SPV: fence syncscope("workgroup") release + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "cluster"); + // CIR: cir.atomic.fence syncscope(system) seq_cst + // LLVM: fence seq_cst + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, ""); +} + +// CIR-LABEL: @_Z9addrspacev +// LLVM-LABEL: @_Z9addrspacev +__device__ void addrspace() { + // CIR: cir.atomic.fence syncscope(workgroup) release {llvm.mmra = #[[MMRA_G]]} + // LLVM: fence syncscope("workgroup") release, !mmra ![[GLOBAL:[0-9]+]] + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup", "global"); + // CIR: cir.atomic.fence syncscope(workgroup) acquire {llvm.mmra = #[[MMRA_L]]} + // LLVM: fence syncscope("workgroup") acquire, !mmra ![[LOCAL:[0-9]+]] + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup", "local"); + // CIR: cir.atomic.fence syncscope(device) seq_cst {llvm.mmra = [#[[MMRA_G]], #[[MMRA_L]]]} + // AMD: fence syncscope("agent") seq_cst, !mmra ![[BOTH:[0-9]+]] + // SPV: fence syncscope("device") seq_cst, !mmra ![[BOTH:[0-9]+]] + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent", "local", "global"); +} + +// LLVM-DAG: ![[GLOBAL]] = !{!"amdgpu-synchronize-as", !"global"} +// LLVM-DAG: ![[LOCAL]] = !{!"amdgpu-synchronize-as", !"local"} +// LLVM-DAG: ![[BOTH]] = !{![[GLOBAL]], ![[LOCAL]]} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
