https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/230829
Translate the AMDGPU scope names to SPIR-V names like classic codegen >From 15653c690a561d160ff8e6204dc62bf0beddadea Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Sat, 10 Oct 2026 20:40:04 +0200 Subject: [PATCH] [CIR][AMDGPU][SPIR-V] Implement __builtin_amdgcn_fence Translate the AMDGPU scope names to SPIR-V names like classic codegen --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 72 ++++++++++++++++++- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 69 +++++++++++++----- .../CIR/CodeGenHIP/builtins-amdgcn-fence.hip | 69 ++++++++++++++++++ 3 files changed, 191 insertions(+), 19 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 0ee8f18551e41..964757b362db6 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; @@ -1017,9 +1020,72 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_fence: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); + // 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{}; + } + cir::SyncScopeKind syncScope; + switch (parsed->first) { + case llvm::AtomicScope::Single: + syncScope = cir::SyncScopeKind::SingleThread; + break; + case llvm::AtomicScope::Wavefront: + syncScope = cir::SyncScopeKind::Wavefront; + break; + case llvm::AtomicScope::Workgroup: + syncScope = cir::SyncScopeKind::Workgroup; + break; + case llvm::AtomicScope::Cluster: + syncScope = cir::SyncScopeKind::Cluster; + break; + case llvm::AtomicScope::Device: + syncScope = cir::SyncScopeKind::Device; + break; + case llvm::AtomicScope::System: + syncScope = cir::SyncScopeKind::System; + break; + } + + // Sorted and unique, like MMRAMetadata::appendTags. + llvm::SmallVector<std::string> addrSpaces; + for (unsigned i = 2, e = expr->getNumArgs(); i < e; ++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; + for (const std::string &as : addrSpaces) + mmras.push_back(mlir::LLVM::MMRATagAttr::get( + &getMLIRContext(), "amdgpu-synchronize-as", as)); + + mlir::Location loc = getLoc(expr->getExprLoc()); + emitAtomicExprWithMemOrder( + expr->getArg(0), /*isStore=*/false, /*isLoad=*/false, + /*isFence=*/true, [&](cir::MemOrder order) { + auto fence = cir::AtomicFenceOp::create( + builder, loc, order, + cir::SyncScopeKindAttr::get(&getMLIRContext(), syncScope)); + 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 7cba2793cb870..000a7913e9913 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -56,6 +56,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; @@ -1301,21 +1302,50 @@ getLLVMMemOrder(std::optional<cir::MemOrder> memorder) { llvm_unreachable("unknown memory order"); } -static llvm::StringRef getLLVMSyncScope(cir::SyncScopeKind syncScope) { +static llvm::AtomicScope getAtomicScope(cir::SyncScopeKind syncScope) { switch (syncScope) { case cir::SyncScopeKind::SingleThread: - return "singlethread"; + case cir::SyncScopeKind::HIPSingleThread: + return llvm::AtomicScope::Single; + case cir::SyncScopeKind::Wavefront: + case cir::SyncScopeKind::HIPWavefront: + case cir::SyncScopeKind::OpenCLSubGroup: + return llvm::AtomicScope::Wavefront; case cir::SyncScopeKind::Workgroup: - return "block"; - default: - return ""; - } + case cir::SyncScopeKind::HIPWorkgroup: + case cir::SyncScopeKind::OpenCLWorkGroup: + return llvm::AtomicScope::Workgroup; + case cir::SyncScopeKind::Cluster: + case cir::SyncScopeKind::HIPCluster: + 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) { +getLLVMSyncScope(std::optional<cir::SyncScopeKind> syncScope, + mlir::ModuleOp moduleOp) { if (syncScope.has_value()) - return getLLVMSyncScope(*syncScope); + return getLLVMSyncScope(*syncScope, moduleOp); return std::nullopt; } @@ -1329,7 +1359,8 @@ mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite( rewriter, op.getLoc(), adaptor.getPtr(), expected, desired, getLLVMMemOrder(adaptor.getSuccOrder()), getLLVMMemOrder(adaptor.getFailOrder()), - getLLVMSyncScope(op.getSyncScope())); + getLLVMSyncScope(op.getSyncScope(), + op->getParentOfType<mlir::ModuleOp>())); cmpxchg.setAlignment(adaptor.getAlignment()); cmpxchg.setWeak(adaptor.getWeak()); @@ -1350,7 +1381,8 @@ mlir::LogicalResult CIRToLLVMAtomicXchgOpLowering::matchAndRewrite( mlir::ConversionPatternRewriter &rewriter) const { assert(!cir::MissingFeatures::atomicSyncScopeID()); mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(adaptor.getMemOrder()); - llvm::StringRef llvmSyncScope = getLLVMSyncScope(adaptor.getSyncScope()); + llvm::StringRef llvmSyncScope = getLLVMSyncScope( + op.getSyncScope(), op->getParentOfType<mlir::ModuleOp>()); rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>( op, mlir::LLVM::AtomicBinOp::xchg, adaptor.getPtr(), adaptor.getVal(), llvmOrder, llvmSyncScope, /*alignment=*/0, op.getIsVolatile()); @@ -1403,7 +1435,11 @@ mlir::LogicalResult CIRToLLVMAtomicFenceOpLowering::matchAndRewrite( mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(adaptor.getOrdering()); auto fence = mlir::LLVM::FenceOp::create(rewriter, op.getLoc(), llvmOrder); - fence.setSyncscope(getLLVMSyncScope(adaptor.getSyncscope())); + 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); @@ -1546,7 +1582,8 @@ mlir::LogicalResult CIRToLLVMAtomicFetchOpLowering::matchAndRewrite( } mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(op.getMemOrder()); - llvm::StringRef llvmSyncScope = getLLVMSyncScope(op.getSyncScope()); + llvm::StringRef llvmSyncScope = getLLVMSyncScope( + op.getSyncScope(), op->getParentOfType<mlir::ModuleOp>()); mlir::LLVM::AtomicBinOp llvmBinOp = getLLVMAtomicBinOp(op.getBinop(), isInt, isSignedInt); auto rmwVal = mlir::LLVM::AtomicRMWOp::create( @@ -2468,8 +2505,8 @@ mlir::LogicalResult CIRToLLVMLoadOpLowering::matchAndRewrite( assert(!cir::MissingFeatures::lowerModeOptLevel()); - std::optional<llvm::StringRef> llvmSyncScope = - getLLVMSyncScope(op.getSyncScope()); + std::optional<llvm::StringRef> llvmSyncScope = getLLVMSyncScope( + op.getSyncScope(), op->getParentOfType<mlir::ModuleOp>()); mlir::LLVM::LoadOp newLoad = mlir::LLVM::LoadOp::create( rewriter, op->getLoc(), llvmTy, adaptor.getAddr(), alignment, @@ -2531,8 +2568,8 @@ mlir::LogicalResult CIRToLLVMStoreOpLowering::matchAndRewrite( op.getValue().getType(), adaptor.getValue()); assert(!cir::MissingFeatures::opLoadStoreTbaa()); - std::optional<llvm::StringRef> llvmSyncScope = - getLLVMSyncScope(op.getSyncScope()); + std::optional<llvm::StringRef> llvmSyncScope = getLLVMSyncScope( + op.getSyncScope(), op->getParentOfType<mlir::ModuleOp>()); mlir::LLVM::StoreOp storeOp = mlir::LLVM::StoreOp::create( rewriter, op->getLoc(), value, adaptor.getAddr(), alignment, 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..2393763d016fa --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip @@ -0,0 +1,69 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir \ +// RUN: -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 amdgcn-amd-amdhsa -x hip -fclangir \ +// RUN: -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 amdgcn-amd-amdhsa -x hip \ +// RUN: -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 \ +// RUN: -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 \ +// RUN: -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 \ +// RUN: -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
