[Mlir-commits] [mlir] [mlir][NVVM] Add support for few more fence Ops (PR #170251)
Pradeep Kumar
llvmlistbot at llvm.org
Tue Dec 2 22:55:52 PST 2025
https://github.com/schwarzschild-radius updated https://github.com/llvm/llvm-project/pull/170251
>From 570f3b8d207f4abf7902ee53c82a8b685920468f Mon Sep 17 00:00:00 2001
From: Pradeep Kumar <pradeepku at nvidia.com>
Date: Mon, 1 Dec 2025 07:14:22 +0000
Subject: [PATCH] [mlir][NVVM] Add support for few more fence Ops
This commit adds support for the following fence Ops:
- fence.sync_restrict
- fence.proxy.sync_restrict
The commit also moves memory.barrier into the Membar/Fence section and migrates fence.mbarrier.init to intrinsics
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 149 +++++++++++++-----
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 20 ++-
.../Dialect/NVVM/NVVMToLLVMIRTranslation.cpp | 35 ++++
.../Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 23 ---
mlir/test/Dialect/LLVMIR/nvvm.mlir | 7 -
.../Target/LLVMIR/nvvm/fence-invalid.mlir | 89 +++++++++++
mlir/test/Target/LLVMIR/nvvm/fence.mlir | 85 ++++++++++
mlir/test/Target/LLVMIR/nvvmir.mlir | 36 -----
8 files changed, 334 insertions(+), 110 deletions(-)
create mode 100644 mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
create mode 100644 mlir/test/Target/LLVMIR/nvvm/fence.mlir
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index a96d65d3fcacd..9f159e8593c3f 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -211,6 +211,27 @@ def SharedSpaceAttr : EnumAttr<NVVM_Dialect, SharedSpace, "shared_space"> {
let assemblyFormat = "`<` $value `>`";
}
+// Attrs describing the Memory Ordering Semantics
+def MemOrderKindWeak : I32EnumAttrCase<"WEAK", 0, "weak">;
+def MemOrderKindRelaxed : I32EnumAttrCase<"RELAXED", 1, "relaxed">;
+def MemOrderKindAcquire : I32EnumAttrCase<"ACQUIRE", 2, "acquire">;
+def MemOrderKindRelease : I32EnumAttrCase<"RELEASE", 3, "release">;
+def MemOrderKindAcqRel : I32EnumAttrCase<"ACQ_REL", 4, "acq_rel">;
+def MemOrderKindSC : I32EnumAttrCase<"SC", 5, "sc">;
+def MemOrderKindMMIO : I32EnumAttrCase<"MMIO", 6, "mmio">;
+def MemOrderKindVolatile : I32EnumAttrCase<"VOLATILE", 8, "volatile">;
+
+def MemOrderKind : I32EnumAttr<"MemOrderKind", "NVVM Memory Ordering kind",
+ [MemOrderKindWeak, MemOrderKindRelaxed, MemOrderKindAcquire,
+ MemOrderKindRelease, MemOrderKindAcqRel, MemOrderKindSC,
+ MemOrderKindMMIO, MemOrderKindVolatile]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def MemOrderKindAttr : EnumAttr<NVVM_Dialect, MemOrderKind, "mem_order"> {
+ let assemblyFormat = "`<` $value `>`";
+}
+
//===----------------------------------------------------------------------===//
// NVVM intrinsic operations
//===----------------------------------------------------------------------===//
@@ -1331,6 +1352,27 @@ def NVVM_ClusterWaitOp : NVVM_Op<"cluster.wait", [NVVMRequiresSM<90>]> {
let assemblyFormat = "attr-dict";
}
+//===----------------------------------------------------------------------===//
+// NVVM Member/Fence
+//===----------------------------------------------------------------------===//
+
+def NVVM_MembarOp : NVVM_Op<"memory.barrier">,
+ Arguments<(ins MemScopeKindAttr:$scope)> {
+ let summary = "Memory barrier operation";
+ let description = [{
+ `membar` operation guarantees that prior memory accesses requested by this
+ thread are performed at the specified `scope`, before later memory
+ operations requested by this thread following the membar instruction.
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar)
+ }];
+
+ let assemblyFormat = "$scope attr-dict";
+ let llvmBuilder = [{
+ createIntrinsicCall(builder, getMembarIntrinsicID($scope));
+ }];
+}
+
def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_fence_sc_cluster);
@@ -1338,6 +1380,38 @@ def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
let assemblyFormat = "attr-dict";
}
+def NVVM_FenceSyncRestrictOp : NVVM_Op<"fence.sync_restrict">,
+ Arguments<(ins MemOrderKindAttr:$order)> {
+ let summary = "Uni-directional thread fence operation";
+ let description = [{
+ The `nvvm.fence.sync_restrict` Op restricts the class of memory
+ operations for which the fence instruction provides the memory ordering guarantees.
+ `sync_restrict` restricts `acquire` memory semantics to `shared_cluster` and
+ `release` memory semantics to `shared_cta` with cluster scope.
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar)
+ }];
+
+ let assemblyFormat = "attr-dict";
+ let llvmBuilder = [{
+ createIntrinsicCall(builder, getFenceSyncRestrictID($order));
+ }];
+
+ let hasVerifier = 1;
+}
+
+def NVVM_FenceMbarrierInitOp : NVVM_Op<"fence.mbarrier.init"> {
+ let description = [{
+ Fence operation that applies on the prior nvvm.mbarrier.init
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar)
+ }];
+
+ let assemblyFormat = "attr-dict";
+ let llvmBuilder = [{
+ createIntrinsicCall(builder, llvm::Intrinsic::nvvm_fence_mbarrier_init_release_cluster);
+ }];
+}
+
def ProxyAlias : I32EnumAttrCase<"alias", 0, "alias">;
def ProxyAsync : I32EnumAttrCase<"async", 1, "async">;
def ProxyAsyncGlobal : I32EnumAttrCase<"async_global", 2, "async.global">;
@@ -1351,10 +1425,15 @@ def ProxyKind : I32EnumAttr<"ProxyKind", "Proxy kind",
}
def ProxyKindAttr : EnumAttr<NVVM_Dialect, ProxyKind, "proxy_kind"> {
+ let description = [{
+ ProxyKind attribute represents a memory proxy which is an abstract label
+ applied to a method of memory access. When two memory operations use distinct
+ methods of memory access, they are said to be different proxies.
+ }];
let assemblyFormat = "`<` $value `>`";
}
-def NVVM_FenceProxyOp : NVVM_PTXBuilder_Op<"fence.proxy">,
+def NVVM_FenceProxyOp : NVVM_Op<"fence.proxy">,
Arguments<(ins ProxyKindAttr:$kind,
OptionalAttr<SharedSpaceAttr>:$space)> {
let description = [{
@@ -1365,16 +1444,11 @@ def NVVM_FenceProxyOp : NVVM_PTXBuilder_Op<"fence.proxy">,
}];
let assemblyFormat = "attr-dict";
- let extraClassDefinition = [{
- std::string $cppClass::getPtx() {
- std::string ptx = "fence.proxy.";
- ptx += stringifyProxyKind(getKind());
- if(getKind() == NVVM::ProxyKind::async_shared)
- { ptx += "::"; ptx += stringifySharedSpace(getSpace().value()); }
- ptx += ";";
- return ptx;
- }
+
+ let llvmBuilder = [{
+ createIntrinsicCall(builder, getFenceProxyID($kind, $space));
}];
+
let hasVerifier = 1;
}
@@ -1411,23 +1485,6 @@ def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
let hasVerifier = 1;
}
-def NVVM_MembarOp : NVVM_Op<"memory.barrier">,
- Arguments<(ins MemScopeKindAttr:$scope)> {
- let summary = "Memory barrier operation";
- let description = [{
- `membar` operation guarantees that prior memory accesses requested by this
- thread are performed at the specified `scope`, before later memory
- operations requested by this thread following the membar instruction.
-
- [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar)
- }];
-
- let assemblyFormat = "$scope attr-dict";
- let llvmBuilder = [{
- createIntrinsicCall(builder, getMembarIntrinsicID($scope), {});
- }];
-}
-
def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
Arguments<(ins MemScopeKindAttr:$scope,
DefaultValuedAttr<ProxyKindAttr,
@@ -1454,6 +1511,28 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
let hasVerifier = 1;
}
+def NVVM_FenceProxySyncRestrictOp : NVVM_Op<"fence.proxy.sync_restrict">,
+ Arguments<(ins MemOrderKindAttr:$order,
+ DefaultValuedAttr<ProxyKindAttr, "ProxyKind::GENERIC">:$fromProxy,
+ DefaultValuedAttr<ProxyKindAttr, "ProxyKind::async">:$toProxy)> {
+ let summary = "Uni-directional proxy fence operation with sync_restrict";
+ let description = [{
+ The `nvvm.fence.proxy.sync_restrict` Op used to establish
+ ordering between a prior memory access performed between proxies. Currently,
+ the ordering is only supported between async and generic proxies. `sync_restrict`
+ restricts `acquire` memory semantics to `shared_cluster` and `release` memory
+ semantics to `shared_cta` with cluster scope.
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar)
+ }];
+
+ let assemblyFormat = "attr-dict";
+ let llvmBuilder = [{
+ createIntrinsicCall(builder, getFenceProxySyncRestrictID($order));
+ }];
+
+ let hasVerifier = 1;
+}
+
def SetMaxRegisterActionIncrease : I32EnumAttrCase<"increase", 0>;
def SetMaxRegisterActionDecrease : I32EnumAttrCase<"decrease", 1>;
def SetMaxRegisterAction : I32EnumAttr<"SetMaxRegisterAction", "NVVM set max register action",
@@ -1476,22 +1555,6 @@ def NVVM_SetMaxRegisterOp : NVVM_Op<"setmaxregister"> {
}];
}
-def NVVM_FenceMbarrierInitOp : NVVM_PTXBuilder_Op<"fence.mbarrier.init"> {
- let arguments = (ins );
- let description = [{
- Fence operation that applies on the prior nvvm.mbarrier.init
-
- [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar)
- }];
-
- let assemblyFormat = "attr-dict";
- let extraClassDefinition = [{
- std::string $cppClass::getPtx() {
- return std::string("fence.mbarrier_init.release.cluster;");
- }
- }];
-}
-
def ShflKindBfly : I32EnumAttrCase<"bfly", 0>;
def ShflKindUp : I32EnumAttrCase<"up", 1>;
def ShflKindDown : I32EnumAttrCase<"down", 2>;
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index ada4223ac12de..54a5406d486ea 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -2094,6 +2094,13 @@ bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
return true; // Has manual mapping
}
+LogicalResult NVVM::FenceSyncRestrictOp::verify() {
+ if (getOrder() != NVVM::MemOrderKind::ACQUIRE &&
+ getOrder() != NVVM::MemOrderKind::RELEASE)
+ return emitOpError("only acquire and release semantics are supported");
+ return success();
+}
+
LogicalResult NVVM::FenceProxyOp::verify() {
if (getKind() == NVVM::ProxyKind::TENSORMAP)
return emitOpError() << "tensormap proxy is not a supported proxy kind";
@@ -2116,7 +2123,6 @@ LogicalResult NVVM::FenceProxyAcquireOp::verify() {
if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
return emitOpError("uni-directional proxies only support tensormap "
"for to_proxy attribute");
-
return success();
}
@@ -2128,7 +2134,19 @@ LogicalResult NVVM::FenceProxyReleaseOp::verify() {
if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
return emitOpError("uni-directional proxies only support tensormap "
"for to_proxy attribute");
+ return success();
+}
+
+LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
+ if (getOrder() != NVVM::MemOrderKind::ACQUIRE &&
+ getOrder() != NVVM::MemOrderKind::RELEASE)
+ return emitOpError("only acquire and release semantics are supported");
+
+ if (getFromProxy() != NVVM::ProxyKind::GENERIC)
+ return emitOpError("only generic is support for from_proxy attribute");
+ if (getToProxy() != NVVM::ProxyKind::async)
+ return emitOpError("only async is supported for to_proxy attribute");
return success();
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
index cecff51e637a5..b7427a559fb79 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
@@ -411,6 +411,41 @@ getTcgen05StIntrinsicID(mlir::NVVM::Tcgen05LdStShape shape, uint32_t num) {
llvm_unreachable("unhandled tcgen05.st lowering");
}
+static llvm::Intrinsic::ID getFenceSyncRestrictID(NVVM::MemOrderKind order) {
+ return order == NVVM::MemOrderKind::ACQUIRE
+ ? llvm::Intrinsic::
+ nvvm_fence_acquire_sync_restrict_space_cluster_scope_cluster
+ : llvm::Intrinsic::
+ nvvm_fence_release_sync_restrict_space_cta_scope_cluster;
+}
+
+static llvm::Intrinsic::ID
+getFenceProxyID(NVVM::ProxyKind kind, std::optional<NVVM::SharedSpace> space) {
+ switch (kind) {
+ case NVVM::ProxyKind::alias:
+ return llvm::Intrinsic::nvvm_fence_proxy_alias;
+ case NVVM::ProxyKind::async:
+ return llvm::Intrinsic::nvvm_fence_proxy_async;
+ case NVVM::ProxyKind::async_global:
+ return llvm::Intrinsic::nvvm_fence_proxy_async_global;
+ case NVVM::ProxyKind::async_shared:
+ return *space == NVVM::SharedSpace::shared_cta
+ ? llvm::Intrinsic::nvvm_fence_proxy_async_shared_cta
+ : llvm::Intrinsic::nvvm_fence_proxy_async_shared_cluster;
+ default:
+ llvm_unreachable("unsupported proxy kind");
+ }
+}
+
+static llvm::Intrinsic::ID
+getFenceProxySyncRestrictID(NVVM::MemOrderKind order) {
+ return order == NVVM::MemOrderKind::ACQUIRE
+ ? llvm::Intrinsic::
+ nvvm_fence_proxy_async_generic_acquire_sync_restrict_space_cluster_scope_cluster
+ : llvm::Intrinsic::
+ nvvm_fence_proxy_async_generic_release_sync_restrict_space_cta_scope_cluster;
+}
+
namespace {
/// Implementation of the dialect interface that converts operations belonging
/// to the NVVM dialect to LLVM IR.
diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index fbf8d9efb3bc7..8fb36ace2c463 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -584,29 +584,6 @@ func.func @cp_async_bulk_wait_group() {
// -----
-func.func @fence_mbarrier_init() {
- //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.mbarrier_init.release.cluster;"
- nvvm.fence.mbarrier.init
- func.return
-}
-// -----
-
-func.func @fence_proxy() {
- //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.proxy.alias;", "" : () -> ()
- nvvm.fence.proxy { kind = #nvvm.proxy_kind<alias>}
- //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.proxy.async;", "" : () -> ()
- nvvm.fence.proxy { kind = #nvvm.proxy_kind<async>}
- //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.proxy.async.global;", "" : () -> ()
- nvvm.fence.proxy { kind = #nvvm.proxy_kind<async.global>}
- //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.proxy.async.shared::cta;", "" : () -> ()
- nvvm.fence.proxy { kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cta>}
- //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.proxy.async.shared::cluster;", "" : () -> ()
- nvvm.fence.proxy { kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cluster>}
- func.return
-}
-
-// -----
-
// CHECK-LABEL: @llvm_nvvm_barrier_arrive
// CHECK-SAME: (%[[barId:.*]]: i32, %[[numberOfThreads:.*]]: i32)
llvm.func @llvm_nvvm_barrier_arrive(%barID : i32, %numberOfThreads : i32) {
diff --git a/mlir/test/Dialect/LLVMIR/nvvm.mlir b/mlir/test/Dialect/LLVMIR/nvvm.mlir
index 6f67a50c1a946..579f0ac3ccad1 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -92,13 +92,6 @@ func.func @llvm_nvvm_cluster_wait() {
llvm.return
}
-// CHECK-LABEL: @llvm_nvvm_fence_sc_cluster
-func.func @llvm_nvvm_fence_sc_cluster() {
- // CHECK: nvvm.fence.sc.cluster
- nvvm.fence.sc.cluster
- llvm.return
-}
-
// CHECK-LABEL: @nvvm_shfl
func.func @nvvm_shfl(
%arg0 : i32, %arg1 : i32, %arg2 : i32,
diff --git a/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
new file mode 100644
index 0000000000000..22578b5581da4
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
@@ -0,0 +1,89 @@
+// RUN: mlir-translate --mlir-to-llvmir -verify-diagnostics -split-input-file %s
+
+llvm.func @fence_sync_restrict() {
+ // expected-error @below {{only acquire and release semantics are supported}}
+ nvvm.fence.sync_restrict {order = #nvvm.mem_order<weak>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_sync_restrict() {
+ // expected-error @below {{only acquire and release semantics are supported}}
+ nvvm.fence.sync_restrict {order = #nvvm.mem_order<mmio>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy() {
+ // expected-error @below {{tensormap proxy is not a supported proxy kind}}
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<tensormap>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy() {
+ // expected-error @below {{generic proxy not a supported proxy kind}}
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<generic>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy() {
+ // expected-error @below {{async_shared fence requires space attribute}}
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.shared>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy() {
+ // expected-error @below {{only async_shared fence can have space attribute}}
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<alias>, space = #nvvm.shared_space<cta>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy_release() {
+ // expected-error @below {{uni-directional proxies only support generic for from_proxy attribute}}
+ nvvm.fence.proxy.release #nvvm.mem_scope<cta> from_proxy = #nvvm.proxy_kind<alias> to_proxy = #nvvm.proxy_kind<tensormap>
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy_release() {
+ // expected-error @below {{uni-directional proxies only support tensormap for to_proxy attribute}}
+ nvvm.fence.proxy.release #nvvm.mem_scope<cta> from_proxy = #nvvm.proxy_kind<generic> to_proxy = #nvvm.proxy_kind<async>
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy_sync_restrict() {
+ // expected-error @below {{only acquire and release semantics are supported}}
+ nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<mmio>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy_sync_restrict() {
+ // expected-error @below {{only async is supported for to_proxy attribute}}
+ nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<acquire>, toProxy = #nvvm.proxy_kind<alias>,
+ fromProxy = #nvvm.proxy_kind<generic>}
+ llvm.return
+}
+
+// -----
+
+llvm.func @fence_proxy_sync_restrict() {
+ // expected-error @below {{only generic is support for from_proxy attribute}}
+ nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<acquire>, toProxy = #nvvm.proxy_kind<async>,
+ fromProxy = #nvvm.proxy_kind<tensormap>}
+ llvm.return
+}
diff --git a/mlir/test/Target/LLVMIR/nvvm/fence.mlir b/mlir/test/Target/LLVMIR/nvvm/fence.mlir
new file mode 100644
index 0000000000000..0ab4cb74b8f54
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/fence.mlir
@@ -0,0 +1,85 @@
+// RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
+
+// CHECK-LABEL: @llvm_nvvm_fence_sc_cluster
+llvm.func @llvm_nvvm_fence_sc_cluster() {
+ // CHECK: nvvm.fence.sc.cluster
+ nvvm.fence.sc.cluster
+ llvm.return
+}
+
+// CHECK-LABEL: @nvvm_fence_sync_restrict
+llvm.func @nvvm_fence_sync_restrict() {
+ // CHECK: call void @llvm.nvvm.fence.acquire.sync_restrict.space.cluster.scope.cluster()
+ nvvm.fence.sync_restrict {order = #nvvm.mem_order<acquire>}
+ // CHECK: call void @llvm.nvvm.fence.release.sync_restrict.space.cta.scope.cluster()
+ nvvm.fence.sync_restrict {order = #nvvm.mem_order<release>}
+ llvm.return
+}
+
+// CHECK-LABEL: @fence_mbarrier_init
+llvm.func @fence_mbarrier_init() {
+ // CHECK: call void @llvm.nvvm.fence.mbarrier_init.release.cluster()
+ nvvm.fence.mbarrier.init
+ llvm.return
+}
+
+// CHECK-LABEL: @nvvm_fence_proxy
+llvm.func @nvvm_fence_proxy() {
+ // CHECK: call void @llvm.nvvm.fence.proxy.alias()
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<alias>}
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.async()
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<async>}
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.async.global()
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.global>}
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.async.shared_cta()
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cta>}
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.async.shared_cluster()
+ nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cluster>}
+ llvm.return
+}
+
+// CHECK-LABEL: @nvvm_fence_proxy_sync_restrict
+llvm.func @nvvm_fence_proxy_sync_restrict() {
+ // CHECK: call void @llvm.nvvm.fence.proxy.async_generic.acquire.sync_restrict.space.cluster.scope.cluster()
+ nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<acquire>}
+ // CHECK: call void @llvm.nvvm.fence.proxy.async_generic.release.sync_restrict.space.cta.scope.cluster()
+ nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<release>}
+ llvm.return
+}
+
+// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_release
+llvm.func @nvvm_fence_proxy_tensormap_generic_release() {
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.cta()
+ nvvm.fence.proxy.release #nvvm.mem_scope<cta>
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.cluster()
+ nvvm.fence.proxy.release #nvvm.mem_scope<cluster>
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.gpu()
+ nvvm.fence.proxy.release #nvvm.mem_scope<gpu>
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.sys()
+ nvvm.fence.proxy.release #nvvm.mem_scope<sys>
+ llvm.return
+}
+
+// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_acquire
+llvm.func @nvvm_fence_proxy_tensormap_generic_acquire(%addr : !llvm.ptr) {
+ %c128 = llvm.mlir.constant(128) : i32
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr {{%[0-9]+}}, i32 128)
+ nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %c128
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr {{%[0-9]+}}, i32 128)
+ nvvm.fence.proxy.acquire #nvvm.mem_scope<cluster> %addr, %c128
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr {{%[0-9]+}}, i32 128)
+ nvvm.fence.proxy.acquire #nvvm.mem_scope<gpu> %addr, %c128
+
+ // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr {{%[0-9]+}}, i32 128)
+ nvvm.fence.proxy.acquire #nvvm.mem_scope<sys> %addr, %c128
+ llvm.return
+}
diff --git a/mlir/test/Target/LLVMIR/nvvmir.mlir b/mlir/test/Target/LLVMIR/nvvmir.mlir
index 5cba5c4fceefd..c4a69097692cb 100644
--- a/mlir/test/Target/LLVMIR/nvvmir.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir.mlir
@@ -699,42 +699,6 @@ llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant},
llvm.return
}
-
-// -----
-// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_release
-llvm.func @nvvm_fence_proxy_tensormap_generic_release() {
- %c128 = llvm.mlir.constant(128) : i32
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.cta()
- nvvm.fence.proxy.release #nvvm.mem_scope<cta>
-
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.cluster()
- nvvm.fence.proxy.release #nvvm.mem_scope<cluster>
-
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.gpu()
- nvvm.fence.proxy.release #nvvm.mem_scope<gpu>
-
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.sys()
- nvvm.fence.proxy.release #nvvm.mem_scope<sys>
- llvm.return
-}
-
-// -----
-// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_acquire
-llvm.func @nvvm_fence_proxy_tensormap_generic_acquire(%addr : !llvm.ptr) {
- %c128 = llvm.mlir.constant(128) : i32
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr {{%[0-9]+}}, i32 128)
- nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %c128
-
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr {{%[0-9]+}}, i32 128)
- nvvm.fence.proxy.acquire #nvvm.mem_scope<cluster> %addr, %c128
-
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr {{%[0-9]+}}, i32 128)
- nvvm.fence.proxy.acquire #nvvm.mem_scope<gpu> %addr, %c128
-
- // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr {{%[0-9]+}}, i32 128)
- nvvm.fence.proxy.acquire #nvvm.mem_scope<sys> %addr, %c128
- llvm.return
-}
// -----
// CHECK-LABEL: @nvvm_exit
More information about the Mlir-commits
mailing list