[Mlir-commits] [mlir] [mlir][gpu] Add address space modifier to gpu.barrier (PR #177425)
Krzysztof Drewniak
llvmlistbot at llvm.org
Thu Jan 22 12:48:34 PST 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/177425
>From caa143eb3f6245426defc44c933ba6f8a814e5c0 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Thu, 22 Jan 2026 00:31:34 +0000
Subject: [PATCH 1/3] [mlir][gpu] Add address space modifier to gpu.barrier
This is a takeover of PR ##110527
This commit adds an optional list of memory fences gpu.barrier,
allowing users to specify which memory scopes they wish to fence
explicitly, while leaving the default semantics (which are equivalent
to calling for a global and local fence by analogy to CUDA's
__syncthreads) unchanged. The new expanded semantics are implemented
for SPIR-V and for the AMDGPU backend.
See also
https://discourse.llvm.org/t/rfc-add-memory-scope-to-gpu-barrier/81021/2?u=fmarno,
where the default behavior of a gpu.barrier was hashed out (though
note that the examples based on VMCNT are outdated for AMDGPU in that
memory fences can now be annotated with the correct set of address
spaces).
This commit also deprecates amdgpu.lds_barrier for usecases that don't
involve targeting a gfx908.
Co-authored-by: Finlay Marno <finlay.marno at codeplay.com>
Assisted-by: Cursor/Claude code (tests and extending
amdgpu.lds_barrier pattern while copying it over)
---
mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPU.td | 9 +++
mlir/include/mlir/Dialect/GPU/IR/GPUBase.td | 2 +
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 26 +++++-
.../Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp | 29 +++++--
mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td | 2 +-
mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td | 4 -
.../GPUToROCDL/LowerGpuOpsToROCDLOps.cpp | 81 ++++++++++++++++++-
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 35 +++++++-
.../GPUToLLVMSPV/gpu-to-llvm-spv.mlir | 19 ++++-
.../GPUToROCDL/gpu-to-rocdl-barrier.mlir | 68 ++++++++++++++++
.../Conversion/GPUToROCDL/gpu-to-rocdl.mlir | 19 ++---
mlir/test/Dialect/GPU/canonicalize.mlir | 48 ++++++++++-
mlir/test/Dialect/GPU/ops.mlir | 6 ++
13 files changed, 311 insertions(+), 37 deletions(-)
create mode 100644 mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir
diff --git a/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPU.td b/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPU.td
index bf0711cc27922..a0ee4dc485561 100644
--- a/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPU.td
+++ b/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPU.td
@@ -956,6 +956,15 @@ def AMDGPU_PermlaneSwapOp : AMDGPU_Op<"permlane_swap", [Pure, AllTypesMatch<["re
def AMDGPU_LDSBarrierOp : AMDGPU_Op<"lds_barrier"> {
let summary = "Barrier that includes a wait for LDS memory operations.";
let description = [{
+ **DEPRECATION NOTICE**: Unless you need the inline-assembly-based workaround
+ for gfx908/MI-100, you should represent this pattern with the equivalent
+
+ ```mlir
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ ```
+
+ instead.
+
`amdgpu.lds_barrier` is both a barrier (all workitems in a workgroup must reach
the barrier before any of them may proceed past it) and a wait for all
operations that affect the Local Data Store (LDS) issued from that wrokgroup
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
index f0086158fb9b6..1a07d506b56d7 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
@@ -99,6 +99,8 @@ def GPU_AddressSpaceEnum : GPU_I32Enum<
def GPU_AddressSpaceAttr :
GPU_I32EnumAttr<"address_space", GPU_AddressSpaceEnum>;
+def GPU_AddressSpaceAttrArray : TypedArrayAttrBase<GPU_AddressSpaceAttr, "GPU Address Space array">;
+
//===----------------------------------------------------------------------===//
// GPU Types.
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index 51565aed92922..b6a1a4e1522df 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -1432,7 +1432,8 @@ def GPU_RotateOp : GPU_Op<
let hasVerifier = 1;
}
-def GPU_BarrierOp : GPU_Op<"barrier"> {
+def GPU_BarrierOp : GPU_Op<"barrier">,
+ Arguments<(ins OptionalAttr<GPU_AddressSpaceAttrArray> :$address_spaces)> {
let summary = "Synchronizes all work items of a workgroup.";
let description = [{
The `barrier` op synchronizes all work items of a workgroup. It is used
@@ -1442,17 +1443,36 @@ def GPU_BarrierOp : GPU_Op<"barrier"> {
gpu.barrier
```
- waits until all work items in the workgroup have reached this point
+ waits until all work items in the workgroup have reached the operation
and all memory accesses made by these work items prior to the op are
visible to all work items in the workgroup. Data hazards between work items
accessing the same memory can be avoided by synchronizing work items
in-between these accesses.
+ If the `memfence` attribute is specified, the set of memory accesses that must
+ by completed after the barrier resolves is limited to only those accesses that
+ read from or write to the specified address spaces (though accesses to other
+ address spaces may be completed as well, especially if a particular combination
+ of address spaces is not supported on a given backend). In particular,
+ specifying `memfence []` creates a barrier that is not required to affect
+ the visibility of any memory operations and is purely used for synchronizing
+ work items.
+
+ ```mlir
+ // only workgroup address spaces accesses required to be visible
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ // no memory accesses required to be visible
+ gpu.barrier memfence []
+ // all memory accesses required to be visible
+ gpu.barrier
+ ```
+
Either none or all work items of a workgroup need to execute this op
in convergence.
}];
- let assemblyFormat = "attr-dict";
+ let assemblyFormat = "(`memfence` $address_spaces^)? attr-dict";
let hasCanonicalizer = 1;
+ let builders = [OpBuilder<(ins)>];
}
def GPU_GPUModuleOp : GPU_Op<"module", [
diff --git a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
index 01b34337647ec..a270f52edd206 100644
--- a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
+++ b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
@@ -121,16 +121,31 @@ struct GPUBarrierConversion final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
lookupOrCreateSPIRVFn(moduleOp, funcName, flagTy, voidTy,
/*isMemNone=*/false, /*isConvergent=*/true);
- // Values used by SPIR-V backend to represent a combination of
- // `CLK_LOCAL_MEM_FENCE` and `CLK_GLOBAL_MEM_FENCE`.
- // See `llvm/lib/Target/SPIRV/SPIRVBuiltins.td`.
+ // Values used by SPIR-V backend to represent `CLK_LOCAL_MEM_FENCE` and
+ // `CLK_GLOBAL_MEM_FENCE`. See `llvm/lib/Target/SPIRV/SPIRVBuiltins.td`.
constexpr int64_t localMemFenceFlag = 1;
constexpr int64_t globalMemFenceFlag = 2;
- constexpr int64_t localGlobalMemFenceFlag =
- localMemFenceFlag | globalMemFenceFlag;
+ int64_t memFenceFlag = 0;
+ std::optional<ArrayAttr> addressSpaces = adaptor.getAddressSpaces();
+ if (addressSpaces) {
+ for (Attribute attr : addressSpaces.value()) {
+ auto addressSpace = cast<gpu::AddressSpaceAttr>(attr).getValue();
+ switch (addressSpace) {
+ case gpu::AddressSpace::Global:
+ memFenceFlag = memFenceFlag | globalMemFenceFlag;
+ break;
+ case gpu::AddressSpace::Workgroup:
+ memFenceFlag = memFenceFlag | localMemFenceFlag;
+ break;
+ case gpu::AddressSpace::Private:
+ break;
+ }
+ }
+ } else {
+ memFenceFlag = localMemFenceFlag | globalMemFenceFlag;
+ }
Location loc = op->getLoc();
- Value flag = LLVM::ConstantOp::create(rewriter, loc, flagTy,
- localGlobalMemFenceFlag);
+ Value flag = rewriter.create<LLVM::ConstantOp>(loc, flagTy, memFenceFlag);
rewriter.replaceOp(op, createSPIRVBuiltinCall(loc, rewriter, func, flag));
return success();
}
diff --git a/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td b/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
index f513bb1a0a826..0fcda38631a9b 100644
--- a/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
+++ b/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
@@ -17,6 +17,6 @@ include "mlir/IR/PatternBase.td"
include "mlir/Dialect/GPU/IR/GPUOps.td"
include "mlir/Dialect/LLVMIR/NVVMOps.td"
-def : Pat<(GPU_BarrierOp), (NVVM_Barrier0Op)>;
+def : Pat<(GPU_BarrierOp : $op $memory_fence), (NVVM_Barrier0Op)>;
#endif // MLIR_CONVERSION_GPUTONVVM_TD
diff --git a/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td b/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
index 8d2f30a9a1683..79ff2e0d39288 100644
--- a/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
+++ b/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
@@ -14,9 +14,5 @@
#define MLIR_CONVERSION_GPUTOROCDL_TD
include "mlir/IR/PatternBase.td"
-include "mlir/Dialect/GPU/IR/GPUOps.td"
-include "mlir/Dialect/LLVMIR/ROCDLOps.td"
-
-def : Pat<(GPU_BarrierOp), (ROCDL_BarrierOp)>;
#endif // MLIR_CONVERSION_GPUTOROCDL_TD
diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
index e5a40cfd090bd..69a83468cfa84 100644
--- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
+++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
@@ -313,6 +313,84 @@ struct GPUShuffleOpLowering : public ConvertOpToLLVMPattern<gpu::ShuffleOp> {
}
};
+struct GPUBarrierOpLowering final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
+ GPUBarrierOpLowering(const LLVMTypeConverter &converter,
+ amdgpu::Chipset chipset)
+ : ConvertOpToLLVMPattern<gpu::BarrierOp>(converter), chipset(chipset) {}
+
+ amdgpu::Chipset chipset;
+
+ LogicalResult
+ matchAndRewrite(gpu::BarrierOp op, gpu::BarrierOp::Adaptor adaptor,
+ ConversionPatternRewriter &rewriter) const override {
+ Location loc = op.getLoc();
+
+ // Analyze the address_spaces attribute to determine fence behavior.
+ bool fenceGlobal = false;
+ bool fenceLDS = false;
+ std::optional<ArrayAttr> addrSpacesToFence = op.getAddressSpaces();
+
+ if (addrSpacesToFence) {
+ for (auto spaceAttr :
+ addrSpacesToFence->getAsRange<gpu::AddressSpaceAttr>()) {
+ switch (spaceAttr.getValue()) {
+ case gpu::AddressSpace::Global:
+ fenceGlobal = true;
+ break;
+ case gpu::AddressSpace::Workgroup:
+ fenceLDS = true;
+ break;
+ case gpu::AddressSpace::Private:
+ break;
+ }
+ }
+ } else {
+ // Default semantics match __syncthreads() and fence both global and LDS.
+ fenceGlobal = true;
+ fenceLDS = true;
+ }
+
+ Attribute mmra;
+ if (fenceLDS && !fenceGlobal) {
+ mmra =
+ rewriter.getAttr<LLVM::MMRATagAttr>("amdgpu-synchronize-as", "local");
+ } else if (fenceGlobal && !fenceLDS) {
+ mmra = rewriter.getAttr<LLVM::MMRATagAttr>("amdgpu-synchronize-as",
+ "global");
+ }
+
+ constexpr llvm::StringLiteral scope = "workgroup";
+
+ bool emitFences = fenceGlobal || fenceLDS;
+ // Emit release fence if needed.
+ if (emitFences) {
+ auto relFence = LLVM::FenceOp::create(
+ rewriter, loc, LLVM::AtomicOrdering::release, scope);
+ if (mmra)
+ relFence->setDiscardableAttr(LLVM::LLVMDialect::getMmraAttrName(),
+ mmra);
+ }
+
+ if (chipset.majorVersion < 12) {
+ ROCDL::SBarrierOp::create(rewriter, loc);
+ } else {
+ ROCDL::BarrierSignalOp::create(rewriter, loc, -1);
+ ROCDL::BarrierWaitOp::create(rewriter, loc, -1);
+ }
+
+ if (emitFences) {
+ auto acqFence = LLVM::FenceOp::create(
+ rewriter, loc, LLVM::AtomicOrdering::acquire, scope);
+ if (mmra)
+ acqFence->setDiscardableAttr(LLVM::LLVMDialect::getMmraAttrName(),
+ mmra);
+ }
+
+ rewriter.eraseOp(op);
+ return success();
+ }
+};
+
/// Import the GPU Ops to ROCDL Patterns.
#include "GPUToROCDL.cpp.inc"
@@ -508,7 +586,8 @@ void mlir::populateGpuToROCDLConversionPatterns(
patterns.add<GPUShuffleOpLowering, GPULaneIdOpToROCDL,
GPUSubgroupBroadcastOpToROCDL>(converter);
- patterns.add<GPUSubgroupSizeOpToROCDL>(converter, chipset);
+ patterns.add<GPUSubgroupSizeOpToROCDL, GPUBarrierOpLowering>(converter,
+ chipset);
populateMathToROCDLConversionPatterns(converter, patterns, chipset);
}
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index 033a94e3f8fce..22a76c458d3c5 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -1472,11 +1472,35 @@ LogicalResult RotateOp::verify() {
/// Remove gpu.barrier after gpu.barrier, the threads are already synchronized!
static LogicalResult eraseRedundantGpuBarrierOps(BarrierOp op,
PatternRewriter &rewriter) {
- if (isa_and_nonnull<BarrierOp>(op->getNextNode())) {
- rewriter.eraseOp(op);
- return success();
+ auto nextOp = dyn_cast_or_null<BarrierOp>(op->getNextNode());
+ if (!nextOp)
+ return failure();
+
+ std::optional<ArrayAttr> thisMemfence = op.getAddressSpaces();
+ std::optional<ArrayAttr> nextMemfence = nextOp.getAddressSpaces();
+
+ if (thisMemfence) {
+ rewriter.modifyOpInPlace(op, [&]() {
+ if (!nextMemfence) {
+ op.removeAddressSpacesAttr();
+ } else {
+ // Fast path - mergge where the two barriers fence the same spaces.
+ if (*thisMemfence == *nextMemfence) {
+ return;
+ }
+ llvm::SmallSetVector<Attribute, 4> mergedSpaces;
+ for (Attribute attr : *thisMemfence)
+ mergedSpaces.insert(attr);
+ for (Attribute attr : *nextMemfence)
+ mergedSpaces.insert(attr);
+ op.setAddressSpacesAttr(
+ rewriter.getArrayAttr(mergedSpaces.takeVector()));
+ }
+ });
}
- return failure();
+
+ rewriter.eraseOp(nextOp);
+ return success();
}
void BarrierOp::getCanonicalizationPatterns(RewritePatternSet &results,
@@ -1484,6 +1508,9 @@ void BarrierOp::getCanonicalizationPatterns(RewritePatternSet &results,
results.add(eraseRedundantGpuBarrierOps);
}
+void BarrierOp::build(mlir::OpBuilder &odsBuilder,
+ mlir::OperationState &odsState) {}
+
//===----------------------------------------------------------------------===//
// GPUFuncOp
//===----------------------------------------------------------------------===//
diff --git a/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir b/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir
index 227287a9adcee..e56f3117e3b3c 100644
--- a/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir
+++ b/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir
@@ -215,14 +215,29 @@ gpu.module @barriers {
// CHECK-LABEL: gpu_barrier
func.func @gpu_barrier() {
- // CHECK: [[FLAGS:%.*]] = llvm.mlir.constant(3 : i32) : i32
- // CHECK: llvm.call spir_funccc @_Z7barrierj([[FLAGS]]) {
+ // CHECK: [[GLOBAL_AND_LOCAL_FLAG:%.*]] = llvm.mlir.constant(3 : i32) : i32
+ // CHECK: llvm.call spir_funccc @_Z7barrierj([[GLOBAL_AND_LOCAL_FLAG]]) {
// CHECK-SAME-DAG: no_unwind
// CHECK-SAME-DAG: convergent
// CHECK-SAME-DAG: will_return
// CHECK-NOT: memory_effects = #llvm.memory_effects
// CHECK-SAME: } : (i32) -> ()
gpu.barrier
+ // CHECK: [[GLOBAL_AND_LOCAL_FLAG2:%.*]] = llvm.mlir.constant(3 : i32) : i32
+ // CHECK: llvm.call spir_funccc @_Z7barrierj([[GLOBAL_AND_LOCAL_FLAG2]])
+ gpu.barrier memfence [#gpu.address_space<global>, #gpu.address_space<workgroup>]
+ // CHECK: [[LOCAL_FLAG:%.*]] = llvm.mlir.constant(1 : i32) : i32
+ // CHECK: llvm.call spir_funccc @_Z7barrierj([[LOCAL_FLAG]])
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ // CHECK: [[GLOBAL_FLAG:%.*]] = llvm.mlir.constant(2 : i32) : i32
+ // CHECK: llvm.call spir_funccc @_Z7barrierj([[GLOBAL_FLAG]])
+ gpu.barrier memfence [#gpu.address_space<global>]
+ // CHECK: [[NONE_FLAG:%.*]] = llvm.mlir.constant(0 : i32) : i32
+ // CHECK: llvm.call spir_funccc @_Z7barrierj([[NONE_FLAG]])
+ gpu.barrier memfence []
+ // CHECK: [[NONE_FLAG2:%.*]] = llvm.mlir.constant(0 : i32) : i32
+ // CHECK: llvm.call spir_funccc @_Z7barrierj([[NONE_FLAG2]])
+ gpu.barrier memfence [#gpu.address_space<private>]
return
}
}
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir
new file mode 100644
index 0000000000000..c1c7d41f66a32
--- /dev/null
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir
@@ -0,0 +1,68 @@
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950' --mlir-print-local-scope | FileCheck %s --check-prefixes=CHECK,GFX9
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx1201' --mlir-print-local-scope | FileCheck %s --check-prefixes=CHECK,GFX12
+
+gpu.module @test_module {
+// CHECK-LABEL: func @barrier_default()
+func.func @barrier_default() {
+ // CHECK: llvm.fence syncscope("workgroup") release{{$}}
+ // GFX9-NEXT: rocdl.s.barrier
+ // GFX12-NEXT: rocdl.s.barrier.signal id = -1
+ // GFX12-NEXT: rocdl.s.barrier.wait id = -1
+ // CHECK-NEXT: llvm.fence syncscope("workgroup") acquire{{$}}
+ gpu.barrier
+ func.return
+}
+
+// CHECK-LABEL: func @barrier_no_fence()
+func.func @barrier_no_fence() {
+ // GFX9-NEXT: rocdl.s.barrier
+ // GFX12-NEXT: rocdl.s.barrier.signal id = -1
+ // GFX12-NEXT: rocdl.s.barrier.wait id = -1
+ // CHECK-NOT: llvm.fence
+ gpu.barrier memfence []
+ func.return
+}
+
+// CHECK-LABEL: func @barrier_workgroup_only()
+func.func @barrier_workgroup_only() {
+ // CHECK: llvm.fence syncscope("workgroup") release {llvm.mmra = #llvm.mmra_tag<"amdgpu-synchronize-as":"local">}
+ // GFX9-NEXT: rocdl.s.barrier
+ // GFX12-NEXT: rocdl.s.barrier.signal id = -1
+ // GFX12-NEXT: rocdl.s.barrier.wait id = -1
+ // CHECK: llvm.fence syncscope("workgroup") acquire {llvm.mmra = #llvm.mmra_tag<"amdgpu-synchronize-as":"local">}
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ func.return
+}
+
+// CHECK-LABEL: func @barrier_global_only()
+func.func @barrier_global_only() {
+ // CHECK-NEXT: llvm.fence syncscope("workgroup") release {llvm.mmra = #llvm.mmra_tag<"amdgpu-synchronize-as":"global">}
+ // GFX9-NEXT: rocdl.s.barrier
+ // GFX12-NEXT: rocdl.s.barrier.signal id = -1
+ // GFX12-NEXT: rocdl.s.barrier.wait id = -1
+ // CHECK-NEXT: llvm.fence syncscope("workgroup") acquire {llvm.mmra = #llvm.mmra_tag<"amdgpu-synchronize-as":"global">}
+ gpu.barrier memfence [#gpu.address_space<global>]
+ func.return
+}
+
+// CHECK-LABEL: func @barrier_both()
+func.func @barrier_both() {
+ // CHECK-NEXT: llvm.fence syncscope("workgroup") release{{$}}
+ // GFX9-NEXT: rocdl.s.barrier
+ // GFX12-NEXT: rocdl.s.barrier.signal id = -1
+ // GFX12-NEXT: rocdl.s.barrier.wait id = -1
+ // CHECK-NEXT: llvm.fence syncscope("workgroup") acquire{{$}}
+ gpu.barrier memfence [#gpu.address_space<global>, #gpu.address_space<workgroup>]
+ func.return
+}
+
+// CHECK-LABEL: func @barrier_private_only()
+func.func @barrier_private_only() {
+ // GFX9-NEXT: rocdl.s.barrier
+ // GFX12-NEXT: rocdl.s.barrier.signal id = -1
+ // GFX12-NEXT: rocdl.s.barrier.wait id = -1
+ // CHECK-NOT: llvm.fence
+ gpu.barrier memfence [#gpu.address_space<private>]
+ func.return
+}
+}
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
index 2fd923743281e..6b20a442c30c5 100755
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
@@ -129,17 +129,6 @@ gpu.module @test_module {
// -----
-gpu.module @test_module {
- // CHECK-LABEL: func @gpu_sync()
- func.func @gpu_sync() {
- // CHECK: rocdl.barrier
- gpu.barrier
- func.return
- }
-}
-
-// -----
-
gpu.module @test_module {
// CHECK-LABEL: func @gpu_sqrt
func.func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) -> (f16, f32, f64) {
@@ -612,7 +601,9 @@ gpu.module @test_module {
// CHECK: llvm.select
// CHECK: llvm.shl
// CHECK: rocdl.ds_bpermute {{.*}}
- // CHECK: rocdl.barrier
+ // CHECK: llvm.fence
+ // CHECK: rocdl.s.barrier
+ // CHECK: llvm.fence
// CHECK: llvm.bitcast
// CHECK: llvm.fadd
%result = gpu.all_reduce add %arg0 uniform {} : (f32) -> (f32)
@@ -636,7 +627,9 @@ gpu.module @test_module {
// CHECK: llvm.select
// CHECK: llvm.shl
// CHECK: rocdl.ds_bpermute {{.*}}
- // CHECK: rocdl.barrier
+ // CHECK: llvm.fence
+ // CHECK: rocdl.s.barrier
+ // CHECK: llvm.fence
%result = gpu.all_reduce %arg0 uniform {
^bb(%lhs : i32, %rhs : i32):
%xor = arith.xori %lhs, %rhs : i32
diff --git a/mlir/test/Dialect/GPU/canonicalize.mlir b/mlir/test/Dialect/GPU/canonicalize.mlir
index 2589976c306d9..1283c1465ca47 100644
--- a/mlir/test/Dialect/GPU/canonicalize.mlir
+++ b/mlir/test/Dialect/GPU/canonicalize.mlir
@@ -13,9 +13,8 @@ func.func @fold_wait_op_test1() {
// -----
-// Erase duplicate barriers.
// CHECK-LABEL: func @erase_barriers
-// CHECK-NEXT: gpu.barrier
+// CHECK-NEXT: gpu.barrier{{$}}
// CHECK-NEXT: return
func.func @erase_barriers() {
gpu.barrier
@@ -23,6 +22,51 @@ func.func @erase_barriers() {
return
}
+// CHECK-LABEL: func @erase_barriers_first_full_fence
+// CHECK-NEXT: gpu.barrier{{$}}
+// CHECK-NEXT: return
+func.func @erase_barriers_first_full_fence() {
+ gpu.barrier
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ return
+}
+
+// CHECK-LABEL: func @erase_barriers_second_full_fence
+// CHECK-NEXT: gpu.barrier{{$}}
+// CHECK-NEXT: return
+func.func @erase_barriers_second_full_fence() {
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ gpu.barrier
+ return
+}
+
+// CHECK-LABEL: func @erase_barriers_merge_memfence
+// CHECK-NEXT: gpu.barrier memfence [#gpu.address_space<workgroup>, #gpu.address_space<global>]
+// CHECK-NEXT: return
+func.func @erase_barriers_merge_memfence() {
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ gpu.barrier memfence [#gpu.address_space<global>]
+ return
+}
+
+// CHECK-LABEL: func @erase_barriers_merge_memfence_same
+// CHECK-NEXT: gpu.barrier memfence [#gpu.address_space<workgroup>]
+// CHECK-NEXT: return
+func.func @erase_barriers_merge_memfence_same() {
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ return
+}
+
+// CHECK-LABEL: func @erase_barriers_empty_memfence
+// CHECK-NEXT: gpu.barrier memfence []
+// CHECK-NEXT: return
+func.func @erase_barriers_empty_memfence() {
+ gpu.barrier memfence []
+ gpu.barrier memfence []
+ return
+}
+
// -----
// Replace uses of gpu.wait op with its async dependency.
diff --git a/mlir/test/Dialect/GPU/ops.mlir b/mlir/test/Dialect/GPU/ops.mlir
index 7772e7a1681c4..1d05268ed4475 100644
--- a/mlir/test/Dialect/GPU/ops.mlir
+++ b/mlir/test/Dialect/GPU/ops.mlir
@@ -181,6 +181,12 @@ module attributes {gpu.container_module} {
%rotate, %pred4 = gpu.rotate %arg0, 3, 16 : f32
"gpu.barrier"() : () -> ()
+ gpu.barrier
+ gpu.barrier memfence [#gpu.address_space<workgroup>]
+ gpu.barrier memfence [#gpu.address_space<global>]
+ gpu.barrier memfence [#gpu.address_space<global>, #gpu.address_space<workgroup>]
+ gpu.barrier memfence [#gpu.address_space<private>]
+ gpu.barrier memfence []
"some_op"(%bIdX, %tIdX) : (index, index) -> ()
%42 = memref.load %arg1[%bIdX] : memref<?xf32, 1>
>From 68519334cd4fe47d565fc8d233c7bb7e92c6518e Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Thu, 22 Jan 2026 18:38:02 +0000
Subject: [PATCH 2/3] Whoops, old-style builder
---
mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
index a270f52edd206..a0eed19f01a8b 100644
--- a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
+++ b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
@@ -145,7 +145,7 @@ struct GPUBarrierConversion final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
memFenceFlag = localMemFenceFlag | globalMemFenceFlag;
}
Location loc = op->getLoc();
- Value flag = rewriter.create<LLVM::ConstantOp>(loc, flagTy, memFenceFlag);
+ Value flag = LLVM::ConstantOp::create(rewriter, loc, flagTy, memFenceFlag);
rewriter.replaceOp(op, createSPIRVBuiltinCall(loc, rewriter, func, flag));
return success();
}
>From 076e376c8a1f8cbe82e5762772d89ebd955059f8 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Thu, 22 Jan 2026 12:48:24 -0800
Subject: [PATCH 3/3] Comment style in .td examples
Co-authored-by: Jakub Kuderski <kubakuderski at gmail.com>
---
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index b6a1a4e1522df..3167388f57e7d 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -1459,11 +1459,11 @@ def GPU_BarrierOp : GPU_Op<"barrier">,
work items.
```mlir
- // only workgroup address spaces accesses required to be visible
+ // Only workgroup address spaces accesses required to be visible.
gpu.barrier memfence [#gpu.address_space<workgroup>]
- // no memory accesses required to be visible
+ // No memory accesses required to be visible.
gpu.barrier memfence []
- // all memory accesses required to be visible
+ // All memory accesses required to be visible.
gpu.barrier
```
More information about the Mlir-commits
mailing list