[clang] 34aba97 - [CIR][AMDGPU][SPIR-V] Complete __builtin_amdgcn_fence lowering implementation (#230829)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 14 01:57:43 PDT 2026
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]]}
More information about the cfe-commits
mailing list