[clang] [CIR][AMDGPU][SPIR-V] Complete __builtin_amdgcn_fence lowering implementation (PR #230829)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 13 02:06:40 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/230829
>From 15653c690a561d160ff8e6204dc62bf0beddadea Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Sat, 10 Oct 2026 20:40:04 +0200
Subject: [PATCH 1/3] [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]]}
>From d31bb3889b486a40cd715aa32e27b73b370032af Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Sun, 11 Oct 2026 07:21:44 +0200
Subject: [PATCH 2/3] address comments
---
.../CIR/CodeGenHIP/builtins-amdgcn-fence.hip | 18 ++++++------------
1 file changed, 6 insertions(+), 12 deletions(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip
index 2393763d016fa..c5e515c864d40 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-fence.hip
@@ -1,21 +1,15 @@
// 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: %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 amdgcn-amd-amdhsa -x hip -fclangir \
-// RUN: -fcuda-is-device -emit-llvm %s -o %t.amd-cir.ll
+// 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 amdgcn-amd-amdhsa -x hip \
-// RUN: -fcuda-is-device -emit-llvm %s -o %t.amd.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 \
-// RUN: -fcuda-is-device -emit-cir %s -o %t.spv.cir
+// 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 \
-// RUN: -fcuda-is-device -emit-llvm %s -o %t.spv-cir.ll
+// 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 \
-// RUN: -fcuda-is-device -emit-llvm %s -o %t.spv.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))
>From 909e57d2c8efbc234606aca8e6481b8663c0779e Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 13 Oct 2026 11:06:26 +0200
Subject: [PATCH 3/3] Address comments
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 47 +++++++++----------
1 file changed, 23 insertions(+), 24 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 30b9b61777fa8..fd5b272fc993b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -61,6 +61,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,
@@ -1136,31 +1154,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
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) {
+ 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));
@@ -1171,6 +1168,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
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));
@@ -1178,7 +1176,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
auto fence = cir::AtomicFenceOp::create(
builder, getLoc(expr->getExprLoc()),
decodeAtomicOrder(expr->getArg(0), getContext()),
- cir::SyncScopeKindAttr::get(&getMLIRContext(), syncScope));
+ cir::SyncScopeKindAttr::get(&getMLIRContext(),
+ convertAtomicScopeToCIR(parsed->first)));
if (!mmras.empty())
fence->setAttr(mlir::LLVM::LLVMDialect::getMmraAttrName(),
mmras.size() == 1 ? mmras.front()
More information about the cfe-commits
mailing list